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

feat(dqn): intent-side magnitude distribution diagnostic (EVAL_INTENT_MAG_DIST)

Browse Source 
Adds a parallel read path that reports the policy's intended mag_idx
BEFORE Kelly/margin caps and before the Hold/Flat dir_idx forces
mag=0. This exposes whether the magnitude Q-head is learning state-
dependent preferences, independent of the Kelly cold-start cap that
was masking it via actual_mag decoding (kelly_position_cap
warmup_floor=0.5 + safety=0.5+0.5*health pinning abs_pos <= 0.375).

Kernel changes (experience_kernels.cu):
- experience_action_select: new trailing optional arg `out_intent_mag`
  (int*, NULL=skip). Populated AFTER the existing mag_idx selection
  via a strict argmax over q_b1, ignoring the Hold/Flat mag=0 forcing,
  with the same higher-bin-wins tie-break used in the b2/b3 paths.
  Uses q_sign so the intent stays consistent with contrarian mode.
- New scatter_intent_chunk kernel: copies step-major chunked intent
  [chunk_len, n_windows] into window-major intent_history
  [n_windows, max_len], mirroring the actions_history layout.

Rust wiring (gpu_backtest_evaluator.rs):
- New fields intent_mag_buf, chunked_intent_mag_buf, scatter_intent_kernel.
  Buffers allocated alongside existing chunked buffers in
  ensure_action_select_ready. intent_mag_buf is zeroed by
  reset_evaluation_state so short rollouts don't read stale data.
- submit_dqn_step_loop_cublas appends the new arg to the action_select
  launch and launches scatter_intent_chunk immediately after, before
  the env_batch_kernel (which never touches intent_mag_buf).
- read_eval_intent_magnitude_distribution mirrors
  read_eval_action_distribution_per_magnitude but decodes raw mag_idx
  (a as usize) rather than the factored action encoding.

Training-path call site (gpu_experience_collector.rs): passes NULL
(0u64) for the new arg — training does not collect intent history.

Trainer wiring:
- new last_eval_intent_magnitude_dist field + accessor; populated in
  metrics.rs::evaluate_on_gpu next to last_eval_magnitude_dist.

Smoke test (magnitude_distribution.rs): adds [EVAL_INTENT_MAG_DIST]
println line; no new assertions.

Diagnostic-only, no new feature flag — production behaviour unchanged.
All 7 files build cleanly with no new warnings.

[testing: smoke compiles + runs, still fails on the existing H10
assertion (EVAL_DIST Quarter=1.000 driven by Kelly cold-start cap),
EVAL_INTENT_MAG_DIST shows Quarter=0.357 Half=0.045 Full=0.599 at
20-epoch smoke on local RTX 3050 Ti — confirming the magnitude head
prefers Full ~60% of the time while the Kelly-capped realised
distribution pins to Quarter.]

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-23 00:22:56 +02:00
parent b8cd4e1d94
commit f9a8a5aa9a
7 changed files with 236 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

