diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index aabaf0115..f38c54a07 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -487,19 +487,32 @@ extern "C" __global__ void backtest_env_step_batch( &trail_triggered ); - /* DIAG (val-collapse): print first 10 bars of window 0. Captures the - * inputs to unified_env_step_core + its outputs so we can localise - * which step zeros target_position for non-Hold picks. Removed once - * the gate is identified. */ - if (w == 0 && current_step < 10) { - int dir_dbg = action_val / (b1_size * b2_size * b3_size); - int mag_dbg = (action_val / (b2_size * b3_size)) % b1_size; + /* 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\n", + "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); + conviction_core, health_k_batch, value, max_equity, cash); } (void)prev_position; (void)prev_position_sign; (void)trail_triggered; diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index cc0263331..a1599885a 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` for first 10 bars of window 0 — tracking val-collapse Long/Short→Flat gate; removed once root cause identified. | — | +| `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_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 | — |