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