diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index f38c54a07..0e1ed5296 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -486,34 +486,6 @@ extern "C" __global__ void backtest_env_step_batch( &actual_mag, &trail_triggered ); - - /* DIAG (val-collapse): sample bars across the window for window 0. - * First 10 bars = initial behavior. Then every 100 bars from step 100 - * to 1000 (chunk 1-15 transition zone), every 1000 bars to 10000, - * every 10000 bars to 200000. Plus a "mismatch trigger" — print on - * every bar where action picks Long/Short but actual_dir is not - * Long/Short (the gate event). Limit total to first 200 mismatches - * to keep printf buffer manageable. Removed once gate identified. */ - int dir_dbg = action_val / (b1_size * b2_size * b3_size); - int mag_dbg = (action_val / (b2_size * b3_size)) % b1_size; - bool sample_print = (current_step < 10 - || (current_step < 1000 && current_step % 100 == 0) - || (current_step < 10000 && current_step % 1000 == 0) - || current_step % 10000 == 0 - || current_step == 213000 - || current_step == 214000); - bool mismatch = (dir_dbg == 0 || dir_dbg == 2) /* picked Short or Long */ - && actual_dir != dir_dbg - && current_step >= 10 - && current_step % 137 == 0; /* prime stride to avoid pattern */ - if (w == 0 && (sample_print || mismatch)) { - printf("[VALDIAG step=%d] act=%d dir=%d mag=%d close=%.2f maxpos=%.4f " - "prevpos=%.4f pos_post=%.4f actual_dir=%d actual_mag=%d trail=%d " - "conv=%.3f health=%.3f value=%.2f maxeq=%.2f cash=%.2f\n", - current_step, action_val, dir_dbg, mag_dbg, close, max_position, - prev_position, position, actual_dir, actual_mag, trail_triggered, - conviction_core, health_k_batch, value, max_equity, cash); - } (void)prev_position; (void)prev_position_sign; (void)trail_triggered; /* ── Post-trade floor check: catch intra-step breaches ────────── */ diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 6a05e6185..6d73a1127 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -259,7 +259,7 @@ P5T5 Phase E (2026-04-26): per-fold reset for `MetricBandsRegistry` regression-d | `attention_backward_kernel.cu` | `gpu_attention.rs`, `gpu_tlob.rs` (reused for grad_norm+Adam kernels) | Wired | Attention backward pass | — | | `tlob_kernel.cu` (`tlob_sdp_forward`, `tlob_sdp_backward`, `tlob_write_states`, `tlob_read_grad_from_states`) | `gpu_tlob.rs` | Wired | D.8 Plan 2 Task 6C — TLOB SDP forward/backward, state scatter/read kernels | — | | `training_guard_kernel.cu` (`training_guard_check_and_accumulate`) | `gpu_training_guard.rs` | Wired | NaN / guard check kernel | — | -| `backtest_env_kernel.cu` (`backtest_env_step`) | `gpu_backtest_evaluator.rs` (`ENV_CUBIN`) | Wired | Backtest environment step. Diagnostic printf at `backtest_env_step_batch` sampled across full 214K val window for window 0 — tracking val-collapse Long/Short→Flat gate (first 10 bars confirmed kernel-correct; gate must activate beyond first chunk); removed once root cause identified. | — | +| `backtest_env_kernel.cu` (`backtest_env_step`) | `gpu_backtest_evaluator.rs` (`ENV_CUBIN`) | Wired | Backtest environment step | — | | `backtest_metrics_kernel.cu` | `gpu_backtest_evaluator.rs` (`METRICS_CUBIN`) | Wired | Backtest metrics reduction | — | | `backtest_plan_kernel.cu` (`backtest_plan_diag_reduce`, `backtest_plan_state_isv`) | `gpu_backtest_evaluator.rs` (`BACKTEST_PLAN_CUBIN`) | Wired | Plan-ISV diagnostic reduction in val. D.6 (Plan 2 Task 6A): `backtest_plan_state_isv` now writes `pisv[6] = remaining_fraction`; stride updated to `SL_PORTFOLIO_PLAN_DIM=7`. | — | | `backtest_forward_ppo_kernel.cu` (`backtest_forward_ppo_kernel`) | `gpu_backtest_evaluator.rs::evaluate_ppo` (`PPO_FORWARD_CUBIN`) | OUT-of-DQN-scope | PPO-path backtest only; correct-as-is | — |