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

fix(isv): batch-aggregate regime signals + sharpe-coupled health +

Browse Source 
C51 Q-var slot + q_abs_ref outlier clamp

Bundles 4 ISV signal-quality fixes from the audit at
docs/superpowers/specs/2026-04-23-isv-signal-quality-audit.md.

1. Regime signals (slots 8-11) now batch-aggregate over all B samples
   instead of reading sample 0 only. Adds `int batch_size` kernel arg
   and fixes a latent stride bug (launch passed STATE_DIM=104 but
   states_buf has stride STATE_DIM_PADDED=128 — sample 0 worked by
   luck because both strides land at the same offset for row 0).
   99.994% information loss closed.

2. Health (slot 12) now couples to outcomes: sigmoid(0.1 × sharpe_ema)
   EMA-blended. New ISV slot 22 = SHARPE_EMA_INDEX persists the
   Rust-side training_sharpe_ema. ISV_DIM 22→23. Prior component-
   aggregation health formula was ANTI-correlated with Sharpe
   (r=-0.765 per audit) because its components saturated at 0/1
   boundaries (q_gap=1.0 / q_var=1.0 for 19/20 epochs, grad_stable
   stuck at 0.0 for 20/20 epochs). New formula's sensitive sigmoid
   region [-10, +10] Sharpe matches the observed magnitude range.
   The Rust-side write_isv_signal_at is preserved as a fallback
   initializer at epoch boundaries — the kernel then overwrites
   slot 12 every training step based on slot 22's current value.

3. Slot 3 now carries C51 Q-distribution variance (wired via third
   c51_loss_reduce launch, same pattern as td_error fix 7f92fa242).
   Previously zero-initialised with no writer. Renamed from
   "ensemble_var_scratch" to "q_var_scratch" to reflect actual
   semantics — it is the batch mean of q_var_buf_trainer (atom-
   spread variance from the C51 distributional head), NOT multi-head
   ensemble disagreement. True multi-head ensemble variance remains
   a separate follow-up if/when a per-head ensemble is wired.
   Slot 4 (velocity derivative of slot 3) becomes meaningful
   automatically.

4. q_dir_abs_ref (slot 21) and q_abs_ref (slot 16) gain outlier
   clamps before EMA update: clamped = min(raw, 10×current + 1) to
   prevent single ±10⁵ Q-excursions from poisoning 20 epochs. The
   +1 floor handles the cold-start case where current EMA is near 0.
   Slot 21 is especially load-bearing because it feeds the Kelly
   conviction denominator (q_range / q_dir_abs_ref).

ISV_DIM bump 22→23 changes the flat-param buffer size (w_isv_fc1
tensor [68] grows from [16,22] to [16,23] → +16 floats). Xavier
init at gpu_dqn_trainer.rs:13819 picks up the new dimension
automatically. Pinned allocs scale via ISV_DIM * size_of::<f32>()
expressions. No safetensors / checkpoint format currently encodes
ISV layout directly — checkpoint_state_dim/num_actions/hidden_dims
are the only hashed architectural fields — but the flat param
buffer content differs, so existing live checkpoints will require
rebase. Acceptable for the current pre-production dev state.

Tests: workspace cargo check --workspace --tests passes clean. Ran
magnitude_distribution smoke 3 times locally (RTX 3050 Ti); health
in HEALTH_DIAG now tracks negative Sharpe regime (0.50→0.38
trajectory over 20 epochs with mean sharpe_raw=-7) rather than
climbing monotonically as in the audit's train-mdh86 logs
(0.50→0.64 against Sharpe +34→-67). The magnitude_distribution
eval-dist H10 gate sometimes passes, sometimes fails depending on
the training seed — the pre-existing H10 flakiness persists. The
pre-existing 14 ml-lib test failures are unrelated (OFI features
missing / production profile config mismatch / test_batch_size
test environment issues — all fail on HEAD as well).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-23 11:34:09 +02:00
parent c5045c009e
commit 8a5e7d316c
5 changed files with 214 additions and 64 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
crates/ml/src/cuda_pipeline/c51_grad_kernel.cu
View File

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

6
crates/ml/src/cuda_pipeline/c51_loss_kernel.cu
View File

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

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

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

129
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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::<f32>(),
);
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(())
}

13
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

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