From a86fba2b1df5884fdb1618e7b3593a9f2472cbf6 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 27 Apr 2026 10:36:09 +0200 Subject: [PATCH] fix(dqn): val_dir_dist + active_frac measurement artifact (-1 sentinel) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit After done_flags[w]=1 (capital floor breach), backtest_env_step early-returns without writing actions_history_buf for remaining slots in [done_step, max_len). The prior zero-init decoded those slots as Short Quarter Market Normal (action 0) via `dir = 0/27 = 0` and inflated val_dir_dist's Short bucket / active_frac to a measurement artifact masking real model behaviour. Two-part fix: 1. `gpu_backtest_evaluator.rs::reset_evaluation_state`: replace `memset_zeros` for actions_history_buf with `cuMemsetD32Async(0xFFFFFFFFu32)` writing -1 sentinel. The Rust readers already filter `if a < 0 { continue; }` so unwritten slots are skipped correctly post-fix. 2. `backtest_metrics_kernel.cu`: add `if (act < 0) continue;` after reading actions_history. The reduce-side metrics (buy_count/sell_count/hold_count → active_frac/dir_entropy + bnd_* trade-boundary detection) now consistently skip unwritten slots. step_returns at those slots are still zero-init (correct) so summing them with r=0 is a no-op. Empirical impact (local 3-fold × 5-epoch smoke, RTX 3050 Ti): val_dir_dist Short: 81-84% → 13-29% (matches val_picked within 5pp) active_frac: 87-91% → 31-50% dir_entropy: 0.57 → 0.83-1.02 The pre-fix "val-Flat-collapse" / "Short-collapse" pathology that motivated substantial subsequent investigation (incl. the 4-plan distributional-RL Thompson rollout draft) was largely a measurement artifact from this bug surfacing differently before vs after the Kelly cap fix (`0c9d1ee39`). Pre-Kelly the Kelly cap clamped most Long/Short → Flat → actions_history was densely written with Flat encoding → 80% Flat reading (real Kelly pathology + small artifact). Post-Kelly the picks survive but the poor smoke-trained model breaches capital floor often → many unwritten Short slots → 83% Short reading (pure artifact). With both fixes, val_dir_dist now reflects real model behaviour. Audit entry updated in docs/dqn-wire-up-audit.md. --- .../cuda_pipeline/backtest_metrics_kernel.cu | 11 +++++++ .../cuda_pipeline/gpu_backtest_evaluator.rs | 32 +++++++++++++++++-- docs/dqn-wire-up-audit.md | 24 ++++++++++++++ 3 files changed, 64 insertions(+), 3 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu index 5e7a5d354..c3cb6a928 100644 --- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -102,6 +102,17 @@ extern "C" __global__ void compute_backtest_metrics( local_cum *= (1.0f + r); int act = actions_history[base + i]; + /* Skip unwritten slots (act < 0): after `done_flags[w]=1` (capital + * floor breach), backtest_env_step early-returns without writing + * actions_history for the remaining slots in [done_step, max_len). + * `reset_evaluation_state` initialises those to -1 sentinel. Without + * this guard, C integer division (`-1 / 9 = 0`) decoded unwritten + * slots as Short Quarter (action 0), inflating sell_count and + * collapsing active_frac/dir_entropy to a measurement artifact + * rather than real model behaviour. step_returns at these slots are + * still zero (memset_zeros'd, never written), so summing them with + * `r=0` is a no-op and remains correct. */ + if (act < 0) continue; int exp_idx = act / (order_actions * urgency_actions); int raw_dir = exp_idx / 3; /* 0=Short, 1=Hold, 2=Long, 3=Flat */ /* Collapse Hold(1) and Flat(3) → no-exposure (-1). diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index f4c94c2d3..2490de21d 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -1258,15 +1258,41 @@ impl GpuBacktestEvaluator { .map_err(|e| MLError::ModelError(format!("reset portfolio upload: {e}")))?; self.portfolio_buf = portfolio_f32; - // Zero done flags, step returns, step rewards, actions history + // Zero done flags, step returns, step rewards self.stream.memset_zeros(&mut self.done_buf) .map_err(|e| MLError::ModelError(format!("reset done_buf: {e}")))?; self.stream.memset_zeros(&mut self.step_returns_buf) .map_err(|e| MLError::ModelError(format!("reset step_returns: {e}")))?; self.stream.memset_zeros(&mut self.step_rewards_buf) .map_err(|e| MLError::ModelError(format!("reset step_rewards: {e}")))?; - self.stream.memset_zeros(&mut self.actions_history_buf) - .map_err(|e| MLError::ModelError(format!("reset actions_history: {e}")))?; + + // Initialise actions_history_buf to -1 sentinel (NOT zero) so the + // reader (`read_eval_action_distribution_per_direction`) can distinguish + // unwritten slots from real action_idx=0 (Short Quarter Market Normal). + // + // Why this matters: after `done_flags[w]=1` (capital floor breach), + // backtest_env_step early-returns without writing actions_history for + // the remaining slots in [done_step, max_len). Zero-initialised + // unwritten slots decoded as Short via `dir = 0/27 = 0` and inflated + // val_dir_dist's Short bucket — a measurement artifact masking the + // real direction distribution. The reader already filters `if a < 0 + // { continue; }`, so writing -1 (i32) = 0xFFFFFFFF (u32) lets it skip + // unwritten slots correctly. + let ah_ptr = self.actions_history_buf.raw_ptr(); + let ah_len = self.actions_history_buf.len(); + unsafe { + let rc = cudarc::driver::sys::cuMemsetD32Async( + ah_ptr, + 0xFFFFFFFFu32, + ah_len, + self.stream.cu_stream(), + ); + if rc != cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS { + return Err(MLError::ModelError(format!( + "reset actions_history (cuMemsetD32Async -1): {rc:?}" + ))); + } + } // Zero the parallel intent-magnitude history (diagnostic) so a short // eval rollout can't read stale values from a prior longer rollout. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index a854b933c..8080acd30 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,30 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +actions_history_buf measurement artifact fix (2026-04-27): two-part fix. (1) +`gpu_backtest_evaluator.rs::reset_evaluation_state` now initialises +`actions_history_buf` to `-1` sentinel via `cuMemsetD32Async(0xFFFFFFFFu32)` +instead of `memset_zeros`. After `done_flags[w]=1` (capital floor breach), +`backtest_env_step` early-returns without writing actions_history for the +remaining slots in `[done_step, max_len)`; the prior zero-init decoded those +as Short Quarter Market Normal (action 0) via `dir = 0/27 = 0`. (2) +`backtest_metrics_kernel.cu` added `if (act < 0) continue;` after reading +`actions_history` so the reduce-side metrics (buy_count/sell_count/hold_count +→ active_frac/dir_entropy + bnd_* trade-boundary detection) skip unwritten +slots consistently with the existing Rust-side reader guards. Empirical impact +on the local 3-fold × 5-epoch smoke: val_dir_dist Short dropped from 83% +(artifact) to 13-29% (matches val_picked_dir_dist within 5pp); active_frac +dropped from 87-91% to 31-50%; dir_entropy increased from 0.57 to 0.83-1.02. +The pre-fix "val-Flat-collapse" / "Short-collapse" pathology that motivated +substantial subsequent investigation was largely a measurement artifact from +this bug surfacing differently before vs after the Kelly cap fix +(`0c9d1ee39`). Pre-Kelly fix, Kelly cap clamped Long/Short → Flat, so +actions_history was densely written with Flat encoding → fewer unwritten +slots → 80% Flat reading (real Kelly pathology + small artifact). Post-Kelly +fix, Long/Short survive but poor smoke-trained model breaches capital floor +often → more unwritten Short slots → 83% Short reading (pure artifact). With +both fixes, val_dir_dist reflects real model behaviour. + Plan A Phase 0 Thompson test kernel (2026-04-27): `cuda_pipeline/thompson_test_kernel.cu` — standalone test-only CUDA kernel implementing the Thompson direction sampling math (inverse-CDF over C51 atoms, uniform-τ over IQN quantiles, argmax of E[Q]) plus