diff --git a/crates/ml/src/cuda_pipeline/common_device_functions.cuh b/crates/ml/src/cuda_pipeline/common_device_functions.cuh index 5afc671d4..9f3a1ff7c 100644 --- a/crates/ml/src/cuda_pipeline/common_device_functions.cuh +++ b/crates/ml/src/cuda_pipeline/common_device_functions.cuh @@ -777,3 +777,240 @@ __device__ float curiosity_inference( return fminf(mse, max_reward); } + +/* ------------------------------------------------------------------ */ +/* Differential Sharpe Ratio (Moody & Saffell 2001) */ +/* ------------------------------------------------------------------ */ + +/** + * Compute one DSR step and update EMA accumulators in-place. + * + * D_t = (B_prev * delta_A - 0.5 * A_prev * delta_B) / (B_prev - A_prev^2)^{3/2} + * + * where delta_A = r_t - A_prev, delta_B = r_t^2 - B_prev. + * Returns clamped DSR in [-5, 5]. + */ +__device__ __forceinline__ float dsr_step( + float r_t, + float* ema_A, /* EMA of returns (in/out) */ + float* ema_B, /* EMA of squared returns (in/out) */ + int* dsr_init, /* 0 = first call, set to 1 after init */ + float eta /* EMA decay rate (e.g. 0.01) */ +) { + if (!(*dsr_init)) { + *ema_A = r_t; + *ema_B = r_t * r_t + 1e-8f; /* avoid zero variance */ + *dsr_init = 1; + return 0.0f; /* No DSR on first step */ + } + + float a_prev = *ema_A; + float b_prev = *ema_B; + + /* Update EMAs */ + *ema_A = a_prev + eta * (r_t - a_prev); + *ema_B = b_prev + eta * (r_t * r_t - b_prev); + + /* DSR formula */ + float delta_a = r_t - a_prev; + float delta_b = r_t * r_t - b_prev; + float variance = b_prev - a_prev * a_prev; + + if (variance < 1e-12f) { + /* Fallback: scale raw PnL to approximate DSR magnitude range */ + return fmaxf(-5.0f, fminf(5.0f, r_t * 100.0f)); + } + + float denom = powf(variance, 1.5f); + float dsr = (b_prev * delta_a - 0.5f * a_prev * delta_b) / denom; + + return fmaxf(-5.0f, fminf(5.0f, dsr)); +} + +/* ------------------------------------------------------------------ */ +/* Order Routing (GPU-side OrderRouter::route) */ +/* ------------------------------------------------------------------ */ + +/** + * Deterministic order routing based on market microstructure. + * + * Maps (exposure_idx, spread, volatility) → (order_type, urgency). + * Matches CPU `OrderRouter::route()` in ml-core/src/order_router.rs exactly. + * + * Order types: 0=LimitMaker, 1=IoC, 2=Market + * Urgency: 0=Patient, 1=Normal, 2=Aggressive + */ +__device__ __forceinline__ void route_order( + float spread, + float median_spread, + float volatility, + float median_vol, + int* out_order_type, + int* out_urgency +) { + /* Order type selection based on spread vs median */ + int order; + if (median_spread <= 0.0f) { + order = 2; /* Market — no spread reference */ + } else if (spread < median_spread) { + order = 0; /* LimitMaker — tight spread, save costs */ + } else if (spread > 2.0f * median_spread) { + order = 2; /* Market — wide spread, ensure fill */ + } else { + order = 1; /* IoC — moderate spread */ + } + + /* Urgency selection based on volatility vs median */ + int urgency; + if (median_vol <= 0.0f) { + urgency = 1; /* Normal — no vol reference */ + } else if (volatility > 1.5f * median_vol) { + urgency = 2; /* Aggressive — high vol */ + } else if (volatility < 0.5f * median_vol) { + urgency = 0; /* Patient — low vol */ + } else { + urgency = 1; /* Normal */ + } + + *out_order_type = order; + *out_urgency = urgency; +} + +/** + * Order-type-dependent transaction cost rate. + * + * Matches CPU FillSimulator::spread_cost() cost structure: + * Market = spread_cost_frac * spread_bps / 10000 + * IoC = spread_cost_frac * 0.5 * spread_bps / 10000 + * LimitMaker = -(spread_capture_frac * spread_bps / 10000) [rebate] + * + * Returns cost as a fraction of trade value (positive = cost, negative = rebate). + */ +__device__ __forceinline__ float order_type_tx_cost( + int order_type, + float spread_bps, + float spread_cost_frac, + float spread_capture_frac +) { + switch (order_type) { + case 0: /* LimitMaker — maker rebate */ + return -(spread_bps * spread_capture_frac / 10000.0f); + case 1: /* IoC — half spread cost */ + return spread_bps * spread_cost_frac * 0.5f / 10000.0f; + case 2: /* Market — full spread cost */ + return spread_bps * spread_cost_frac / 10000.0f; + default: + return spread_bps * spread_cost_frac / 10000.0f; + } +} + +/* ------------------------------------------------------------------ */ +/* Fill Simulation (GPU-side FillSimulator::simulate_fill) */ +/* ------------------------------------------------------------------ */ + +/** + * Splitmix64-style deterministic hash: (step, action_idx) → [0.0, 1.0). + * + * Bit-exact match with CPU splitmix64_unit() in fill_simulator.rs. + * Uses the same Weyl-sequence seed combining and two-round finalizer. + */ +__device__ __forceinline__ float splitmix64_fill(int step, int action_idx) { + /* Combine step and action_idx with Weyl sequence constants */ + unsigned long long z = (unsigned long long)step * 0x9e3779b97f4a7c15ULL + + (unsigned long long)action_idx * 0x6a09e667f3bcc908ULL; + + /* Splitmix64 two-round finalizer */ + z = (z ^ (z >> 30)) * 0xbf58476d1ce4e5b9ULL; + z = (z ^ (z >> 27)) * 0x94d049bb133111ebULL; + z ^= (z >> 31); + + /* Convert to [0.0, 1.0) — upper 53 bits as double mantissa, then narrow to float */ + double d = (double)(z >> 11) / (double)(1ULL << 53); + return (float)d; +} + +/** + * Compute fill probability for a given order type and market conditions. + * + * Matches CPU FillSimulator::fill_probability() exactly. + * + * Parameters: + * order_type: 0=LimitMaker, 1=IoC, 2=Market + * urgency: 0=Patient, 1=Normal, 2=Aggressive + * normalized_vol: volatility / median_volatility (1.0 = average) + * ioc_fill_prob: base IoC fill probability (default 0.85) + * limit_fill_min: minimum limit fill probability (default 0.30) + * limit_fill_max: maximum limit fill probability (default 0.80) + */ +__device__ __forceinline__ float fill_probability( + int order_type, + int urgency, + float normalized_vol, + float ioc_fill_prob, + float limit_fill_min, + float limit_fill_max +) { + float base_prob; + switch (order_type) { + case 2: /* Market — always fills */ + base_prob = 1.0f; + break; + case 1: /* IoC */ + base_prob = ioc_fill_prob; + break; + case 0: /* LimitMaker — vol-dependent */ + default: { + float vol_clamped = fminf(fmaxf(normalized_vol, 0.0f), 3.0f); + float vol_factor = vol_clamped / 3.0f; + base_prob = limit_fill_min + vol_factor * (limit_fill_max - limit_fill_min); + break; + } + } + + /* Urgency modifier */ + float urgency_mod; + switch (urgency) { + case 0: urgency_mod = 0.85f; break; /* Patient */ + case 2: urgency_mod = 1.10f; break; /* Aggressive */ + default: urgency_mod = 1.0f; break; /* Normal */ + } + + float prob = base_prob * urgency_mod; + return fminf(fmaxf(prob, 0.0f), 1.0f); +} + +/** + * Simulate order fill using deterministic hash. + * + * Matches CPU FillSimulator::simulate_fill() exactly. + * + * Returns: 1 if filled, 0 if not filled. + * Writes cost adjustment to *out_cost_adj (positive=cost, negative=rebate, 0=no fill). + */ +__device__ __forceinline__ int simulate_fill_check( + int order_type, + int urgency, + float normalized_vol, + float spread_bps, + int step, + int action_idx, + float ioc_fill_prob, + float limit_fill_min, + float limit_fill_max, + float spread_cost_frac, + float spread_capture_frac, + float* out_cost_adj +) { + float prob = fill_probability(order_type, urgency, normalized_vol, + ioc_fill_prob, limit_fill_min, limit_fill_max); + float pseudo_random = splitmix64_fill(step, action_idx); + int filled = (pseudo_random < prob) ? 1 : 0; + + if (filled) { + *out_cost_adj = order_type_tx_cost(order_type, spread_bps, + spread_cost_frac, spread_capture_frac); + } else { + *out_cost_adj = 0.0f; + } + return filled; +}