From fa8d5466140801bd78049ac18e2d48ff1d3dc30c Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 22 Apr 2026 20:12:20 +0200 Subject: [PATCH] =?UTF-8?q?diag+fix(dqn):=20Task=202.X=20ISV-adaptive=20ma?= =?UTF-8?q?gnitude=20mechanism=20=E2=80=94=20reveals=20direction-branch=20?= =?UTF-8?q?is=20the=20real=20blocker?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Per feedback_adaptive_not_tuned.md: adaptive signal-driven mechanism, zero static tuning knobs. Extends the existing ISV bus with per-magnitude Q-mean EMAs and an absolute-scale reference; C51 loss + gradient kernels now read ISV at zero hot-path cost to modulate per-bin weight in response to observed collapse severity. Weight is 1.0 when Q is healthy; scales up per-bin when collapse signal fires; self-disables as training stabilises. Additions: * ISV_DIM 13 → 17. New slots: [13] Q_MAG_MEAN_QUARTER: ema(mean Q(Quarter), tau=0.05) [14] Q_MAG_MEAN_HALF: ema(mean Q(Half), tau=0.05) [15] Q_MAG_MEAN_FULL: ema(mean Q(Full), tau=0.05) [16] Q_ABS_REF: ema(max(|Q_mean[k]|), tau=0.05) — scale-invariant reference * q_mag_bin_means_reduce kernel (q_stats_kernel.cu) — one-block reduce computing per-mag Q-means from q_out_buf; output written to pinned scratch slots; drives the EMAs in isv_signal_update. * c51_loss_kernel::get_magnitude_bin_weight helper + matching inlined logic in c51_grad_kernel: composite collapse signal = min(1, frac_bin + (1 - learning_health)); bin_weight = 1.0 + collapse * mag_bias_signal[k] (bounded in [1, 2]); mag_bias_signal[k] = (k+1)/b1_size. * isv_signal_update extended with q_mag_means_ptr + q_abs_ref_ptr + mag_size kernel args. Diagnostics (keystone finding below): * gpu_backtest_evaluator::read_eval_action_distribution_per_direction — 4-bin per-direction count at eval (Short/Hold/Long/Flat fractions). This diagnostic flipped the task diagnosis. * DQNTrainer::last_eval_direction_dist accessor. * last_isv_magnitude_bin_q_means accessor. * EVAL_DIR_DIST + ISV_BIN_MEANS debug prints in magnitude_distribution smoke. * ef >= 0.05 smoke gate added (currently unreachable behind pre-existing eh+ef >= 0.30 gate; kept for future use). Training-time outcome: Pre-fix MAG_DIST: Quarter=0.60 Half=0.10 Full=0.23 Post-fix MAG_DIST: Quarter=0.46 Half=0.24 Full=0.28 (2.4× Half lift, Full unchanged) Pre-fix EVAL_DIST: eq=1.000 eh=0.000 ef=0.000 Post-fix EVAL_DIST: eq=0.981 eh=0.019 ef=0.000 Root cause revealed (why the adaptive fix couldn't lift ef off 0): EVAL_DIR_DIST: Short=0.045 Hold=0.115 Long=0.070 Flat=0.771 ~88% of eval states have direction ∈ {Hold, Flat}. Kernel at experience_kernels.cu:~896 FORCES mag_idx=0 (Quarter) in those cases as a structural ABI invariant. Only ~11.5% of eval samples have a free magnitude choice. Upper bound on ef regardless of magnitude mechanism: ~0.11. The magnitude branch mechanism works as designed — it correctly rebalances per-bin Q-means and lifts the training-time Half share 2.4×. But direction-branch collapse to Flat masks everything downstream. Task 2.X's magnitude-only scope cannot unblock eval ef. The real fix target is direction-branch eval collapse. Follow-up task "Task 2.Y make direction-branch trade" extends the same ISV-driven composite-signal mechanism to branch 0 (Short/Long vs Hold/Flat). Scoping doc to be written. Smoke validation: magnitude_distribution FAIL (pre-existing eh+ef >= 0.30 gate; same fail mode as HEAD before this commit) reward_component_audit PASS controller_activity PASS exploration_coverage PASS multi_fold_convergence PASS (avg best_val_metric=0.039, within ±15%) No config fields. No static tuning knobs. No feature flags. All modulation flows through the ISV bus. Shape constants documented: eps=1e-6 (numerical guard), alpha=0.05 (ema tau matching existing ISV pattern), MAX_MAG=4 (branch-size ceiling, already established), mag_bias_signal[k]=(k+1)/b1_size (architectural monotonicity w.r.t. bin index as stake size). --- .../ml/src/cuda_pipeline/c51_grad_kernel.cu | 49 +++- .../ml/src/cuda_pipeline/c51_loss_kernel.cu | 113 +++++++++ .../src/cuda_pipeline/experience_kernels.cu | 50 +++- .../cuda_pipeline/gpu_backtest_evaluator.rs | 42 ++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 226 +++++++++++++++++- crates/ml/src/cuda_pipeline/q_stats_kernel.cu | 94 ++++++++ crates/ml/src/trainers/dqn/fused_training.rs | 10 + .../dqn/smoke_tests/magnitude_distribution.rs | 56 +++++ crates/ml/src/trainers/dqn/trainer/mod.rs | 22 ++ 9 files changed, 652 insertions(+), 10 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu index 6a3f4dacd..d375b16b6 100644 --- a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu @@ -27,7 +27,12 @@ extern "C" __global__ void c51_grad_kernel( const float* __restrict__ per_sample_support, /* [B, 3] per-sample [v_min, v_max, delta_z] */ const float* __restrict__ liquid_mod, /* [4] pinned device-mapped per-branch modulators */ const float* __restrict__ atom_positions, /* [4, num_atoms] adaptive positions. NULL = linear. */ - const float* __restrict__ q_mean_ema_ptr) /* [1] pinned device-mapped Q-mean EMA */ + const float* __restrict__ q_mean_ema_ptr, /* [1] pinned device-mapped Q-mean EMA */ + /* Task 2.X adaptive magnitude fix — ISV signal bus [ISV_DIM=15]. Slots + * [13]=q_mag_spread_ema, [14]=q_dir_spread_ema drive an adaptive bin + * weight applied to the magnitude branch (d==1). NULL → identity + * (safe: reduces to pre-fix behaviour). */ + const float* __restrict__ isv_signals) { int tid = blockIdx.x * blockDim.x + threadIdx.x; int total_elems = batch_size * num_atoms; @@ -67,6 +72,48 @@ extern "C" __global__ void c51_grad_kernel( d_combined += inv_batch * entropy_coeff * (1.0f + lp_clamped); } + /* Task 2.X "make Full useful": scale magnitude-branch gradient by + * the same signal-driven bin weight applied to branch_ce in + * c51_loss_batched. Consistency is required — the forward CE and + * the backward gradient must share the same scalar multiplier so + * the update step follows the true gradient of the (scaled) loss. + * + * Composite signal — see c51_loss_kernel.cu::get_magnitude_bin_weight + * for the full rationale. Component A reads ISV slots [13..16] + * (per-mag Q-mean EMAs + |Q| ref) and fires when the taken bin's + * Q-mean lags max. Component B reads ISV slot [12] (learning_health) + * and fires when training has not stabilised. collapse_frac = + * min(1, frac_bin + frac_health). Bounded bin_weight ∈ [1, 2]. */ + if (d == 1 && isv_signals != NULL) { + int a1 = branch_actions[1]; + if (b1_size > 0 && a1 >= 0 && a1 < b1_size && a1 < 3) { + /* Component A: bin-specific bias-vs-argmax */ + float q_q = isv_signals[13]; + float q_h = isv_signals[14]; + float q_f = isv_signals[15]; + float q_abs_ref = isv_signals[16]; + float means[3] = { q_q, q_h, q_f }; + float max_mean = means[0]; + for (int k = 1; k < 3; k++) + max_mean = fmaxf(max_mean, means[k]); + float frac_bin = 0.0f; + if (q_abs_ref > 1e-6f) { + float bias_gap = fmaxf(0.0f, max_mean - means[a1]); + frac_bin = fminf(1.0f, bias_gap / fmaxf(q_abs_ref, 1e-6f)); + } + /* Component B: meta-health stress signal */ + float health = fminf(1.0f, fmaxf(0.0f, isv_signals[12])); + float frac_health = 1.0f - health; + /* Composite */ + float collapse_frac = fminf(1.0f, frac_bin + frac_health); + if (collapse_frac > 0.0f) { + float mag_bias_signal = (float)(a1 + 1) / (float)b1_size; + float bin_weight = 1.0f + collapse_frac * mag_bias_signal; + d_combined *= bin_weight; + } + } + } + d_val_sum += d_combined; float branch_scale = branch_scales[b * 4 + d]; diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index 05e1fc1a2..c53a2dcd4 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -199,6 +199,107 @@ __device__ __forceinline__ float get_learning_health(const float* __restrict__ i return (isv_signals_ptr != NULL) ? isv_signals_ptr[12] : 0.5f; } +/* Task 2.X "make Full useful" — signal-driven magnitude bin weight. + * + * Composite collapse signal — combines two ISV-driven components so the + * mechanism has actual leverage in ALL smoke regimes observed during + * Task-2.X iteration (earlier formulations dormant because the chosen + * reference always matched or exceeded the target): + * + * ── Component A: bin-specific bias-vs-argmax (from per-bin Q-means) ── + * ISV slots [13..15] = per-magnitude-bin Q-value mean EMAs + * [13] Q_mean_ema(Quarter), [14] Q_mean_ema(Half), [15] Q_mean_ema(Full) + * ISV slot [16] = absolute-Q-scale reference EMA (max |Q_mean| across bins). + * + * max_mean = max(Q_mean[k]) + * bias_gap_bin = max(0, max_mean − Q_mean[a1]) + * frac_bin = min(1, bias_gap_bin / max(Q_abs_ref, eps)) + * + * `frac_bin` is positive iff the TAKEN bin's Q-mean is below the best + * bin's Q-mean (self-disables when the taken bin IS the argmax). + * + * ── Component B: meta-health stress signal (from ISV learning health) ── + * ISV slot [12] = learning_health ∈ [0, 1]. 1 = healthy, 0 = collapsed. + * + * frac_health = (1 − learning_health) ∈ [0, 1] + * + * This engages UNIFORMLY across all magnitude samples (not only the + * ones where Q(a1) lags max), providing leverage even when the per-bin + * component A is zero for that specific sample. Meta-signal-driven — + * fades to zero as the model's training dynamics stabilise. Same slot + * that already drives label-smoothing and Boltzmann temperature scaling + * in this kernel (see LABEL_SMOOTHING_HEALTH_INDEX); reusing it keeps + * the ISV surface coherent. + * + * ── Composite ── + * collapse_frac = min(1, frac_bin + frac_health) + * bin_weight = 1 + collapse_frac * mag_bias_signal(a1) ∈ [1, 2] + * + * mag_bias_signal[a1] = (a1 + 1) / b1_size ∈ {0.33, 0.67, 1.00} for + * Quarter/Half/Full respectively. Higher-index bins correspond to + * higher intrinsic PnL variance — exactly the dimension C51's + * expected-Q under-prices. Bin-proportional weighting concentrates + * the correction on the under-represented tail. + * + * Properties: + * * Bounded ∈ [1, 2] — keeps magnitude-branch gradient in the same + * ballpark as other loss terms' budgets (direction, order, urgency, + * CQL, IQN, ensemble). The fix is a correction, not a replacement. + * * Dual self-disabling: + * - Component A: when Full becomes the argmax bin, frac_bin(Full) + * drops to 0; the mechanism fades as the Q-landscape rebalances. + * - Component B: when learning_health → 1 (training stabilises), + * frac_health → 0; the mechanism fades as training converges. + * * No hard-coded tuning target — every term is an ISV-driven signal + * or an architectural property (bin count, bin index). + * + * Earlier Task-2.X iterations (dir-branch Q-spread reference; per-bin + * Q-mean bias alone) had smoke-observed collapse_frac=0 throughout + * training, so the mechanism never fired. The composite signal above + * uses an AND/OR ensemble of two independent collapse indicators — + * either fires when its specific collapse mode is active, so the + * mechanism reliably engages in practice. + * + * Returns 1.0 if ISV is unavailable (safe identity). `eps=1e-6f` is the + * standard numerical-guard epsilon matching existing patterns in this + * file (see block_bellman_project_f, barrier_gradient_direction, etc.). + */ +__device__ __forceinline__ float get_magnitude_bin_weight( + const float* __restrict__ isv_signals_ptr, + int a1, + int b1_size +) { + if (isv_signals_ptr == NULL) return 1.0f; + if (b1_size <= 0 || a1 < 0 || a1 >= b1_size || a1 >= 3) return 1.0f; + + /* Component A: bin-specific bias-vs-argmax */ + float q_q = isv_signals_ptr[13]; + float q_h = isv_signals_ptr[14]; + float q_f = isv_signals_ptr[15]; + float q_abs_ref = isv_signals_ptr[16]; + float mean_by_bin[3] = { q_q, q_h, q_f }; + float max_mean = mean_by_bin[0]; + for (int k = 1; k < 3; k++) + max_mean = fmaxf(max_mean, mean_by_bin[k]); + float frac_bin = 0.0f; + if (q_abs_ref > 1e-6f) { + float bias_gap = fmaxf(0.0f, max_mean - mean_by_bin[a1]); + frac_bin = fminf(1.0f, bias_gap / fmaxf(q_abs_ref, 1e-6f)); + } + + /* Component B: meta-health stress signal */ + float health = fminf(1.0f, fmaxf(0.0f, isv_signals_ptr[12])); + float frac_health = 1.0f - health; + + /* Composite — bounded ∈ [0, 1]. Reaches 1 when either component + * saturates OR both components are partially active and sum ≥ 1. */ + float collapse_frac = fminf(1.0f, frac_bin + frac_health); + if (collapse_frac <= 0.0f) return 1.0f; + + float mag_bias_signal = (float)(a1 + 1) / (float)b1_size; + return 1.0f + collapse_frac * mag_bias_signal; +} + /* ══════════════════════════════════════════════════════════════════════ * MAIN KERNEL: c51_loss_batched (float arithmetic, BF16 I/O) * ══════════════════════════════════════════════════════════════════════ */ @@ -713,6 +814,18 @@ extern "C" __global__ void c51_loss_batched( for (int j = tid; j < num_atoms; j += BLOCK_THREADS) local_ce -= shmem_lp[j] * shmem_current_lp[j]; float branch_ce = block_reduce_sum_f(local_ce, shmem_reduce, tid); + + /* Task 2.X adaptive magnitude fix: scale magnitude-branch CE by a + * signal-driven bin weight derived from the ISV bus. Identity + * (bin_weight=1.0) when ISV reports magnitude Q-spread ≥ direction + * Q-spread — the mechanism self-disables once the magnitude head's + * Q-values differentiate. See get_magnitude_bin_weight. */ + if (d == 1) { + float mag_bin_weight = get_magnitude_bin_weight( + isv_signals, branch_action[1], branch_sizes[1]); + branch_ce *= mag_bin_weight; + } + total_ce += branch_ce; /* ═══ STEP f: Online-target Q-divergence for adaptive tau ══ diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index fef789d82..2740efe79 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -5536,8 +5536,8 @@ extern "C" __global__ void risk_budget_backward( /* ================================================================== */ extern "C" __global__ void isv_signal_update( - float* __restrict__ isv_signals, /* [12] pinned device-mapped */ - float* __restrict__ isv_history, /* [K*12] pinned device-mapped */ + float* __restrict__ isv_signals, /* [ISV_DIM=17] pinned device-mapped */ + float* __restrict__ isv_history, /* [K*12] pinned — history only rotates slots [0..11] */ float* __restrict__ lagged_td_error, /* [1] pinned — for recursive confidence target */ const float* __restrict__ q_mean_ptr, /* [1] pinned — batch Q-mean */ const float* __restrict__ q_mean_ema_ptr, /* [1] pinned — Q-mean EMA */ @@ -5549,7 +5549,13 @@ extern "C" __global__ void isv_signal_update( const float* __restrict__ loss_ptr, /* [1] pinned — total loss */ int K, /* temporal history length (4) */ const float* __restrict__ states_ptr, /* [B, state_dim] — for ADX/CUSUM regime signals. NULL = skip regime. */ - int state_dim /* to index ADX at [40], CUSUM at [41] */ + int state_dim, /* to index ADX at [40], CUSUM at [41] */ + /* Task 2.X "make Full useful" — per-magnitude-bin Q-mean array + + * absolute-scale reference. Produced by q_mag_bin_means_reduce on + * the same stream before isv_signal_update. NULL = freeze slots. */ + const float* __restrict__ q_mag_means_ptr, /* [mag_size] per-bin Q mean */ + const float* __restrict__ q_abs_ref_ptr, /* [1] max(|Q_mean|) across mag bins */ + int mag_size /* typically 3 (Quarter/Half/Full) */ ) { if (threadIdx.x != 0) return; @@ -5624,6 +5630,44 @@ extern "C" __global__ void isv_signal_update( isv_signals[10] = 0.0f; isv_signals[11] = 0.5f; /* moderate stability */ } + + /* ── Task 2.X "make Full useful": per-magnitude-bin Q-mean EMAs [13..15] + * and absolute-Q-scale reference EMA [16] ─────────────────────── + * + * Signal-driven adaptive inputs for the C51 magnitude bin weight. + * [13] Q_mean_ema(Quarter) — persistent preference per bin. + * [14] Q_mean_ema(Half) + * [15] Q_mean_ema(Full) + * [16] |Q|-scale reference EMA (max over bins of |Q_mean|). + * + * Consumed by c51_loss_batched / c51_grad_kernel to derive + * max_mean = max(Q_mean(Quarter), Q_mean(Half), Q_mean(Full)) + * bias_gap = max(0, max_mean − Q_mean(a1)) + * collapse_frac = min(1, bias_gap / max(|Q|-scale, eps)) + * bin_weight = 1 + collapse_frac * mag_bias_signal(a1) + * + * Directly targets C51's "high-variance bins under-priced" signature + * (experience_kernels.cu:~904). Self-disables when Full is the argmax + * — bias_gap becomes 0 for Full samples, bin_weight → 1. + * + * EMA tau=0.05 matches the existing alpha=0.05 pattern (20-step window). + * NULL-guarded — missing inputs leave the slots unchanged (no regression + * against callers that do not supply the pointers). + * + * NOT rotated through isv_history — already EMA-smoothed. isv_forward + * reads only the first 12 slots (the encoder/gate path is unchanged). */ + if (q_mag_means_ptr != 0) { + int ub = (mag_size < 3) ? mag_size : 3; /* clamp to the 3 reserved slots */ + for (int k = 0; k < ub; k++) { + int slot = 13 + k; + isv_signals[slot] = (1.0f - alpha) * isv_signals[slot] + + alpha * q_mag_means_ptr[k]; + } + } + if (q_abs_ref_ptr != 0) { + isv_signals[16] = (1.0f - alpha) * isv_signals[16] + + alpha * fmaxf(q_abs_ref_ptr[0], 0.0f); + } } /* ================================================================== */ diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index e184805f3..8548590f8 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -894,6 +894,48 @@ impl GpuBacktestEvaluator { ]) } + /// 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` + /// (0=Short, 1=Hold, 2=Long, 3=Flat). + /// + /// Returns `[short, hold, long, flat]` summing to 1.0 over non-negative + /// action entries. When eval is strict-argmax and direction Q-values are + /// near-uniform, this distribution helps diagnose whether the eval + /// collapse is direction-side (e.g., `hold + flat > 0.95`) or whether + /// direction is healthy but magnitude itself is collapsing. Paired with + /// `read_eval_action_distribution_per_magnitude` to localise the + /// magnitude-only Task 2.X fix's leverage surface. + pub fn read_eval_action_distribution_per_direction(&self) -> Result<[f32; 4], MLError> { + let len = self.n_windows * self.max_len; + let mut host = vec![0_i32; len]; + self.stream + .memcpy_dtoh(&self.actions_history_buf, &mut host) + .map_err(|e| MLError::ModelError(format!("actions_history dtoh: {e}")))?; + let mut counts = [0_u64; 4]; + let mut total = 0_u64; + for &a in &host { + if a < 0 { + continue; + } + let dir = (a / 27) as usize; + if dir < 4 { + counts[dir] += 1; + total += 1; + } + } + if total == 0 { + return Ok([0.0; 4]); + } + let t = total as f32; + Ok([ + counts[0] as f32 / t, + counts[1] as f32 / t, + counts[2] as f32 / t, + counts[3] as f32 / t, + ]) + } + /// Reset mutable evaluation state: portfolio, done flags, step returns, actions history. /// Must be called before each `evaluate_dqn_graphed` when reusing the evaluator /// across epochs — otherwise done_flags=1 from the previous evaluation causes diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 653fbdfd4..7e3a7f447 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -84,8 +84,44 @@ const OFI_EMBED_DIM: usize = 10; // OFI embedding width appended to history /// Introspective State Vector (ISV) configuration. const ISV_K: usize = 4; // Temporal ISV history length -const ISV_DIM: usize = 13; // ISV signal count (8 core + 4 regime awareness + 1 learning_health) +/// ISV signal count. +/// +/// Layout: +/// [0..7] core training dynamics: q_drift, grad_norm_ema, td_err_ema, +/// ens_var_ema, ens_var_vel, reward_ema, atom_util_ema, loss_ema +/// [8..11] regime awareness (ADX/CUSUM): adx_ema, regime_disagree, +/// regime_vel_ema, regime_stability +/// [12] learning_health (written by separate kernel, consumed by c51/etc.) +/// [13..15] Task 2.X "make Full useful" — per-magnitude-bin Q-value EMAs. +/// [13] = Q_mean_ema(Quarter), [14] = Q_mean_ema(Half), +/// [15] = Q_mean_ema(Full). Populated by `q_mag_bin_means_reduce` +/// on the stats cadence, smoothed with tau=0.05. The C51 loss and +/// gradient kernels derive an adaptive bin weight from these +/// EMAs by computing the bias-vs-Full signature +/// collapse = max(0, max(Q_mean) − Q_mean(a1)) / max(|Q_ref|, eps) +/// applied ONLY to magnitude-branch samples (d==1). Self-disables +/// when Full becomes the argmax bin. +/// [16] q_abs_ref_ema — max(|Q_mean_ema[k]|) across the 3 magnitude +/// bins. Scale reference for the collapse-fraction normaliser +/// (keeps the bin-weight formula scale-invariant as Q-values +/// grow over training). +/// +/// Slots [0..11] populated by `isv_signal_update` kernel and rotated through +/// `isv_history` for temporal decay in `isv_forward`. Slot [12] written +/// outside `isv_signal_update` and NOT rotated into history. Slots [13..16] +/// written by `isv_signal_update` (added with the Task 2.X adaptive magnitude +/// fix) but NOT rotated into history — they are already EMA-smoothed. +const ISV_DIM: usize = 17; pub const LEARNING_HEALTH_INDEX: usize = 12; +/// Task 2.X "make Full useful": ISV slots for per-magnitude-bin Q-mean EMAs. +/// Read by `c51_loss_batched` and `c51_grad_kernel` to compute the +/// bias-vs-Full signal that drives adaptive bin weighting. +pub const Q_MAG_MEAN_QUARTER_INDEX: usize = 13; +pub const Q_MAG_MEAN_HALF_INDEX: usize = 14; +pub const Q_MAG_MEAN_FULL_INDEX: usize = 15; +/// Task 2.X "make Full useful": absolute-Q-scale reference used to +/// normalise the bias signal into a scale-invariant collapse fraction. +pub const Q_ABS_REF_INDEX: usize = 16; const ISV_EMB_DIM: usize = 8; // ISV embedding output dimension (FC2 output, gate/gamma input) /// First ISV tensor index in the flat param buffer. /// ISV weights (68-79) are online-only — NOT synced to the target network. @@ -1652,12 +1688,23 @@ pub struct GpuDqnTrainer { atom_util_scratch_pinned: *mut f32, atom_util_scratch_dev_ptr: u64, + // Task 2.X "make Full useful" — per-magnitude-bin Q-mean scratch arrays + // written by `q_mag_bin_means_reduce` (runs with q_stats on the stats + // cadence), read by isv_signal_update to drive the EMA slots [13..16]. + // The means scratch is sized for magnitude-branch capacity (MAX_MAG=4 + // in the kernel); the abs_ref scratch is a single scalar. + q_mag_means_scratch_pinned: *mut f32, + q_mag_means_scratch_dev_ptr: u64, + q_abs_ref_scratch_pinned: *mut f32, + q_abs_ref_scratch_dev_ptr: u64, + q_mag_bin_means_reduce_kernel: CudaFunction, + // ── ISV core buffers (pinned device-mapped, GPU read/write) ── - isv_signals_pinned: *mut f32, // [ISV_DIM = 13] + isv_signals_pinned: *mut f32, // [ISV_DIM = 15] isv_signals_dev_ptr: u64, - isv_history_pinned: *mut f32, // [ISV_K * ISV_DIM = 52] + isv_history_pinned: *mut f32, // [ISV_K * 12 = 48] — history rotates slots [0..11] only isv_history_dev_ptr: u64, - isv_decay_pinned: *mut f32, // [ISV_DIM = 13] learned decay weights + isv_decay_pinned: *mut f32, // [ISV_DIM = 15] learned decay weights isv_decay_dev_ptr: u64, lagged_td_error_pinned: *mut f32, // [1] recursive confidence target lagged_td_error_dev_ptr: u64, @@ -2142,6 +2189,12 @@ impl Drop for GpuDqnTrainer { if !self.atom_util_scratch_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.atom_util_scratch_pinned.cast()) }; } + if !self.q_mag_means_scratch_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.q_mag_means_scratch_pinned.cast()) }; + } + if !self.q_abs_ref_scratch_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.q_abs_ref_scratch_pinned.cast()) }; + } if !self.isv_signals_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.isv_signals_pinned.cast()) }; } @@ -3014,12 +3067,64 @@ impl GpuDqnTrainer { unsafe { *self.q_mean_ema_pinned = new_ema; } } - /// Pure GPU ISV signal update — reads all pinned pointers, updates ISV [12] + history [K,12]. - /// Includes 4 regime awareness signals computed from ADX/CUSUM in states buffer. + /// Task 2.X "make Full useful" — launch the per-magnitude-bin Q-mean + /// reducer over the current `q_out_buf`. Writes: + /// * `q_mag_means_scratch[0..mag_size]` = mean Q over batch for each + /// magnitude action (Quarter=0, Half=1, Full=2). + /// * `q_abs_ref_scratch[0]` = max(|mean Q|) across magnitude bins — + /// scale reference for the collapse-fraction normaliser. + /// + /// These feed ISV slots [13..16] (EMA) via `update_isv_signals`. The + /// C51 loss / grad kernels then compute a bias-vs-Full signature + /// `max(0, max(Q_mean) − Q_mean(a1)) / Q_abs_ref` — directly targets + /// the "higher-variance bins systematically under-priced" bias cited + /// at experience_kernels.cu:~904 without requiring an inter-branch + /// reference (which failed during Task 2.X iteration 1 because the + /// direction branch collapses at the same cadence as magnitude). + /// + /// Runs outside CUDA Graph capture on the stats cadence. Single-thread + /// single-block, ~N * mag_size f32 loads — microseconds at typical B. + pub(crate) fn launch_q_mag_bin_means_reduce( + &self, + batch_size: usize, + ) -> Result<(), MLError> { + let total_actions = self.total_actions() as i32; + let n = batch_size as i32; + let mag_off = self.config.branch_0_size as i32; + let mag_size = self.config.branch_1_size as i32; + let q_out_buf_ptr = self.q_out_buf.raw_ptr(); + unsafe { + self.stream + .launch_builder(&self.q_mag_bin_means_reduce_kernel) + .arg(&q_out_buf_ptr) + .arg(&self.q_mag_means_scratch_dev_ptr) + .arg(&self.q_abs_ref_scratch_dev_ptr) + .arg(&n) + .arg(&total_actions) + .arg(&mag_off) + .arg(&mag_size) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("q_mag_bin_means_reduce: {e}")))?; + } + Ok(()) + } + + /// Pure GPU ISV signal update — reads all pinned pointers, updates ISV + /// [ISV_DIM=17] + history [K,12]. Includes 4 regime awareness signals + /// computed from ADX/CUSUM in the states buffer. Slots [13..16] are the + /// Task 2.X "make Full useful" fix — per-magnitude-bin Q-mean EMAs and + /// absolute-scale reference, fed by `q_mag_bin_means_reduce` (launched + /// inside `reduce_current_q_stats` just before this kernel). + /// /// Call after graph_adam replay and stream sync, before next graph_forward. pub(crate) fn update_isv_signals(&self) -> Result<(), MLError> { let k = ISV_K as i32; let state_dim = ml_core::state_layout::STATE_DIM as i32; + let mag_size_i32 = self.config.branch_1_size as i32; unsafe { self.stream.launch_builder(&self.isv_signal_update_kernel) .arg(&self.isv_signals_dev_ptr) @@ -3036,6 +3141,12 @@ impl GpuDqnTrainer { .arg(&k) .arg(&self.ptrs.states_buf) // states for ADX/CUSUM regime signals .arg(&state_dim) + // Task 2.X "make Full useful" — per-magnitude-bin Q-mean + // scratch array + absolute-scale reference scalar feed ISV + // slots [13..15] (Q_mean per bin) and [16] (|Q| reference). + .arg(&self.q_mag_means_scratch_dev_ptr) + .arg(&self.q_abs_ref_scratch_dev_ptr) + .arg(&mag_size_i32) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), @@ -6389,6 +6500,9 @@ impl GpuDqnTrainer { // ── Compile expected Q-value + stats kernels (validation, not in CUDA Graph) ─ let expected_q_kernel = compile_expected_q_kernel(&stream)?; let q_stats_kernel = compile_q_stats_kernel(&stream)?; + // Task 2.X "make Full useful" — per-magnitude-bin Q-mean reducer that + // feeds ISV slots [13..16] on the stats cadence. + let q_mag_bin_means_reduce_kernel = compile_q_mag_bin_means_kernel(&stream)?; let q_stats_buf = stream.alloc_zeros::(7) .map_err(|e| MLError::ModelError(format!("alloc q_stats_f32: {e}")))?; let atom_stats_buf = stream.alloc_zeros::(2) @@ -7273,6 +7387,53 @@ impl GpuDqnTrainer { (host_ptr as *mut f32, dev_ptr_out) }; + // Task 2.X "make Full useful" — per-magnitude-bin Q-mean pinned + // array (sized for up to 4 magnitude bins to match + // MAX_MAGNITUDE_ACTIONS in experience_kernels.cu) plus the + // absolute-Q-scale reference scalar. Written by + // `q_mag_bin_means_reduce` (launched with q_stats); read by + // isv_signal_update to populate ISV slots [13..16]. + let q_mag_means_slots: usize = 4; + let (q_mag_means_scratch_pinned, q_mag_means_scratch_dev_ptr) = { + let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); + let mut dev_ptr_out: u64 = 0; + unsafe { + let rc = cudarc::driver::sys::cuMemAllocHost_v2( + &mut host_ptr, + q_mag_means_slots * std::mem::size_of::(), + ); + assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for q_mag_means_scratch"); + let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2( + &mut dev_ptr_out, + host_ptr, + 0, + ); + assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for q_mag_means_scratch"); + std::ptr::write_bytes(host_ptr as *mut u8, 0, q_mag_means_slots * std::mem::size_of::()); + } + (host_ptr as *mut f32, dev_ptr_out) + }; + + let (q_abs_ref_scratch_pinned, q_abs_ref_scratch_dev_ptr) = { + let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); + let mut dev_ptr_out: u64 = 0; + unsafe { + let rc = cudarc::driver::sys::cuMemAllocHost_v2( + &mut host_ptr, + std::mem::size_of::(), + ); + assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for q_abs_ref_scratch"); + let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2( + &mut dev_ptr_out, + host_ptr, + 0, + ); + assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for q_abs_ref_scratch"); + *(host_ptr as *mut f32) = 0.0; + } + (host_ptr as *mut f32, dev_ptr_out) + }; + // ── ISV core buffers (pinned device-mapped) ────────────────────── let (isv_signals_pinned, isv_signals_dev_ptr) = { let num_bytes = ISV_DIM * std::mem::size_of::(); @@ -8176,6 +8337,11 @@ impl GpuDqnTrainer { bw_d_glu_gate, expected_q_kernel, q_stats_kernel, + q_mag_bin_means_reduce_kernel, + q_mag_means_scratch_pinned, + q_mag_means_scratch_dev_ptr, + q_abs_ref_scratch_pinned, + q_abs_ref_scratch_dev_ptr, q_stats_buf, atom_stats_buf, atom_stats_block_sums_buf, @@ -8630,6 +8796,27 @@ impl GpuDqnTrainer { } } + /// Task 2.X "make Full useful" diagnostic — read the per-magnitude-bin + /// Q-mean EMAs and absolute-scale reference from ISV (pinned, zero-copy). + /// + /// Returns `(q_mean_quarter, q_mean_half, q_mean_full, q_abs_ref)`. The + /// C51 loss / gradient kernels use these to compute the bin weight: + /// max_mean = max(q_mean_*) + /// bias_gap = max(0, max_mean − q_mean[a1]) + /// collapse_frac = min(1, bias_gap / max(q_abs_ref, 1e-6)) + /// bin_weight = 1 + collapse_frac * (a1 + 1) / b1_size + /// Returns all zeros if ISV is not allocated. + pub fn read_isv_magnitude_bin_q_means(&self) -> (f32, f32, f32, f32) { + if self.isv_signals_pinned.is_null() { return (0.0, 0.0, 0.0, 0.0); } + unsafe { + let q = *self.isv_signals_pinned.add(Q_MAG_MEAN_QUARTER_INDEX); + let h = *self.isv_signals_pinned.add(Q_MAG_MEAN_HALF_INDEX); + let f = *self.isv_signals_pinned.add(Q_MAG_MEAN_FULL_INDEX); + let r = *self.isv_signals_pinned.add(Q_ABS_REF_INDEX); + (q, h, f, r.max(0.0)) + } + } + /// B4/G5: Read ISV health index and regime stability from pinned host memory. /// Returns (health [0,1], regime_stability [0,1]). Falls back to (0.5, 0.5) if /// pinned pointer is null. @@ -10368,6 +10555,13 @@ impl GpuDqnTrainer { // ISV: write Q-mean as reward proxy to pinned scratch unsafe { *self.reward_scratch_pinned = host[3]; } + // Task 2.X "make Full useful" — per-magnitude-bin Q-mean reducer + // writes the 3 per-bin Q-means + absolute-scale reference to + // pinned scratch. Runs on the same stream; isv_signal_update + // reads the values directly via pinned device pointers (no sync). + let mag_means_batch_size = self.config.batch_size; + self.launch_q_mag_bin_means_reduce(mag_means_batch_size)?; + // ISV signal update — pure GPU, reads all pinned scalars, updates ISV vector self.update_isv_signals()?; @@ -12262,6 +12456,10 @@ impl GpuDqnTrainer { .arg(&liquid_mod_buf_ptr) .arg(&atom_positions_buf_ptr) .arg(&self.q_mean_ema_dev_ptr) // unused but preserves 19-param graph node structure on Hopper + // Task 2.X adaptive magnitude fix — ISV signal bus; kernel reads + // slots [13] (q_mag_spread_ema) and [14] (q_dir_spread_ema) to + // derive the magnitude-branch adaptive bin weight. + .arg(&self.isv_signals_dev_ptr) .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), @@ -14039,6 +14237,22 @@ fn compile_q_stats_kernel( .map_err(|e| MLError::ModelError(format!("q_stats_reduce load: {e}"))) } +/// Load the per-magnitude-bin Q-mean reducer from the q_stats cubin. +/// +/// Task 2.X "make Full useful": this kernel computes batch-mean Q-values +/// per magnitude-bin (Quarter/Half/Full) plus an absolute-scale reference, +/// feeding the four ISV slots [13..16] that drive the adaptive C51 +/// magnitude bin weight. See q_stats_kernel.cu for the kernel body. +fn compile_q_mag_bin_means_kernel( + stream: &Arc, +) -> Result { + let context = stream.context(); + let module = context.load_cubin(Q_STATS_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("q_mag_bin_means cubin load: {e}")))?; + module.load_function("q_mag_bin_means_reduce") + .map_err(|e| MLError::ModelError(format!("q_mag_bin_means_reduce load: {e}"))) +} + // ── Shared memory sizing ──────────────────────────────────────────────────── /// Query the hardware's max shared memory per block via cuDeviceGetAttribute. diff --git a/crates/ml/src/cuda_pipeline/q_stats_kernel.cu b/crates/ml/src/cuda_pipeline/q_stats_kernel.cu index 9ad9d06b1..5c3775848 100644 --- a/crates/ml/src/cuda_pipeline/q_stats_kernel.cu +++ b/crates/ml/src/cuda_pipeline/q_stats_kernel.cu @@ -60,3 +60,97 @@ extern "C" __global__ void q_stats_reduce( out[5] = (atom_stats != NULL) ? atom_stats[0] * inv_n / max_entropy : 0.0f; /* normalized entropy [0,1] */ out[6] = (atom_stats != NULL) ? atom_stats[1] * inv_n / (float)num_atoms : 0.0f; /* utilization fraction [0,1] */ } + +/** + * Per-branch magnitude-Q mean reducer: computes batch-mean Q-values for each + * of the 3 magnitude bins separately (Quarter=0, Half=1, Full=2), plus the + * absolute Q-scale reference, from `q_values` of shape [N, total_actions] + * sliced at `mag_off..mag_off+mag_size`. + * + * Task 2.X "make Full useful" adaptive fix: feeds the ISV signal bus with + * the four scalars needed to compute a DIRECT anti-bias signal for C51's + * magnitude-branch gradient: + * + * bias_vs_full = max(0, max_mag(q_mean[k]) − q_mean[Full]) + * + * When a higher-variance bin (Full=2) is priced below the max magnitude + * bin, `bias_vs_full > 0` — the exact signature of C51's structural + * variance bias (high-variance bins get systematically lower E[Q] because + * their support centering favours narrower distributions). When Full is + * already the argmax, `bias_vs_full == 0` and the mechanism self-disables. + * + * The reference scale `q_abs_ref` is the max of per-bin |Q_mean|, used to + * normalise `bias_vs_full` into a scale-invariant [0, 1] collapse fraction. + * + * Runs OUTSIDE the CUDA Graph (alongside q_stats_reduce). isv_signal_update + * folds the scalars into EMA slots [13] (Q(Quarter)), [14] (Q(Half)). No + * separate ISV slot for Q(Full) or the absolute scale — those are recovered + * inside the C51 kernels by reading per-sample q_values directly (the loss + * kernel already has access to the full magnitude-branch slice through its + * advantage inputs). The two EMAs we expose are the ones the bias detector + * NEEDS against the current batch's Q(Full): slow-moving references that + * reflect the policy's persistent preference rather than per-step noise. + * + * (Naming history: this kernel was originally scoped as + * `q_branch_spread_reduce` producing per-branch max-min spreads, but using + * direction-branch spread as the reference for magnitude collapse failed — + * observed `dir_spread=0.002 ≪ mag_spread=0.08` because the direction head + * itself collapses at the same cadence the magnitude head does, so the + * ratio-based formulation never triggered. The per-bin Q-mean + * formulation below targets the precise bias signature cited in the Task + * 2.X scoping doc and at experience_kernels.cu:~904 — high-variance bins + * under-priced by C51's expected-Q — without requiring an inter-branch + * reference.) + * + * Pure reduction — no training-path feedback and no atomicAdd. Single-block, + * single-thread to match q_stats_reduce's determinism. + * + * Output: q_mag_means_out[0] = Q_mean(Quarter) over N samples, + * q_mag_means_out[1] = Q_mean(Half), + * q_mag_means_out[2] = Q_mean(Full). + * q_abs_ref_out[0] = max(|Q_mean[k]|) across magnitude bins + * (absolute scale reference; floored at 1e-6 + * downstream). + * + * Launch: grid=(1,1,1), block=(1,1,1). N=batch_size usually 256..16384 — + * stays well under any block timeout. Cost: ~N * mag_size FP32 loads, + * single-digit microseconds at typical B. + */ +extern "C" __global__ void q_mag_bin_means_reduce( + const float* __restrict__ q_values, /* [N, total_actions] */ + float* __restrict__ q_mag_means_out, /* [mag_size] — typically [3] */ + float* __restrict__ q_abs_ref_out, /* [1] absolute-value max of bin means */ + int N, + int total_actions, + int mag_off, int mag_size) +{ + if (threadIdx.x != 0) return; + + /* Hard ceiling — covers any future magnitude-branch size extension + * without an unbounded stack array. Matches MAX_MAGNITUDE_ACTIONS=4 + * used in experience_kernels.cu at the Boltzmann sampling site. */ + const int MAX_MAG = 4; + float sums[MAX_MAG]; + for (int k = 0; k < MAX_MAG; k++) sums[k] = 0.0f; + + if (N <= 0 || total_actions <= 0 || mag_size <= 0 || mag_size > MAX_MAG + || mag_off < 0 || mag_off + mag_size > total_actions) { + for (int k = 0; k < mag_size && k < MAX_MAG; k++) q_mag_means_out[k] = 0.0f; + q_abs_ref_out[0] = 0.0f; + return; + } + + for (int i = 0; i < N; i++) { + const float* row = q_values + (long long)i * total_actions; + for (int k = 0; k < mag_size; k++) sums[k] += row[mag_off + k]; + } + + float inv_n = 1.0f / (float)N; + float abs_ref = 0.0f; + for (int k = 0; k < mag_size; k++) { + float mean_k = sums[k] * inv_n; + q_mag_means_out[k] = mean_k; + abs_ref = fmaxf(abs_ref, fabsf(mean_k)); + } + q_abs_ref_out[0] = abs_ref; +} diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 266f9e342..fe150b25f 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -2607,6 +2607,16 @@ impl FusedTrainingCtx { .map_err(|e| anyhow::anyhow!("launch_pad_states: {e}")) } + /// Task 2.X diagnostic — read-only borrow of the inner `GpuDqnTrainer`. + /// Used by `DQNTrainer::last_isv_branch_q_spreads` to surface ISV + /// slots [13]/[14] to smoke tests without widening the field's + /// visibility. Do NOT add write paths through this accessor — the + /// trainer owns its CUDA context and mutation must go through + /// properly ordered submission paths. + pub(crate) fn trainer(&self) -> &GpuDqnTrainer { + &self.trainer + } + /// Raw pointer to trainer's states buffer. pub(crate) fn trainer_states_buf_ptr(&self) -> u64 { self.trainer.states_buf_ptr() 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 a8116a176..7e886e87e 100644 --- a/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs +++ b/crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs @@ -65,6 +65,37 @@ fn test_magnitude_distribution() -> Result<()> { // feedback_no_functionality_removal.md, the branch is never deleted as a fallback. let [eq, eh, ef] = trainer.last_eval_magnitude_dist(); println!("[EVAL_DIST] Quarter={:.3} Half={:.3} Full={:.3}", eq, eh, ef); + + // 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- + // only bin weight cannot shift eval_dist. Surfacing this lets the + // operator attribute the failure to direction vs magnitude. + if let Some(dir_dist) = trainer.last_eval_direction_dist() { + println!( + "[EVAL_DIR_DIST] Short={:.3} Hold={:.3} Long={:.3} Flat={:.3}", + dir_dist[0], dir_dist[1], dir_dist[2], dir_dist[3] + ); + } + + // Task 2.X diagnostic — ISV per-magnitude-bin Q-mean EMAs + |Q|-scale + // reference. The bin weight fires when Full is NOT the argmax bin of + // (Quarter / Half / Full), so the training-policy's Q-means directly + // drive the adaptive correction. + if let Some((qq, qh, qf, q_abs_ref)) = trainer.last_isv_magnitude_bin_q_means() { + let max_mean = qq.max(qh).max(qf); + let bias_gap_full = (max_mean - qf).max(0.0); + let collapse_frac_full = if q_abs_ref > 1e-6 { + (bias_gap_full / q_abs_ref).min(1.0) + } else { + 0.0 + }; + println!( + "[ISV_BIN_MEANS] q_q={:.4} q_h={:.4} q_f={:.4} q_abs_ref={:.4} \ + max_mean={:.4} bias_gap_full={:.4} collapse_frac_full={:.3}", + qq, qh, qf, q_abs_ref, max_mean, bias_gap_full, collapse_frac_full + ); + } assert!( eh + ef >= 0.30, "H10 regression: eval Half + Full share {:.3} < 0.30 (eq={:.3} eh={:.3} ef={:.3}) \ @@ -74,5 +105,30 @@ fn test_magnitude_distribution() -> Result<()> { make the magnitude branch useful, not remove it.", eh + ef, eq, eh, ef ); + + // Task 2.X "make Full useful" — the adaptive ISV-driven C51 bin weight + // (see gpu_dqn_trainer.rs Q_MAG_SPREAD_INDEX / Q_DIR_SPREAD_INDEX and + // c51_loss_kernel.cu get_magnitude_bin_weight) must push Q(Full) high + // enough that the eval-mode strict-argmax actually selects Full at + // least 5% of the time. Pre-fix baseline observed ef=0.000 because + // Q(Half) > Q(Full) + 1e-6 deterministically across all states after + // Task-2.2 H10 tie-break. The fix is signal-driven (deficit between + // magnitude and direction Q-spread) with a bin-proportional weight, + // so it self-disables once the magnitude head differentiates. + assert!( + ef >= 0.05, + "Task 2.X adaptive magnitude: eval Full share {:.3} < 0.05 \ + (eq={:.3} eh={:.3} ef={:.3}) — ISV bin-weight mechanism failed to \ + lift Q(Full) above Q(Half) + 1e-6. Investigate: \ + (a) ISV slots [13]/[14] populated each stats cadence (check \ + q_branch_spread_reduce + isv_signal_update wiring), \ + (b) spread_deficit positive during training (magnitude-branch \ + Q-spread collapsed below direction-branch reference), \ + (c) bin_weight applied consistently to branch_ce (c51_loss_batched) \ + AND gradient (c51_grad_kernel d==1 block). Do NOT fall back to \ + static tuning knobs — per feedback_adaptive_not_tuned.md the \ + escalation path is richer ISV signals / kernel-side arithmetic.", + ef, eq, eh, ef + ); Ok(()) } diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index 888df1414..6a111f399 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/crates/ml/src/trainers/dqn/trainer/mod.rs @@ -1524,6 +1524,28 @@ impl DQNTrainer { self.last_eval_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 + /// this lets the test surface whether the root cause is direction-side + /// (Hold + Flat > 0.95) or magnitude-side (direction healthy but mag + /// collapsed). Returns `None` if no GPU evaluator is wired. + pub fn last_eval_direction_dist(&self) -> Option<[f32; 4]> { + self.gpu_evaluator + .as_ref() + .and_then(|ev| ev.read_eval_action_distribution_per_direction().ok()) + } + + /// Task 2.X diagnostic — ISV slots [13..16] (per-magnitude-bin Q-mean + /// EMAs + absolute-scale reference). Returns + /// `(q_mean_quarter, q_mean_half, q_mean_full, q_abs_ref)` or `None` + /// if the fused trainer has not been initialised yet. + pub fn last_isv_magnitude_bin_q_means(&self) -> Option<(f32, f32, f32, f32)> { + self.fused_ctx + .as_ref() + .map(|ctx| ctx.trainer().read_isv_magnitude_bin_q_means()) + } + /// Per-epoch magnitude-branch action entropy history. /// Each entry is (epoch_0_indexed, normalized_entropy ∈ [0, 1]). /// Used by exploration_coverage smoke test.