69
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -760,7 +760,12 @@ extern "C" __global__ void experience_action_select(
int timestep, /* current timestep for stateless RNG */
const float* __restrict__ per_sample_epsilon, /* [N] IQL expectile gap epsilon, NULL=use cosine schedule */
const float* __restrict__ isv_signals_ptr, /* [8] pinned ISV signals for adaptive hold. NULL = static. */
int contrarian_active /* D7/N7: when non-zero, negate Q values before Boltzmann → argmin sampling */
int contrarian_active, /* D7/N7: when non-zero, negate Q values before Boltzmann → argmin sampling */
int* __restrict__ out_intent_mag /* [N] optional: pure-policy magnitude Q-head argmax BEFORE
* Hold/Flat forcing and BEFORE Kelly/margin caps applied in
* env_step. NULL = skip. Diagnostic-only — lets operators see
* what the magnitude head WOULD pick if direction gating and
* post-policy physics didn't mask it. */
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -1095,6 +1100,26 @@ extern "C" __global__ void experience_action_select(
}
}
/* Pure-policy magnitude intent — diagnostic only, bypasses the Hold/Flat
* dir_idx gate that forces mag_idx=0 and the downstream Kelly/margin
* clamps inside env_step. Computed from q_b1 with strict argmax and the
* same tie-break convention used elsewhere (prefer higher bin index on
* ties within 1e-6). Uses q_sign to stay consistent with contrarian
* mode. Written to out_intent_mag[i] if non-NULL. */
if (out_intent_mag != NULL) {
int intent = 0;
float q_intent_max = q_sign * q_b1[0];
for (int a = 1; a < b1_size; a++) {
float qv_intent = q_sign * q_b1[a];
if (qv_intent > q_intent_max + 1e-6f
|| (fabsf(qv_intent - q_intent_max) <= 1e-6f && a > intent)) {
q_intent_max = qv_intent;
intent = a;
}
}
out_intent_mag[i] = intent;
}
/* Compose factored action: 4-branch encoding
* action = dir * (b1*b2*b3) + mag * (b2*b3) + order * b3 + urgency */
int action_idx = dir_idx * b1_size * b2_size * b3_size
@@ -6185,3 +6210,45 @@ extern "C" __global__ void backtest_state_gather(
for (int k = SL_STATE_DIM; k < padded_sd; k++)
out[k] = 0.0f;
}
/* ================================================================== */
/* Kernel: scatter_intent_chunk */
/* ================================================================== */
/**
* Scatter per-window, per-step intent magnitude values from the chunked
* action-select output buffer into a window-major history buffer matching
* the layout of actions_history_buf.
*
* Input layout (chunk_intent): [chunk_len, n_windows] — step-major, same
* as ch_actions (written by experience_action_select
* with batch = n_windows * chunk_len and thread
* index i = s * n_windows + w).
* Output layout (intent_history): [n_windows, max_len] — window-major, mirrors
* actions_history[w * max_len + current_step].
*
* Each thread handles one (window, step_offset) pair. Thread index
* t = s * n_windows + w with w in [0, n), s in [0, chunk_len).
*
* Diagnostic-only kernel. Writes within [start_step, start_step + chunk_len);
* all other positions in intent_history remain zero from reset_evaluation_state.
*/
extern "C" __global__ void scatter_intent_chunk(
const int* __restrict__ chunk_intent, /* [chunk_len, n_windows] step-major */
int* __restrict__ intent_history, /* [n_windows, max_len] window-major */
int n_windows,
int max_len,
int start_step,
int chunk_len
) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int total = n_windows * chunk_len;
if (tid >= total) return;
int s = tid / n_windows; /* step offset within chunk */
int w = tid - s * n_windows;
int step = start_step + s;
if (step < 0 || step >= max_len) return;
intent_history[w * max_len + step] = chunk_intent[s * n_windows + w];
}

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

