gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

diag(dqn): kernel printf for first 10 bars of val window 0

Browse Source 
Captures action_val + decoded dir/mag, max_position, prev_position,
position post-trade, actual_dir, actual_mag, trail_triggered,
conviction, health, value, max_equity inside backtest_env_step_batch.
Localises which step in unified_env_step_core zeros target_position
(or position) for non-Hold picks.

After tracing the math, every analytical candidate (margin cap, Kelly
cap with health-coupled warmup floor, trail stop) returns a positive
target at cold start with my fix. val_dir_dist still locks at 0% Long
while val_picked_dir_dist shows 17% Long — only direct kernel-side
observation can show what's actually clamping.

Removed once the gate is identified.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-26 21:12:07 +02:00
parent e15adecef1
commit 7d29c5ee48
2 changed files with 16 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
View File

@@ -486,6 +486,21 @@ extern "C" __global__ void backtest_env_step_batch(
&actual_mag,
&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;
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",
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);
}
(void)prev_position; (void)prev_position_sign; (void)trail_triggered;
/* ── Post-trade floor check: catch intra-step breaches ────────── */

2
docs/dqn-wire-up-audit.md
View File

@@ -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 | — |
| `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_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 | — |