From f9a8a5aa9a076eedbe6d6fe1140bc7c382dbe319 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 23 Apr 2026 00:22:56 +0200 Subject: [PATCH] feat(dqn): intent-side magnitude distribution diagnostic (EVAL_INTENT_MAG_DIST) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- .../src/cuda_pipeline/experience_kernels.cu | 69 +++++++++- .../cuda_pipeline/gpu_backtest_evaluator.rs | 126 +++++++++++++++++- .../cuda_pipeline/gpu_experience_collector.rs | 1 + .../dqn/smoke_tests/magnitude_distribution.rs | 11 ++ .../src/trainers/dqn/trainer/constructor.rs | 1 + crates/ml/src/trainers/dqn/trainer/metrics.rs | 10 ++ crates/ml/src/trainers/dqn/trainer/mod.rs | 21 +++ 7 files changed, 236 insertions(+), 3 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 9ef950e38..38c8e8467 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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]; +} diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 8548590f8..0e0c9614f 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -288,6 +288,14 @@ pub struct GpuBacktestEvaluator { done_buf: CudaSlice, // [n_windows] actions_buf: CudaSlice, // [n_windows] actions_history_buf: CudaSlice, // [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>, // Gather kernel output buffer (overwritten every step) states_buf: CudaSlice, // [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>, + /// 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>, + + /// `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, } 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::(cn) .map_err(|e| MLError::ModelError(format!("alloc chunked actions_buf: {e}")))?; + let ch_intent_mag_buf = self.stream.alloc_zeros::(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::(n * self.max_len) + .map_err(|e| MLError::ModelError(format!("alloc intent_mag_buf: {e}")))?; let ch_q_gaps = self.stream.alloc_zeros::(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(()) } diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 70d46eb54..c5286dd71 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -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}" diff --git a/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs b/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs index 5e234d242..085aa288b 100644 --- a/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs +++ b/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs @@ -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- diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs index 13158e001..571442737 100644 --- a/crates/ml/src/trainers/dqn/trainer/constructor.rs +++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs @@ -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], diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs index ca3c43d70..85d3104f0 100644 --- a/crates/ml/src/trainers/dqn/trainer/metrics.rs +++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs @@ -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) diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index 948b4a8dc..b839f4338 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/crates/ml/src/trainers/dqn/trainer/mod.rs @@ -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