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

fix(dqn): val_dir_dist + active_frac measurement artifact (-1 sentinel)

Browse Source 
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.
This commit is contained in:
jgrusewski
2026-04-27 10:36:09 +02:00
parent a3bb040bc2
commit a86fba2b1d
3 changed files with 64 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu
View File

@@ -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).

32
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -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.

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

@@ -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