diff --git a/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu b/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu index ce0c410a1..a0d119c37 100644 --- a/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu +++ b/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu @@ -1085,6 +1085,23 @@ extern "C" __global__ void dqn_full_experience_kernel( const float* __restrict__ rms_v_gamma, /* [VALUE_H] */ const float* __restrict__ rms_a_gamma, /* [ADV_H] */ + /* ---- Fill simulation config ---- */ + float fill_median_spread, + float fill_median_vol, + float fill_ioc_fill_prob, + float fill_limit_fill_min, + float fill_limit_fill_max, + float fill_spread_cost_frac, + float fill_spread_capture_frac, + int fill_simulation_enabled, + + /* ---- DSR (Differential Sharpe Ratio) config ---- */ + int use_dsr, + float dsr_eta, + + /* ---- N-step returns config ---- */ + int n_steps, + /* ---- RNG states [N] ---- */ unsigned int* rng_states, @@ -1166,6 +1183,18 @@ extern "C" __global__ void dqn_full_experience_kernel( float ema_var = 1.0f; int ema_init = 0; + /* ---- DSR accumulators (per-episode) ---- */ + float dsr_A = 0.0f; + float dsr_B = 1e-8f; + int dsr_initialized = 0; + + /* ---- N-step ring buffer (per-episode) ---- */ + float nstep_ring[N_STEPS_MAX]; + int nstep_ring_idx = 0; + int nstep_ring_len = 0; + int effective_n = (n_steps > 0 && n_steps <= N_STEPS_MAX) ? n_steps : 1; + nstep_reset(nstep_ring, &nstep_ring_idx, &nstep_ring_len, effective_n); + /* ---- Main episode loop ---- */ for (int t = 0; t < L; t++) { int global_bar = ep_start + t; @@ -1309,13 +1338,54 @@ extern "C" __global__ void dqn_full_experience_kernel( action_counts[action_idx]++; total_action_count++; - out_actions[out_off] = action_idx; online_q_selected = q_values[action_idx]; + /* ---- Step 5b: Order routing + fill simulation ---- */ + /* Route action through spread/vol-dependent order type selection, + * then check fill probability. Unfilled orders override to Flat. */ + if (fill_simulation_enabled) { + int order_type, urgency; + route_order(spread, fill_median_spread, 0.0f /* vol from price */, fill_median_vol, + &order_type, &urgency); + + float norm_vol = (fill_median_vol > 0.0f) ? (spread / fill_median_spread) : 1.0f; + /* Clamp normalized vol to a reasonable range for fill probability */ + norm_vol = fminf(fmaxf(norm_vol, 0.0f), 3.0f); + + float cost_adj; + int filled = simulate_fill_check( + order_type, urgency, norm_vol, + spread * 10000.0f, /* convert spread fraction to bps */ + global_bar, action_idx, + fill_ioc_fill_prob, fill_limit_fill_min, fill_limit_fill_max, + fill_spread_cost_frac, fill_spread_capture_frac, + &cost_adj + ); + + if (!filled) { + action_idx = 2; /* Override to Flat — order didn't fill */ + } + } + + out_actions[out_off] = action_idx; + /* ---- Step 6: Portfolio simulation ---- */ float target_exposure = action_to_exposure(action_idx); float target_position = target_exposure * max_position; - float tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier; + /* Use order-type-dependent tx cost when fill sim enabled, + * otherwise fall back to fixed market rate. */ + float tx_rate; + if (fill_simulation_enabled) { + int order_type_for_cost, urgency_unused; + route_order(spread, fill_median_spread, 0.0f, fill_median_vol, + &order_type_for_cost, &urgency_unused); + float spread_bps = spread * 10000.0f; + tx_rate = fabsf(order_type_tx_cost(order_type_for_cost, spread_bps, + fill_spread_cost_frac, fill_spread_capture_frac)); + tx_rate *= tx_cost_multiplier; + } else { + tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier; + } /* Detect reversal (sign change) */ int is_reversal = (position > 0.0f && target_position < 0.0f) || @@ -1488,27 +1558,37 @@ extern "C" __global__ void dqn_full_experience_kernel( + diversity_scale * div_penalty + curiosity_scale * curiosity_reward; - /* ---- EMA reward normalization (matches CPU RewardNormalizer) ---- - * Normalize-then-update: use OLD stats to normalize, then update - * stats with the raw value. Clamp to [-3, 3] like the CPU path. */ - float raw_reward = combined_reward; - if (ema_init) { - float std = sqrtf(ema_var); - if (std > 1e-8f) - combined_reward = (combined_reward - ema_mean) / std; - combined_reward = fmaxf(-3.0f, fminf(3.0f, combined_reward)); - } - /* Update EMA running stats with the raw (un-normalized) value */ - if (!ema_init) { - ema_mean = raw_reward; - ema_var = 1.0f; - ema_init = 1; + /* DSR: replace EMA normalization when enabled */ + if (use_dsr) { + combined_reward = dsr_step(combined_reward, + &dsr_A, &dsr_B, &dsr_initialized, dsr_eta); } else { - ema_mean = reward_norm_alpha * raw_reward - + (1.0f - reward_norm_alpha) * ema_mean; - float diff = raw_reward - ema_mean; - ema_var = reward_norm_alpha * diff * diff - + (1.0f - reward_norm_alpha) * ema_var; + /* ---- EMA reward normalization (matches CPU RewardNormalizer) ---- */ + float raw_reward = combined_reward; + if (ema_init) { + float std = sqrtf(ema_var); + if (std > 1e-8f) + combined_reward = (combined_reward - ema_mean) / std; + combined_reward = fmaxf(-3.0f, fminf(3.0f, combined_reward)); + } + if (!ema_init) { + ema_mean = raw_reward; + ema_var = 1.0f; + ema_init = 1; + } else { + ema_mean = reward_norm_alpha * raw_reward + + (1.0f - reward_norm_alpha) * ema_mean; + float diff = raw_reward - ema_mean; + ema_var = reward_norm_alpha * diff * diff + + (1.0f - reward_norm_alpha) * ema_var; + } + } + + /* N-step: accumulate discounted return */ + if (effective_n > 1) { + combined_reward = nstep_push_and_sum( + nstep_ring, &nstep_ring_idx, &nstep_ring_len, + combined_reward, gamma, effective_n); } /* ---- Step 13: Episode done check ---- */ @@ -1565,6 +1645,14 @@ extern "C" __global__ void dqn_full_experience_kernel( ema_mean = 0.0f; ema_var = 1.0f; ema_init = 0; + + /* Reset DSR accumulators */ + dsr_A = 0.0f; + dsr_B = 1e-8f; + dsr_initialized = 0; + + /* Reset n-step ring */ + nstep_reset(nstep_ring, &nstep_ring_idx, &nstep_ring_len, effective_n); } } /* end if (!skip_data) */ } /* end timestep loop */ @@ -1710,6 +1798,16 @@ extern "C" __global__ void dqn_full_experience_kernel_warp( const float* __restrict__ rms_v_gamma, /* [VALUE_H] */ const float* __restrict__ rms_a_gamma, /* [ADV_H] */ + /* ---- Fill simulation config ---- */ + float fill_median_spread, + float fill_median_vol, + float fill_ioc_fill_prob, + float fill_limit_fill_min, + float fill_limit_fill_max, + float fill_spread_cost_frac, + float fill_spread_capture_frac, + int fill_simulation_enabled, + /* ---- RNG states [N] ---- */ unsigned int* rng_states, @@ -1955,13 +2053,49 @@ extern "C" __global__ void dqn_full_experience_kernel_warp( action_idx = __shfl_sync(0xFFFFFFFF, action_idx, 0); if (lane_id == 0) { - out_actions[out_off] = action_idx; online_q_selected = q_values[action_idx]; + /* ---- Step 5b: Order routing + fill simulation ---- */ + if (fill_simulation_enabled) { + int order_type, urgency; + route_order(spread, fill_median_spread, 0.0f, fill_median_vol, + &order_type, &urgency); + + float norm_vol = (fill_median_vol > 0.0f) ? (spread / fill_median_spread) : 1.0f; + norm_vol = fminf(fmaxf(norm_vol, 0.0f), 3.0f); + + float cost_adj; + int filled = simulate_fill_check( + order_type, urgency, norm_vol, + spread * 10000.0f, + global_bar, action_idx, + fill_ioc_fill_prob, fill_limit_fill_min, fill_limit_fill_max, + fill_spread_cost_frac, fill_spread_capture_frac, + &cost_adj + ); + + if (!filled) { + action_idx = 2; /* Override to Flat */ + } + } + + out_actions[out_off] = action_idx; + /* ---- Step 6: Portfolio simulation ---- */ float target_exposure = action_to_exposure(action_idx); float target_position = target_exposure * max_position; - float tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier; + float tx_rate; + if (fill_simulation_enabled) { + int order_type_for_cost, urgency_unused; + route_order(spread, fill_median_spread, 0.0f, fill_median_vol, + &order_type_for_cost, &urgency_unused); + float spread_bps = spread * 10000.0f; + tx_rate = fabsf(order_type_tx_cost(order_type_for_cost, spread_bps, + fill_spread_cost_frac, fill_spread_capture_frac)); + tx_rate *= tx_cost_multiplier; + } else { + tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier; + } int is_reversal = (position > 0.0f && target_position < 0.0f) || (position < 0.0f && target_position > 0.0f);