@@ -288,6 +288,14 @@ pub struct GpuBacktestEvaluator {
done_buf: CudaSlice<i32>, // [n_windows]
actions_buf: CudaSlice<i32>, // [n_windows]
actions_history_buf: CudaSlice<i32>, // [n_windows * max_len]
/// Intent magnitude history — pure-policy magnitude argmax from each
/// eval step, recorded BEFORE the Hold/Flat dir_idx forcing in the
/// action-select kernel and BEFORE Kelly/margin caps modify exposure
/// inside env_step. Parallel diagnostic layout to `actions_history_buf`:
/// `[n_windows * max_len]`, indexed `w * max_len + step`. Populated
/// only when `evaluate_dqn_graphed` runs — stays zero otherwise.
/// `None` until `ensure_action_select_ready` allocates it.
intent_mag_buf: Option<CudaSlice<i32>>,
// Gather kernel output buffer (overwritten every step)
states_buf: CudaSlice<f32>, // [n_windows * state_dim] (8-aligned, f32 for cublasSgemm)
@@ -350,6 +358,15 @@ pub struct GpuBacktestEvaluator {
/// Batched forward actions buffer for chunked action select: [n_windows * CHUNK_SIZE].
chunked_actions_buf: Option<CudaSlice<i32>>,
/// Chunked intent-magnitude buffer mirroring `chunked_actions_buf`:
/// `[chunk_len, n_windows]` step-major layout (thread `i = s*n + w`).
/// Filled by `experience_action_select` on each chunk, then scattered
/// into the window-major `intent_mag_buf` by `scatter_intent_kernel`.
chunked_intent_mag_buf: Option<CudaSlice<i32>>,
/// `scatter_intent_chunk` kernel handle — copies `chunked_intent_mag_buf`
/// into `intent_mag_buf` at the current chunk's step range.
scatter_intent_kernel: Option<CudaFunction>,
}
impl Drop for GpuBacktestEvaluator {
@@ -566,6 +583,7 @@ impl GpuBacktestEvaluator {
done_buf,
actions_buf,
actions_history_buf,
intent_mag_buf: None,
states_buf,
metrics_buf,
n_windows,
@@ -588,6 +606,8 @@ impl GpuBacktestEvaluator {
chunked_q_values: None,
chunked_states_buf: None,
chunked_actions_buf: None,
chunked_intent_mag_buf: None,
scatter_intent_kernel: None,
})
}
@@ -894,6 +914,51 @@ impl GpuBacktestEvaluator {
])
}
/// Pure-policy magnitude intent distribution from the most recent
/// eval rollout. Parallel to `read_eval_action_distribution_per_magnitude`
/// but reads `intent_mag_buf` — raw `mag_idx` values (0=Quarter,
/// 1=Half, 2=Full) written by the action-select kernel BEFORE the
/// Hold/Flat dir_idx forcing and BEFORE Kelly/margin caps inside
/// env_step decode `actual_mag`. This answers: "what magnitude does
/// the policy's Q-head actually prefer when Kelly cold-start is masking
/// everything to Quarter via `actual_mag`?"
///
/// Returns `Ok([0.0; 3])` if `intent_mag_buf` has not been allocated
/// yet (evaluator hasn't run a DQN eval), or if all entries are
/// skipped (zero rollout).
pub fn read_eval_intent_magnitude_distribution(&self) -> Result<[f32; 3], MLError> {
let buf = match self.intent_mag_buf.as_ref() {
Some(b) => b,
None => return Ok([0.0; 3]),
};
let len = self.n_windows * self.max_len;
let mut host = vec![0_i32; len];
self.stream
.memcpy_dtoh(buf, &mut host)
.map_err(|e| MLError::ModelError(format!("intent_mag_buf dtoh: {e}")))?;
let mut counts = [0_u64; 3];
let mut total = 0_u64;
for &a in &host {
if a < 0 {
continue;
}
let mag = a as usize;
if mag < 3 {
counts[mag] += 1;
total += 1;
}
}
if total == 0 {
return Ok([0.0; 3]);
}
let t = total as f32;
Ok([
counts[0] as f32 / t,
counts[1] as f32 / t,
counts[2] as f32 / t,
])
}
/// Task 2.X "make Full useful" diagnostic — read back the per-direction
/// action distribution from the most recent evaluation. Action encoding
/// is `dir*27 + mag*9 + ord*3 + urg`, so direction = `action / 27`
@@ -957,6 +1022,13 @@ impl GpuBacktestEvaluator {
self.stream.memset_zeros(&mut self.actions_history_buf)
.map_err(|e| MLError::ModelError(format!("reset actions_history: {e}")))?;
// Zero the parallel intent-magnitude history (diagnostic) so a short
// eval rollout can't read stale values from a prior longer rollout.
if let Some(ref mut intent) = self.intent_mag_buf {
self.stream.memset_zeros(intent)
.map_err(|e| MLError::ModelError(format!("reset intent_mag_buf: {e}")))?;
}
// Reset Kelly stats — each evaluation starts cold. Kelly priors +
// warmup floor in trade_physics.cuh handle the cold-start period.
self.stream.memset_zeros(&mut self.kelly_stats_buf)
@@ -1002,6 +1074,15 @@ impl GpuBacktestEvaluator {
let ch_q_gaps = self.chunked_q_gaps_buf.as_ref().ok_or_else(|| {
MLError::ModelError("chunked_q_gaps_buf unexpectedly None".to_owned())
})?;
let ch_intent = self.chunked_intent_mag_buf.as_ref().ok_or_else(|| {
MLError::ModelError("chunked_intent_mag_buf unexpectedly None".to_owned())
})?;
let intent_history = self.intent_mag_buf.as_ref().ok_or_else(|| {
MLError::ModelError("intent_mag_buf unexpectedly None".to_owned())
})?;
let scatter_intent = self.scatter_intent_kernel.as_ref().ok_or_else(|| {
MLError::ModelError("scatter_intent_kernel unexpectedly None".to_owned())
})?;
let n = self.n_windows;
let state_dim = self.state_dim;
@@ -1107,20 +1188,49 @@ impl GpuBacktestEvaluator {
.arg(&0u64) // per_sample_epsilon: NULL → use cosine schedule
.arg(&self.isv_signals_ptr) // ISV signals for adaptive hold
.arg(&0i32) // F6: contrarian_active — always off during backtest
.arg(ch_intent) // out_intent_mag: pure-policy magnitude intent diagnostic
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"chunked action_select chunk_start={chunk_start}: {e}"
)))?;
}
// ── Phase 4b: Scatter chunked intent into window-major history
//
// experience_action_select writes ch_intent[s * n + w] for batch
// index i = s * n + w. scatter_intent_chunk copies into the
// window-major intent_mag_buf at [w * max_len + (start + s)].
// Diagnostic-only — writes are additive with the Phase 5 env
// kernel (which never touches intent_mag_buf).
let start_step_i32 = chunk_start as i32;
let chunk_len_i32 = chunk_len as i32;
let scatter_threads = n * chunk_len;
let scatter_blocks = ((scatter_threads as u32) + 255) / 256;
unsafe {
self.stream
.launch_builder(scatter_intent)
.arg(ch_intent)
.arg(intent_history)
.arg(&(n as i32))
.arg(&(self.max_len as i32))
.arg(&start_step_i32)
.arg(&chunk_len_i32)
.launch(LaunchConfig {
grid_dim: (scatter_blocks.max(1), 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"scatter_intent_chunk chunk_start={chunk_start}: {e}"
)))?;
}
// ── Phase 5: Batched env_step (single launch for entire chunk) ──
//
// All chunk_len steps processed in one kernel launch. The batched
// kernel loops over steps internally, reading actions directly from
// the chunked_actions buffer. Eliminates chunk_len individual launches
// (was 512 launches → 1 launch per chunk).
let start_step_i32 = chunk_start as i32;
let chunk_len_i32 = chunk_len as i32;
let env_blocks = ((n as u32) + 255) / 256;
let feature_dim_i32 = self.feature_dim as i32;
unsafe {
@@ -1359,6 +1469,9 @@ impl GpuBacktestEvaluator {
let as_kernel = dqn_module
.load_function("experience_action_select")
.map_err(|e| MLError::ModelError(format!("experience_action_select load: {e}")))?;
let scatter_intent = dqn_module
.load_function("scatter_intent_chunk")
.map_err(|e| MLError::ModelError(format!("scatter_intent_chunk load: {e}")))?;
// ── Chunked buffers (batch_size = n_windows * CHUNK_SIZE) ────────────
let chunk = DQN_BACKTEST_CHUNK_SIZE;
@@ -1372,6 +1485,12 @@ impl GpuBacktestEvaluator {
.map_err(|e| MLError::ModelError(format!("alloc chunked states_buf: {e}")))?;
let ch_actions_buf = self.stream.alloc_zeros::<i32>(cn)
.map_err(|e| MLError::ModelError(format!("alloc chunked actions_buf: {e}")))?;
let ch_intent_mag_buf = self.stream.alloc_zeros::<i32>(cn)
.map_err(|e| MLError::ModelError(format!("alloc chunked intent_mag_buf: {e}")))?;
// Intent-magnitude history: parallel layout to actions_history_buf.
let intent_mag_buf = self.stream.alloc_zeros::<i32>(n * self.max_len)
.map_err(|e| MLError::ModelError(format!("alloc intent_mag_buf: {e}")))?;
let ch_q_gaps = self.stream.alloc_zeros::<f32>(cn)
.map_err(|e| MLError::ModelError(format!("alloc chunked q_gaps: {e}")))?;
@@ -1391,7 +1510,10 @@ impl GpuBacktestEvaluator {
self.chunked_q_values = Some(ch_q_values);
self.chunked_states_buf = Some(ch_states_buf);
self.chunked_actions_buf = Some(ch_actions_buf);
self.chunked_intent_mag_buf = Some(ch_intent_mag_buf);
self.chunked_q_gaps_buf = Some(ch_q_gaps);
self.intent_mag_buf = Some(intent_mag_buf);
self.scatter_intent_kernel = Some(scatter_intent);
Ok(())
}

1
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -2867,6 +2867,7 @@ impl GpuExperienceCollector {
.arg(&self.per_sample_epsilon_ptr) // IQL expectile gap epsilon (0=cosine schedule)
.arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0=NULL=static)
.arg(&(self.contrarian_active_cache as i32)) // D7/N7: Q-negation flag
.arg(&0u64) // out_intent_mag: NULL — training path does not collect
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_action_select t={t}: {e}"

11
crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs
View File

@@ -66,6 +66,17 @@ fn test_magnitude_distribution() -> Result<()> {
let [eq, eh, ef] = trainer.last_eval_magnitude_dist();
println!("[EVAL_DIST] Quarter={:.3} Half={:.3} Full={:.3}", eq, eh, ef);
// Diagnostic — pure-policy magnitude intent BEFORE Kelly/margin caps and
// BEFORE the Hold/Flat dir gate force mag=0. When EVAL_DIST collapses to
// Quarter=1 because kelly_position_cap's warmup_floor pins abs_pos ≤
// 0.375×max_pos for the first ~10 trades, this line reveals whether the
// magnitude Q-head actually learned state-dependent preferences.
let [iq, ih, if_] = trainer.last_eval_intent_magnitude_dist();
println!(
"[EVAL_INTENT_MAG_DIST] Quarter={:.3} Half={:.3} Full={:.3}",
iq, ih, if_
);
// Task 2.X diagnostic — eval direction distribution. When eval collapses
// to Quarter==1 AND Hold+Flat dominate, the root cause is direction-side
// (magnitude gets forced to 0 by kernel for Hold/Flat), so a magnitude-

1
crates/ml/src/trainers/dqn/trainer/constructor.rs
View File

@@ -583,6 +583,7 @@ impl DQNTrainer {
controller_total_epochs: 0,
grad_clip_kicked_this_epoch: false,
last_eval_magnitude_dist: [0.0_f32; 3],
last_eval_intent_magnitude_dist: [0.0_f32; 3],
prev_reward_contrib_popart_var: 0.0,
last_reward_contrib: [0.0_f32; 4],

10
crates/ml/src/trainers/dqn/trainer/metrics.rs
View File

@@ -643,6 +643,16 @@ impl DQNTrainer {
Ok(dist) => self.last_eval_magnitude_dist = dist,
Err(e) => tracing::warn!("eval magnitude dist readback failed (non-fatal): {e}"),
}
// Parallel pure-policy intent readout — sees what the magnitude
// Q-head WOULD pick before Hold/Flat gating and before Kelly/margin
// caps collapse the realised `actual_mag` to Quarter during the
// cold-start warmup. Diagnostic-only; failure is non-fatal.
match ev.read_eval_intent_magnitude_distribution() {
Ok(dist) => self.last_eval_intent_magnitude_dist = dist,
Err(e) => tracing::warn!(
"eval intent magnitude dist readback failed (non-fatal): {e}"
),
}
}
Ok(-val_sharpe)

21
crates/ml/src/trainers/dqn/trainer/mod.rs
View File

@@ -292,6 +292,17 @@ pub struct DQNTrainer {
/// Layout: [Quarter, Half, Full] in [0, 1].
pub(crate) last_eval_magnitude_dist: [f32; 3],
/// Pure-policy magnitude-intent distribution from the most recent
/// VALIDATION backtest (eval-mode rollout), read from the parallel
/// `intent_mag_buf` populated by the action-select kernel BEFORE the
/// Hold/Flat dir_idx forcing and BEFORE Kelly/margin caps mask the
/// realised `actual_mag` inside env_step. Layout: [Quarter, Half, Full]
/// in [0, 1]. Diagnostic-only — used by the magnitude_distribution
/// smoke test to surface what the magnitude head WOULD pick independent
/// of the Kelly cold-start cap (kelly_position_cap warmup_floor=0.5 +
/// safety=0.5+0.5×health pins abs_pos ≤ 0.375×max_pos).
pub(crate) last_eval_intent_magnitude_dist: [f32; 3],
/// Task 0.9 — prior-epoch values of adaptive-controller outputs, used to
/// detect "fire" = value changed from the prior epoch. Controllers audited:
/// anti-LR, adaptive tau, adaptive gamma, adaptive grad-clip, CQL alpha
@@ -1524,6 +1535,16 @@ impl DQNTrainer {
self.last_eval_magnitude_dist
}
/// Pure-policy magnitude-intent distribution from the most recent
/// VALIDATION backtest — what the magnitude Q-head WOULD pick if
/// direction gating (Hold/Flat → mag=0) and Kelly/margin caps didn't
/// mask the realised `actual_mag`. Layout: [Quarter, Half, Full] in
/// [0, 1]. Diagnostic-only — complements `last_eval_magnitude_dist`
/// when Kelly cold-start pins realised magnitude to Quarter.
pub fn last_eval_intent_magnitude_dist(&self) -> [f32; 3] {
self.last_eval_intent_magnitude_dist
}
/// Task 2.X diagnostic — per-direction eval action distribution, readback
/// from the most recent validation backtest. Layout: [Short, Hold, Long,
/// Flat] in [0, 1]. When magnitude_distribution's `ef` collapses to zero