diff --git a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu index 3cd594823..f9fe2f683 100644 --- a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu @@ -29,10 +29,12 @@ extern "C" __global__ void c51_grad_kernel( 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 */ /* Task 2.X adaptive magnitude fix + Task 2.Y adaptive direction fix — - * ISV signal bus [ISV_DIM=22]. Task 2.X slots [13..16] drive an adaptive + * ISV signal bus [ISV_DIM=23]. Task 2.X slots [13..16] drive an adaptive * bin weight applied to the magnitude branch (d==1); Task 2.Y slots * [17..21] drive an adaptive bin weight applied to the direction branch - * (d==0). NULL → identity (safe: reduces to pre-fix behaviour). */ + * (d==0). Slot 22 is SHARPE_EMA, consumed on-GPU by isv_signal_update to + * drive slot 12 (learning_health). NULL → identity (safe: reduces to + * pre-fix behaviour). */ const float* __restrict__ isv_signals) { int tid = blockIdx.x * blockDim.x + threadIdx.x; diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index dbf403291..f08fec235 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -587,7 +587,7 @@ extern "C" __global__ void c51_loss_batched( const float* __restrict__ cvar_alpha_buf, /* [B] per-sample CVaR alpha from risk branch. NULL = use iqn_readiness. */ /* ── B3/G4: ISV signals for health-scaled Expected SARSA temperature ── */ - const float* __restrict__ isv_signals /* [ISV_DIM=13] pinned device-mapped. isv_signals[12] = learning_health. NULL = 0.5 fallback. */ + const float* __restrict__ isv_signals /* [ISV_DIM=23] pinned device-mapped. isv_signals[12] = learning_health. NULL = 0.5 fallback. */ ) { extern __shared__ float shmem_f[]; @@ -1151,7 +1151,7 @@ extern "C" __global__ void barrier_gradient_direction( const float* __restrict__ adv_logits, /* [B, total_actions, num_atoms] f32 (direction is first b0_size slots) */ const float* __restrict__ v_logits, /* [B, num_atoms] f32 */ const float* __restrict__ z_vals, /* [num_atoms] f32 */ - const float* __restrict__ isv_signals, /* [ISV_DIM=13] pinned; [12]=health */ + const float* __restrict__ isv_signals, /* [ISV_DIM=23] pinned; [12]=health */ float* __restrict__ d_adv_logits, /* [B, total_actions, num_atoms] — atomicAdd direction slice only */ float* __restrict__ d_v_logits, /* [B, num_atoms] — atomicAdd here */ int batch_size, @@ -1302,7 +1302,7 @@ extern "C" __global__ void ib_gradient_direction( const float* __restrict__ adv_logits, /* [B, total_actions, num_atoms] f32 */ const float* __restrict__ v_logits, /* [B, num_atoms] f32 */ const float* __restrict__ z_vals, /* [num_atoms] f32 */ - const float* __restrict__ isv_signals, /* [ISV_DIM=13] pinned; [12]=health */ + const float* __restrict__ isv_signals, /* [ISV_DIM=23] pinned; [12]=health */ float* __restrict__ d_adv_logits, /* [B, total_actions, num_atoms] — atomicAdd direction slice */ float* __restrict__ d_v_logits, /* [B, num_atoms] — atomicAdd */ int batch_size, diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index da13f31cb..ecee6f703 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -5620,20 +5620,21 @@ extern "C" __global__ void risk_budget_backward( /* ================================================================== */ extern "C" __global__ void isv_signal_update( - float* __restrict__ isv_signals, /* [ISV_DIM=22] pinned device-mapped */ + float* __restrict__ isv_signals, /* [ISV_DIM=23] 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 */ const float* __restrict__ grad_norm_ptr, /* [1] pinned — gradient norm² */ const float* __restrict__ td_error_ptr, /* [1] pinned — batch mean |TD-error| */ - const float* __restrict__ ens_var_ptr, /* [1] pinned — batch mean ensemble var */ + const float* __restrict__ q_var_ptr, /* [1] pinned — batch mean C51 Q-distribution variance */ const float* __restrict__ reward_ptr, /* [1] pinned — reward proxy */ const float* __restrict__ atom_util_ptr, /* [1] pinned — atom utilization */ 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] */ + const float* __restrict__ states_ptr, /* [B, state_dim_padded] — for ADX/CUSUM regime signals. NULL = skip regime. */ + int state_dim, /* stride of states_ptr — use STATE_DIM_PADDED, NOT logical STATE_DIM. ADX at [40], CUSUM at [41]. */ + int batch_size, /* B — used to batch-aggregate regime signals across all samples, not just sample 0. */ /* 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. */ @@ -5673,12 +5674,20 @@ extern "C" __global__ void isv_signal_update( /* [2] TD-error EMA */ isv_signals[2] = (1.0f - alpha) * isv_signals[2] + alpha * td_error_ptr[0]; - /* [3] Ensemble variance EMA (save old for delta) */ - float old_ens_ema = isv_signals[3]; - isv_signals[3] = (1.0f - alpha) * old_ens_ema + alpha * ens_var_ptr[0]; + /* [3] C51 Q-distribution variance EMA (save old for delta) + * + * Semantics: batch-mean of the atom-spread variance computed by the C51 + * distributional head — NOT multi-head ensemble variance. The source is + * `q_var_buf_trainer` (size b*12, reduced by a dedicated c51_loss_reduce + * call that scales by b*12). Slot 3's consumers (ISV encoder MLP only) + * want a scale signal for training uncertainty, which the C51 Q-variance + * proxies correctly. True multi-head ensemble disagreement would be a + * separate scratch path if/when a per-head ensemble is wired. */ + float old_qvar_ema = isv_signals[3]; + isv_signals[3] = (1.0f - alpha) * old_qvar_ema + alpha * q_var_ptr[0]; - /* [4] Ensemble variance velocity */ - isv_signals[4] = (isv_signals[3] - old_ens_ema) / fmaxf(old_ens_ema, 1e-6f); + /* [4] C51 Q-distribution variance velocity (derivative of slot 3) */ + isv_signals[4] = (isv_signals[3] - old_qvar_ema) / fmaxf(old_qvar_ema, 1e-6f); /* [5] Reward EMA */ isv_signals[5] = (1.0f - alpha) * isv_signals[5] + alpha * reward_ptr[0]; @@ -5689,10 +5698,29 @@ extern "C" __global__ void isv_signal_update( /* [7] Loss EMA */ isv_signals[7] = (1.0f - alpha) * isv_signals[7] + alpha * loss_ptr[0]; - /* ── Regime awareness signals [8]-[11] from ADX/CUSUM ──────────── */ - if (states_ptr != 0 && state_dim > 41) { - float adx = states_ptr[0 * state_dim + 40]; - float cusum = states_ptr[0 * state_dim + 41]; + /* ── Regime awareness signals [8]-[11] from ADX/CUSUM ──────────── * + * + * Batch-aggregated across all B samples (ISV audit 2026-04-23): the + * prior implementation read only sample 0 and threw away 99.994% of the + * regime information, making the signal noisy at the batch-shuffle + * level (sample 0 is effectively random after per-epoch shuffles). + * + * The launch passes `state_dim = STATE_DIM_PADDED = 128` (the true + * stride of `states_buf`). The prior code passed the logical + * `STATE_DIM = 104`, which only produced correct values for row 0 by + * accident (both strides land at the same offset for sample 0). */ + if (states_ptr != 0 && state_dim > 41 && batch_size > 0) { + /* Batch-mean ADX and CUSUM across all B rows (single-thread loop — + * kernel already runs on 1 thread, microsecond cost at B=16384). */ + float adx_sum = 0.0f; + float cusum_sum = 0.0f; + for (int b = 0; b < batch_size; b++) { + adx_sum += states_ptr[(long long)b * state_dim + 40]; + cusum_sum += states_ptr[(long long)b * state_dim + 41]; + } + float inv_b = 1.0f / (float)batch_size; + float adx = adx_sum * inv_b; + float cusum = cusum_sum * inv_b; /* [8] Regime velocity: EMA of ADX (used for delta computation next step) */ float prev_adx = isv_signals[8]; @@ -5721,6 +5749,47 @@ extern "C" __global__ void isv_signal_update( isv_signals[11] = 0.5f; /* moderate stability */ } + /* ── [12] Learning health EMA driven by sharpe_ema (slot 22) ────── + * + * ISV audit 2026-04-23: the prior Rust-side HealthEma aggregator + * fed slot 12 with seven-component health (q_gap, q_var, atom_util, + * grad_norm, ens_disagreement, grad_consistency, spectral_gap). + * Observed on train-mdh86 (20 epochs): health climbed monotonically + * 0.50 → 0.64 while Sharpe cratered +34 → −67 (Pearson r = −0.765, + * strongly ANTI-correlated). Components q_gap=1.0 and q_var=1.0 were + * saturated at the upper bound for 19/20 epochs; grad_stable was + * stuck at 0.0 for 20/20 epochs. The aggregate could not track + * outcomes because its inputs were clipped / dead. + * + * Replacement: sigmoid(0.1 × sharpe_ema) EMA-blended with the + * previous slot 12 value. The 0.1 coefficient places the sigmoid's + * sensitive region in the [-10, +10] Sharpe range, which matches + * the typical magnitudes observed in this system (Sharpe=0 → + * health=0.5, Sharpe=±10 → health≈0.27/0.73). No external scale + * reference needed — sigmoid saturates gracefully at extremes. + * + * Host Rust writes `training_sharpe_ema` into slot 22 + * (SHARPE_EMA_INDEX) every epoch via `write_isv_signal_at`; this + * kernel reads it here and couples it into slot 12 which is the + * heaviest ISV consumer in the codebase (C51 bin weighting, reward- + * bias damping, CQL gate, distillation alpha, Kelly warmup floor). + * + * PRIOR HEALTH FORMULA (replaced 2026-04-23 per ISV audit): + * Rust-side training_loop.rs:1954: + * let health_value = self.learning_health.update(&raw); + * fused.write_isv_signal_at(LEARNING_HEALTH_INDEX, health_value); + * where raw contained 7 saturating components. See HealthEma in + * crates/ml/src/cuda_pipeline/learning_health.rs for the old + * aggregation; the Rust-side writer is preserved but is now + * overwritten on-GPU every epoch by this block. */ + { + float sharpe_ema_val = isv_signals[22]; + float s_scaled = 0.1f * sharpe_ema_val; + float new_health_target = 1.0f / (1.0f + expf(-s_scaled)); + isv_signals[12] = (1.0f - alpha) * isv_signals[12] + + alpha * new_health_target; + } + /* ── Task 2.X "make Full useful": per-magnitude-bin Q-mean EMAs [13..15] * and absolute-Q-scale reference EMA [16] ─────────────────────── * @@ -5755,8 +5824,18 @@ extern "C" __global__ void isv_signal_update( } } if (q_abs_ref_ptr != 0) { - isv_signals[16] = (1.0f - alpha) * isv_signals[16] - + alpha * fmaxf(q_abs_ref_ptr[0], 0.0f); + /* Outlier clamp (ISV audit 2026-04-23): a single Q-scale spike of + * ±10^5 at α=0.05 previously took ~20 epochs to decay, silently + * inflating the C51 bin_weight denominator and disabling the + * adaptive mechanism for the entire run. Clamp the raw input to + * 10×current EMA + 1 so one bad epoch's pollution is bounded to + * 10× its own magnitude rather than leaking for 20+ epochs. + * +1 floor handles the cold-start case where EMA is near 0. */ + float cur = isv_signals[16]; + float raw = fmaxf(q_abs_ref_ptr[0], 0.0f); + float cap = 10.0f * cur + 1.0f; + float clamped = (raw < cap) ? raw : cap; + isv_signals[16] = (1.0f - alpha) * cur + alpha * clamped; } /* ── Task 2.Y "make direction branch useful at eval": per-direction Q-mean @@ -5792,8 +5871,19 @@ extern "C" __global__ void isv_signal_update( } } if (q_dir_abs_ref_ptr != 0) { - isv_signals[21] = (1.0f - alpha) * isv_signals[21] - + alpha * fmaxf(q_dir_abs_ref_ptr[0], 0.0f); + /* Outlier clamp (ISV audit 2026-04-23): same rationale as slot 16 + * above. Slot 21 is especially load-bearing because it feeds the + * Kelly conviction denominator (`conviction = q_range / + * q_dir_abs_ref` at experience_kernels.cu:1189), so a single + * ±10^5 direction-Q excursion previously pushed conviction → ~0 + * for ~20 epochs, triggering the Kelly warmup-floor collapse + * path and shrinking positions exactly when the policy was + * diverging. */ + float cur = isv_signals[21]; + float raw = fmaxf(q_dir_abs_ref_ptr[0], 0.0f); + float cap = 10.0f * cur + 1.0f; + float clamped = (raw < cap) ? raw : cap; + isv_signals[21] = (1.0f - alpha) * cur + alpha * clamped; } } diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 8ab8e0957..e46b7fa1b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -124,13 +124,16 @@ const ISV_K: usize = 4; // Temporal ISV history length /// collapse-fraction normaliser. /// /// 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. -/// Slots [17..21] likewise written by `isv_signal_update` (added with the -/// Task 2.Y direction-branch fix) and NOT rotated into history. -const ISV_DIM: usize = 22; +/// `isv_history` for temporal decay in `isv_forward`. Slot [12] is now written +/// by `isv_signal_update` from `sharpe_ema` (slot 22) 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. Slots [17..21] likewise written by +/// `isv_signal_update` (added with the Task 2.Y direction-branch fix) and +/// NOT rotated into history. Slot [22] (`SHARPE_EMA_INDEX`) added by the +/// ISV-audit bundle (2026-04-23) — Rust host broadcasts `training_sharpe_ema` +/// each epoch; the kernel reads it to drive slot 12 (learning_health). +const ISV_DIM: usize = 23; 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 @@ -154,6 +157,13 @@ pub const Q_DIR_MEAN_FLAT_INDEX: usize = 20; /// `Q_ABS_REF_INDEX` for magnitude). Keeps the direction-branch /// collapse-fraction scale-invariant as Q-values grow over training. pub const Q_DIR_ABS_REF_INDEX: usize = 21; +/// ISV-audit bundle (2026-04-23): slot carrying host-side `training_sharpe_ema`. +/// Broadcast each epoch by `training_loop.rs` via `write_isv_signal_at`, then +/// read on-GPU by `isv_signal_update` to drive slot 12 (learning_health). +/// Couples health directly to realised training outcomes, replacing the +/// component-aggregation formula whose Sharpe-correlation measured at +/// r=-0.765 on the train-mdh86 logs that triggered the audit. +pub const SHARPE_EMA_INDEX: usize = 22; 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. @@ -1768,22 +1778,21 @@ pub struct GpuDqnTrainer { risk_adam_step: i32, // ── ISV scratch scalars (pinned device-mapped, written by GPU kernels) ── - // Zero-initialised scratch. No writer currently emits batch mean - // |TD-error| into this slot, so `isv_signals[2]` (the TD-error EMA - // in `isv_signal_update`) accumulates a constant-zero signal and - // effectively decays to zero. Downstream consumers treat ISV[2] as - // an unavailable signal. Keeping the buffer allocated preserves the - // `isv_signal_update` kernel ABI — a future C51-loss writeback can - // populate it without a signature change. + // Batch mean |TD-error| scratch. Written by the second + // `c51_loss_reduce` launch at `launch_loss_reduce` (commit + // 7f92fa242). Feeds ISV[2] TD-error EMA in `isv_signal_update`. td_error_scratch_pinned: *mut f32, td_error_scratch_dev_ptr: u64, - // Zero-initialised scratch for mean ensemble variance. Same - // situation as `td_error_scratch_pinned`: no kernel currently - // writes into it, so `isv_signals[3]` (ensemble-variance EMA) and - // `isv_signals[4]` (variance velocity) stay at zero. Buffer kept - // for `isv_signal_update` ABI stability. - ensemble_var_scratch_pinned: *mut f32, - ensemble_var_scratch_dev_ptr: u64, + // C51 Q-distribution variance scratch (ISV audit 2026-04-23: renamed + // from `ensemble_var_scratch_pinned`). Written by a third + // `c51_loss_reduce` launch at `launch_loss_reduce` that reduces + // `q_var_buf_trainer` (size b*12, atom-spread variance per action). + // Feeds ISV[3] "C51 Q-variance EMA" and ISV[4] velocity in + // `isv_signal_update`. NOT multi-head ensemble variance — the + // original misnomer led to the signal being treated as a zeroed + // dead path for the 20-epoch window analysed in the audit. + q_var_scratch_pinned: *mut f32, + q_var_scratch_dev_ptr: u64, reward_scratch_pinned: *mut f32, reward_scratch_dev_ptr: u64, atom_util_scratch_pinned: *mut f32, @@ -1812,11 +1821,11 @@ pub struct GpuDqnTrainer { q_dir_bin_means_reduce_kernel: CudaFunction, // ── ISV core buffers (pinned device-mapped, GPU read/write) ── - isv_signals_pinned: *mut f32, // [ISV_DIM = 15] + isv_signals_pinned: *mut f32, // [ISV_DIM = 23] isv_signals_dev_ptr: u64, 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 = 15] learned decay weights + isv_decay_pinned: *mut f32, // [ISV_DIM = 23] learned decay weights isv_decay_dev_ptr: u64, lagged_td_error_pinned: *mut f32, // [1] recursive confidence target lagged_td_error_dev_ptr: u64, @@ -2292,8 +2301,8 @@ impl Drop for GpuDqnTrainer { if !self.td_error_scratch_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.td_error_scratch_pinned.cast()) }; } - if !self.ensemble_var_scratch_pinned.is_null() { - let _ = unsafe { cudarc::driver::result::free_host(self.ensemble_var_scratch_pinned.cast()) }; + if !self.q_var_scratch_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.q_var_scratch_pinned.cast()) }; } if !self.reward_scratch_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.reward_scratch_pinned.cast()) }; @@ -3310,16 +3319,27 @@ impl GpuDqnTrainer { } /// 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). + /// [ISV_DIM=23] + history [K,12]. Includes 4 regime awareness signals + /// computed from ADX/CUSUM in the states buffer (batch-aggregated across + /// all B rows; ISV audit 2026-04-23 replaced the prior sample-0 read + /// that threw away 99.994% of the regime information). 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). + /// + /// The `state_dim` kernel argument is passed as + /// `ml_core::state_layout::STATE_DIM_PADDED` (128) — the actual stride + /// of `states_buf`. The previous launch passed the logical + /// `STATE_DIM` (104), which only produced correct values for row 0 by + /// accident because both strides place sample 0 at the same byte + /// offset. With the batch-aggregate fix all B rows must be strided + /// correctly. /// /// 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 state_dim = ml_core::state_layout::STATE_DIM_PADDED as i32; + let batch_size_i32 = self.config.batch_size as i32; let mag_size_i32 = self.config.branch_1_size as i32; let dir_size_i32 = self.config.branch_0_size as i32; unsafe { @@ -3331,13 +3351,14 @@ impl GpuDqnTrainer { .arg(&self.q_mean_ema_dev_ptr) .arg(&self.grad_norm_dev_ptr) .arg(&self.td_error_scratch_dev_ptr) - .arg(&self.ensemble_var_scratch_dev_ptr) + .arg(&self.q_var_scratch_dev_ptr) .arg(&self.reward_scratch_dev_ptr) .arg(&self.atom_util_scratch_dev_ptr) .arg(&self.total_loss_dev_ptr) .arg(&k) .arg(&self.ptrs.states_buf) // states for ADX/CUSUM regime signals .arg(&state_dim) + .arg(&batch_size_i32) // 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). @@ -7622,7 +7643,7 @@ impl GpuDqnTrainer { (host_ptr as *mut f32, dev_ptr_out) }; - let (ensemble_var_scratch_pinned, ensemble_var_scratch_dev_ptr) = { + let (q_var_scratch_pinned, q_var_scratch_dev_ptr) = { let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); let mut dev_ptr_out: u64 = 0; unsafe { @@ -7630,13 +7651,13 @@ impl GpuDqnTrainer { &mut host_ptr, std::mem::size_of::(), ); - assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for ensemble_var_scratch"); + assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for q_var_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 ensemble_var_scratch"); + assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for q_var_scratch"); *(host_ptr as *mut f32) = 0.0; } (host_ptr as *mut f32, dev_ptr_out) @@ -8879,8 +8900,8 @@ impl GpuDqnTrainer { risk_adam_step: 0, td_error_scratch_pinned, td_error_scratch_dev_ptr, - ensemble_var_scratch_pinned, - ensemble_var_scratch_dev_ptr, + q_var_scratch_pinned, + q_var_scratch_dev_ptr, reward_scratch_pinned, reward_scratch_dev_ptr, atom_util_scratch_pinned, @@ -12907,15 +12928,25 @@ impl GpuDqnTrainer { /// Deterministic loss reduction: sequential sum of per_sample_loss → total_loss / batch_size. /// Grid=(1,1,1), Block=(1,1,1). Must stay single-thread inside graph_mega on Hopper. /// - /// Also emits the batch mean |TD-error| into `td_error_scratch_dev_ptr` - /// by running the same reduction kernel over `td_errors_buf`. This feeds - /// ISV[2] (the TD-error EMA in `isv_signal_update`), which downstream - /// consumers previously saw as a constant-zero signal because the scratch - /// was never written. + /// Also emits two additional pinned scratch scalars on the same stream: + /// * Second launch writes batch mean |TD-error| into + /// `td_error_scratch_dev_ptr` (reducing `td_errors_buf` [B]). + /// Feeds ISV[2] (TD-error EMA in `isv_signal_update`). + /// * Third launch writes batch mean C51 Q-distribution variance into + /// `q_var_scratch_dev_ptr` (reducing `q_var_buf_trainer` [B*12]). + /// Feeds ISV[3] and ISV[4] (variance EMA + velocity in + /// `isv_signal_update`). Added per ISV audit 2026-04-23 which + /// identified slot 3 as having no writer and reading constant-zero + /// into the ISV encoder MLP. + /// + /// The reducer kernel divides by its `n` argument, so we pass + /// `b * 12` for the q_var reduction to get the full-tensor mean. fn launch_loss_reduce(&self, total_loss_dev_ptr: u64) -> Result<(), MLError> { let b = self.config.batch_size as i32; + let b_qvar = (self.config.batch_size * 12) as i32; let per_sample_loss_buf_ptr = self.per_sample_loss_buf.raw_ptr(); let td_errors_buf_ptr = self.td_errors_buf.raw_ptr(); + let q_var_buf_trainer_ptr = self.q_var_buf_trainer.raw_ptr(); unsafe { self.stream .launch_builder(&self.c51_loss_reduce_kernel) @@ -12941,6 +12972,20 @@ impl GpuDqnTrainer { shared_mem_bytes: 0, }) .map_err(|e| MLError::ModelError(format!("c51_td_error_reduce launch: {e}")))?; + /* Third launch: batch mean C51 Q-distribution variance into the + * ISV scratch. Feeds slot 3 EMA + slot 4 velocity. Same generic + * reducer; n=b*12 reduces the full per-action variance tensor. */ + self.stream + .launch_builder(&self.c51_loss_reduce_kernel) + .arg(&q_var_buf_trainer_ptr) + .arg(&self.q_var_scratch_dev_ptr) + .arg(&b_qvar) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("c51_q_var_reduce launch: {e}")))?; } Ok(()) } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 487edc75e..bc16def12 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -3530,6 +3530,19 @@ impl DQNTrainer { let alpha = (0.05_f32 + 0.3 * err).min(0.5); self.training_sharpe_ema = (1.0 - alpha) * self.training_sharpe_ema + alpha * sharpe_f32; } + + // ISV audit 2026-04-23: broadcast training_sharpe_ema into ISV + // slot 22 (SHARPE_EMA_INDEX). Consumed on-GPU by + // `isv_signal_update` to drive slot 12 (learning_health) via + // sigmoid(0.1 × sharpe_ema) EMA. Replaces the prior + // component-aggregation health formula whose Sharpe-correlation + // measured at r=-0.765 on the train-mdh86 logs. + if let Some(ref fused) = self.fused_ctx { + fused.write_isv_signal_at( + crate::cuda_pipeline::gpu_dqn_trainer::SHARPE_EMA_INDEX, + self.training_sharpe_ema, + ); + } } // Limit history vectors to prevent unbounded growth