From a57459e73694584ba7706dbc62eecdbf39c8150c Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 8 Apr 2026 12:44:23 +0200 Subject: [PATCH] feat(4branch): backtest env_step uses 4-branch position mapping Co-Authored-By: Claude Sonnet 4.6 --- .../src/cuda_pipeline/backtest_env_kernel.cu | 47 ++++++++++++------- .../cuda_pipeline/gpu_backtest_evaluator.rs | 5 ++ 2 files changed, 35 insertions(+), 17 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index f498eb27c..e78cd8258 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -25,12 +25,14 @@ __device__ void handle_capital_floor_breach( float new_capital, float step_ret, float* step_rewards, float* step_returns, int* actions_history, int* done_flags, int w, int max_len, int current_step, - int b0_size, int b1_size, int b2_size + int b0_size, int b1_size, int b2_size, int b3_size ) { step_rewards[w] = step_ret; step_returns[w * max_len + current_step] = step_ret; - // Full factored action encoding for Flat (exposure midpoint) - actions_history[w * max_len + current_step] = (b0_size / 2) * b1_size * b2_size; + // Flat = direction=1, magnitude=1, order=any(0), urgency=any(0) + // Flat action = 1 * b1*b2*b3 + 1 * b2*b3 + 0 * b3 + 0 + actions_history[w * max_len + current_step] = 1 * b1_size * b2_size * b3_size + + 1 * b2_size * b3_size; // Full episode restart — metrics captured before this reset portfolio_state[ps + 0] = new_capital; portfolio_state[ps + 1] = 0.0f; @@ -130,7 +132,7 @@ extern "C" __global__ void backtest_env_step( : 0.0f; handle_capital_floor_breach(portfolio_state, ps, liq_value, liq_ret, step_rewards, step_returns, actions_history, done_flags, - w, max_len, current_step, b0_size, b1_size, b2_size); + w, max_len, current_step, b0_size, b1_size, b2_size, b2_size); return; } @@ -229,7 +231,7 @@ extern "C" __global__ void backtest_env_step( : 0.0f; handle_capital_floor_breach(portfolio_state, ps, new_value, step_ret, step_rewards, step_returns, actions_history, done_flags, - w, max_len, current_step, b0_size, b1_size, b2_size); + w, max_len, current_step, b0_size, b1_size, b2_size, b2_size); return; } @@ -298,6 +300,7 @@ extern "C" __global__ void backtest_env_step_batch( int b0_size, int b1_size, int b2_size, + int b3_size, int min_hold_bars, float contract_multiplier, float margin_pct @@ -347,7 +350,7 @@ extern "C" __global__ void backtest_env_step_batch( : 0.0f; handle_capital_floor_breach(portfolio_state, ps, liq_value, liq_ret, step_rewards, step_returns, actions_history, done_flags, - w, max_len, current_step, b0_size, b1_size, b2_size); + w, max_len, current_step, b0_size, b1_size, b2_size, b3_size); /* Bug 9 fix: return instead of break to avoid tail write */ return; } @@ -359,9 +362,10 @@ extern "C" __global__ void backtest_env_step_batch( int action_val = chunked_actions[s * n_windows + w]; /* ── Decode action ────────────────────────────────────────────── */ - int exposure_idx = decode_exposure_index(action_val, b0_size, b1_size, b2_size); - float target_exposure = compute_target_position(exposure_idx, b0_size, max_position); - int order_type_idx = decode_order_type(action_val, b1_size, b2_size); + int dir_idx = decode_direction_4b(action_val, b1_size, b2_size, b3_size); + int mag_idx = decode_magnitude_4b(action_val, b1_size, b2_size, b3_size); + float target_exposure = compute_target_position_4branch(dir_idx, mag_idx, max_position); + int order_type_idx = decode_order_4b(action_val, b2_size, b3_size); /* ── Margin-aware position cap ────────────────────────────────── */ float margin_per_contract = close * contract_multiplier * margin_pct; @@ -439,7 +443,7 @@ extern "C" __global__ void backtest_env_step_batch( ? (new_value - value) / value : 0.0f; handle_capital_floor_breach(portfolio_state, ps, new_value, step_ret, step_rewards, step_returns, actions_history, done_flags, - w, max_len, current_step, b0_size, b1_size, b2_size); + w, max_len, current_step, b0_size, b1_size, b2_size, b3_size); return; /* Exit kernel — episode over */ } @@ -457,13 +461,22 @@ extern "C" __global__ void backtest_env_step_batch( step_rewards[w] = step_ret; step_returns[w * max_len + current_step] = step_ret; { - float f_actual_frac = position / fmaxf(max_position, 0.01f); - int actual_exp_idx = (int)roundf((f_actual_frac + 1.0f) * 0.5f * (float)(b0_size - 1)); - if (actual_exp_idx < 0) actual_exp_idx = 0; - if (actual_exp_idx >= b0_size) actual_exp_idx = b0_size - 1; - int orig_order = (b1_size > 0 && b2_size > 0) ? (action_val / b2_size) % b1_size : 0; - int orig_urgency = (b2_size > 0) ? action_val % b2_size : 0; - int actual_action = actual_exp_idx * b1_size * b2_size + orig_order * b2_size + orig_urgency; + /* Re-encode actual position as 4-branch factored action. + * Recover direction and magnitude from position, preserve order/urgency. */ + float pos_sign = position; + int actual_dir = (pos_sign > 0.001f) ? 2 : (pos_sign < -0.001f) ? 0 : 1; + float abs_pos = fabsf(position) / fmaxf(max_position, 0.01f); + int actual_mag; + if (abs_pos < 0.375f) actual_mag = 0; /* ~0.25 → Quarter */ + else if (abs_pos < 0.75f) actual_mag = 1; /* ~0.50 → Half */ + else actual_mag = 2; /* ~1.00 → Full */ + if (actual_dir == 1) actual_mag = 0; /* Flat: magnitude irrelevant */ + int orig_order = decode_order_4b(action_val, b2_size, b3_size); + int orig_urgency = decode_urgency_4b(action_val, b3_size); + int actual_action = actual_dir * b1_size * b2_size * b3_size + + actual_mag * b2_size * b3_size + + orig_order * b3_size + + orig_urgency; actions_history[w * max_len + current_step] = actual_action; } } /* end step loop */ diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index e82ceb23e..e62e82b0a 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -319,6 +319,7 @@ pub struct GpuBacktestEvaluator { b0_size: i32, b1_size: i32, b2_size: i32, + b3_size: i32, /// Q-value gap threshold for trade conviction filter (0.0 = disabled). q_gap_threshold: f32, @@ -605,6 +606,7 @@ impl GpuBacktestEvaluator { b0_size: 0, b1_size: 0, b2_size: 0, + b3_size: 0, q_gap_threshold: 0.0, rng_states: None, q_gaps_buf: None, @@ -1028,6 +1030,7 @@ impl GpuBacktestEvaluator { let b0 = dqn_cfg.branch_0_size as i32; let b1 = dqn_cfg.branch_1_size as i32; let b2 = dqn_cfg.branch_2_size as i32; + let b3 = dqn_cfg.branch_3_size as i32; let v_min_f = dqn_cfg.v_min; let v_max_f = dqn_cfg.v_max; let state_row_bytes = n * state_dim * std::mem::size_of::(); @@ -1178,6 +1181,7 @@ impl GpuBacktestEvaluator { .arg(&b0) .arg(&b1) .arg(&b2) + .arg(&b3) .arg(&self.config.min_hold_bars) .arg(&self.config.contract_multiplier) .arg(&self.config.margin_pct) @@ -1516,6 +1520,7 @@ impl GpuBacktestEvaluator { self.b0_size = dqn_cfg.branch_0_size as i32; self.b1_size = dqn_cfg.branch_1_size as i32; self.b2_size = dqn_cfg.branch_2_size as i32; + self.b3_size = dqn_cfg.branch_3_size as i32; self.cublas_forward = Some(cublas); self.cublas_params_flat = Some(params_flat); self.cublas_h_s1 = Some(h_s1);