diff --git a/crates/ml/src/cuda_pipeline/batched_backward.rs b/crates/ml/src/cuda_pipeline/batched_backward.rs index d3504dff7..0551c9770 100644 --- a/crates/ml/src/cuda_pipeline/batched_backward.rs +++ b/crates/ml/src/cuda_pipeline/batched_backward.rs @@ -1952,6 +1952,42 @@ impl CublasBackwardSet { ) } + /// Compute only upstream gradient dX = dY @ W^T with a custom leading + /// dimension for X (Plan 4 Task 1B-iv). + /// + /// Mirrors `launch_dx_only` but threads `x_lda` into cuBLAS — needed for + /// the bottleneck Linear backward when its dX must be written into a + /// slice of a wider buffer (`d_vsn_gated_state[B, state_dim_padded]`'s + /// market slice `[0..market_dim]` with row stride `state_dim_padded`). + /// `beta=0.0` overwrites the destination slice; the surrounding + /// `vsn_d_gated_state_portfolio_pad_kernel` zero-initialises the market + /// slice so callers can rely on `dx` being fully defined regardless of + /// kernel/GEMM ordering. + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch_dx_only_lda( + &self, + stream: &Arc, + dy: u64, + w: u64, + dx: u64, + out_dim: usize, + in_dim: usize, + x_lda: usize, + batch: usize, + beta: f32, + ) -> Result<(), MLError> { + self.sgemm_f32( + stream, + cublas_sys::cublasOperation_t::CUBLAS_OP_N, + cublas_sys::cublasOperation_t::CUBLAS_OP_N, + in_dim as i32, batch as i32, out_dim as i32, + 1.0, w, in_dim as i32, + dy, out_dim as i32, + beta, dx, x_lda as i32, + "dX_only_lda", + ) + } + /// Launch 2-phase bias gradient reduction: db[j] += sum_b(dY[b * out_dim + j]). /// /// Phase 1: block-level shared-memory reduce into `bias_grad_partials_buf`. diff --git a/crates/ml/src/cuda_pipeline/batched_forward.rs b/crates/ml/src/cuda_pipeline/batched_forward.rs index e8583fe9f..1ff0d3587 100644 --- a/crates/ml/src/cuda_pipeline/batched_forward.rs +++ b/crates/ml/src/cuda_pipeline/batched_forward.rs @@ -278,12 +278,13 @@ pub struct CublasGemmSet { /// `vsn_mask_buf` consumed both by `vsn_mask_ema_update` (1B-i ISV kernel) /// and by 1B-iv's softmax-Jacobian backward. vsn_softmax_and_gate_fwd_kernel: CudaFunction, - /// Plan 4 Task 1B-iv consumer (loaded here in 1B-iii so the field is in - /// place; consumer wired in 1B-iv). `vsn_softmax_and_gate_backward( + /// Plan 4 Task 1B-iv consumer: `vsn_softmax_and_gate_backward( /// d_gated_state, state, mask, B, state_dim_padded, num_groups, /// group_begins, group_ends, d_logits, d_state)` implements the FULL - /// softmax-over-groups Jacobian. - #[allow(dead_code)] + /// softmax-over-groups Jacobian. Wired into the backward chain as the + /// final step of `launch_cublas_backward_to` (post-bottleneck Linear bwd + /// + portfolio passthrough assembly), writing the 24 VSN dW/dB slots at + /// grad_buf indices [95..119). vsn_softmax_and_gate_bwd_kernel: CudaFunction, /// Plan 4 Task 1B-iii: scatter handle for column-g writes of per-group /// `Linear_2_g` outputs into the assembled `vsn_logits_buf [B, num_groups]`. @@ -1486,6 +1487,204 @@ impl CublasGemmSet { Ok(()) } + /// Plan 4 Task 1B-iv: VSN feature-selection backward. + /// + /// Consumes the gradient flowing into the VSN-gated state (`d_gated_state`, + /// assembled by the trainer from the bottleneck Linear backward dX into + /// the market slice + the bottleneck-concat portfolio passthrough into + /// `[market_dim..STATE_DIM]` + zero pad into `[STATE_DIM..STATE_DIM_PADDED]`) + /// and writes per-group MLP weight + bias gradients into 24 grad_buf slots + /// at indices `[95..119)` (the layout from 1B-ii). + /// + /// Sequence (all GPU; no atomicAdd, no DtoH): + /// 1. `vsn_softmax_and_gate_backward` kernel: + /// d_gated_state, state, mask → d_logits[B, num_groups], + /// d_state[B, state_dim_padded] (discarded) + /// 2. For each group g in [0, num_groups): + /// a. `strided_gather` extracts column g of d_logits[B, num_groups] + /// into `d_logit_scratch[B]` (contiguous tight buffer). + /// b. cuBLAS Linear_2 backward (`launch_dw_only` + `launch_dx_only`): + /// dY = d_logit_scratch[B, 1] (contiguous after gather), + /// X = save_h1_g[g] [B, VSN_HIDDEN] + /// dW + dB into grad_buf[95+4g+2 / +3]; dX → d_h1_scratch[B, VSN_HIDDEN] + /// c. `relu_mask` on d_h1_scratch using save_h1_g[g] (post-ReLU activation + /// from forward — `relu_mask`'s "activation" arg is post-ReLU per kernel + /// contract). + /// d. cuBLAS Linear_1 backward via `backward_fc_layer_lda`: + /// dY = d_h1_scratch[B, VSN_HIDDEN], + /// X = state[B, group_dim_g] sliced from state[B, state_dim_padded] + /// with row stride state_dim_padded + /// dW + dB into grad_buf[95+4g+0 / +1]; no dX (`d_state` was already + /// computed by step 1, and the value is discarded since the state + /// buffer is not trainable). + /// + /// `state_ptr` and `mask_ptr` are the saved buffers from the matching online + /// `vsn_forward` call (the trainer keeps them in `states_buf` and + /// `vsn_mask_buf` respectively). `save_h1_g_ptrs` are the 6 per-group + /// post-ReLU saves the online forward populated. + /// + /// `d_logit_scratch_ptr` (tight `[B]`) and `d_h1_scratch_ptr` + /// (`[B, VSN_HIDDEN_DIM]`) are reused across the 6 groups since the loop + /// runs sequentially on the main stream. + #[allow(clippy::too_many_arguments)] + pub fn vsn_backward( + &mut self, + stream: &Arc, + cublas_backward: &super::batched_backward::CublasBackwardSet, + d_gated_state_ptr: u64, // INPUT [B, STATE_DIM_PADDED] + state_ptr: u64, // INPUT [B, STATE_DIM_PADDED] saved + mask_ptr: u64, // INPUT [B, num_groups] saved + save_h1_g_ptrs: &[u64; ml_core::state_layout::SL_NUM_FEATURE_GROUPS], + w_ptrs: &[u64; super::gpu_dqn_trainer::NUM_WEIGHT_TENSORS], + grad_buf_base: u64, + param_sizes: &[usize], + d_logits_ptr: u64, // SCRATCH [B, num_groups] + d_state_ptr: u64, // SCRATCH [B, STATE_DIM_PADDED] (kernel writes; discarded) + d_logit_scratch_ptr: u64, // SCRATCH [B] (per-group, reused) + d_h1_scratch_ptr: u64, // SCRATCH [B, VSN_HIDDEN_DIM] (per-group, reused) + group_begins_dev: u64, + group_ends_dev: u64, + gather_kernel: &CudaFunction, // strided_gather (column extraction) + batch_size: usize, + ) -> Result<(), MLError> { + use super::gpu_dqn_trainer::{padded_byte_offset, VSN_HIDDEN_DIM}; + const VSN_TENSORS_BASE: usize = 95; + const NUM_GROUPS: usize = ml_core::state_layout::SL_NUM_FEATURE_GROUPS; + + let group_ranges = ml_core::state_layout::FEATURE_GROUP_RANGES; + let state_dim_padded = self.state_dim_padded; + let f32_size = std::mem::size_of::() as u64; + + // ── Step 1: softmax-and-gate backward ──────────────────────────────── + // Writes d_logits[B, num_groups] (consumed by per-group Linear_2 bwd + // via gather) and d_state[B, state_dim_padded] (discarded — state buf + // not trainable). Block size 256, grid size B; shared mem layout + // (per kernel comment): 2*num_groups + block floats. + { + let b_i32 = batch_size as i32; + let state_dim_padded_i32 = state_dim_padded as i32; + let num_groups_i32 = NUM_GROUPS as i32; + let smem_floats = 2 * NUM_GROUPS + 256; + let smem_bytes = (smem_floats * std::mem::size_of::()) as u32; + unsafe { + stream + .launch_builder(&self.vsn_softmax_and_gate_bwd_kernel) + .arg(&d_gated_state_ptr) + .arg(&state_ptr) + .arg(&mask_ptr) + .arg(&b_i32) + .arg(&state_dim_padded_i32) + .arg(&num_groups_i32) + .arg(&group_begins_dev) + .arg(&group_ends_dev) + .arg(&d_logits_ptr) + .arg(&d_state_ptr) + .launch(LaunchConfig { + grid_dim: (batch_size as u32, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: smem_bytes, + }) + .map_err(|e| MLError::ModelError(format!( + "vsn_softmax_and_gate_backward: {e}" + )))?; + } + } + + // ── Step 2: per-group Linear_2 backward + ReLU mask + Linear_1 bwd ─── + for g in 0..NUM_GROUPS { + let (gb, ge) = group_ranges[g]; + let group_dim = ge - gb; + debug_assert!( + group_dim > 0, + "vsn_backward: feature group {g} has zero dim — group ranges invariant violated" + ); + + // (a) Gather column g of d_logits[B, num_groups] → d_logit_scratch[B]. + // Column-extraction via strided_gather: dst_stride=1, col=g, + // src_lda=num_groups, total = B * dst_stride = B. + let src_lda_i32 = NUM_GROUPS as i32; + let col_i32 = g as i32; + let dst_stride_i32: i32 = 1; + let total_i32 = batch_size as i32; + let blocks = ((batch_size as u32 + 255) / 256).max(1); + unsafe { + stream + .launch_builder(gather_kernel) + .arg(&d_logits_ptr) + .arg(&d_logit_scratch_ptr) + .arg(&src_lda_i32) + .arg(&col_i32) + .arg(&dst_stride_i32) + .arg(&total_i32) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "vsn_backward strided_gather group {g}: {e}" + )))?; + } + + // (b) Linear_2 backward: dY = d_logit_scratch[B, 1] (tight), + // X = save_h1_g[g] [B, VSN_HIDDEN] (tight), + // dW [1, VSN_HIDDEN] into grad_buf[95+4g+2], + // dB [1] into grad_buf[95+4g+3], + // dX [B, VSN_HIDDEN] into d_h1_scratch (overwrite, beta=0). + let goff_w2 = padded_byte_offset(param_sizes, VSN_TENSORS_BASE + 4 * g + 2); + let goff_b2 = padded_byte_offset(param_sizes, VSN_TENSORS_BASE + 4 * g + 3); + cublas_backward.launch_dw_only( + stream, + d_logit_scratch_ptr, // dY [B, 1] + save_h1_g_ptrs[g], // X [B, VSN_HIDDEN] + grad_buf_base + goff_w2, // dW [1, VSN_HIDDEN] + grad_buf_base + goff_b2, // dB [1] + 1, VSN_HIDDEN_DIM, batch_size, + )?; + cublas_backward.launch_dx_only( + stream, + d_logit_scratch_ptr, // dY [B, 1] + w_ptrs[VSN_TENSORS_BASE + 4 * g + 2], // W2 [1, VSN_HIDDEN] + d_h1_scratch_ptr, // dX [B, VSN_HIDDEN] + 1, VSN_HIDDEN_DIM, batch_size, 0.0_f32, + )?; + + // (c) ReLU mask on d_h1_scratch using saved post-ReLU h1. + cublas_backward.relu_mask( + stream, + d_h1_scratch_ptr, + save_h1_g_ptrs[g], + batch_size * VSN_HIDDEN_DIM, + )?; + + // (d) Linear_1 backward (state slice with x_lda = state_dim_padded): + // dY = d_h1_scratch [B, VSN_HIDDEN] (tight), + // X = state + gb*sizeof(f32) (group-slice column offset), + // x_lda = state_dim_padded (row stride of full state row), + // dW [VSN_HIDDEN, group_dim_g] into grad_buf[95+4g+0], + // dB [VSN_HIDDEN] into grad_buf[95+4g+1], + // no dX (d_state was computed in step 1; state buf untrainable). + let goff_w1 = padded_byte_offset(param_sizes, VSN_TENSORS_BASE + 4 * g + 0); + let goff_b1 = padded_byte_offset(param_sizes, VSN_TENSORS_BASE + 4 * g + 1); + let state_slice_ptr = state_ptr + (gb as u64) * f32_size; + cublas_backward.backward_fc_layer_lda( + stream, + d_h1_scratch_ptr, // dY [B, VSN_HIDDEN] + state_slice_ptr, // X [B, group_dim] @ stride=state_dim_padded + w_ptrs[VSN_TENSORS_BASE + 4 * g + 0], // W1 [VSN_HIDDEN, group_dim] (unused for dX-skip path) + grad_buf_base + goff_w1, // dW [VSN_HIDDEN, group_dim] + grad_buf_base + goff_b1, // dB [VSN_HIDDEN] + 0, // dX skipped (state buf untrainable) + VSN_HIDDEN_DIM, + group_dim, + state_dim_padded, // x_lda + batch_size, + )?; + } + + Ok(()) + } + /// Run a single advantage branch's decoder (FC head + adv-logits) on the /// caller-supplied stream, sequentially. /// diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu index 17978adae..c5ffe606e 100644 --- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu +++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu @@ -1186,6 +1186,63 @@ extern "C" __global__ void bn_tanh_backward_kernel( d_bn[idx] = d_upstream * (1.0f - tanh_val * tanh_val); } +/* ══════════════════════════════════════════════════════════════════════ + * VSN-AWARE BOTTLENECK CONCAT BACKWARD: PORTFOLIO PASSTHROUGH (Plan 4 Task 1B-iv) + * + * Inverse of the portfolio-passthrough path inside `bn_tanh_concat_kernel`: + * forward (concat): concat_out[b, bn_dim + k] = states[b, market_dim + k] + * for k in [0, portfolio_dim) + * backward (here): d_gated_state[b, market_dim + k] = d_concat[b, bn_dim + k] + * + * The market slice of `d_gated_state[b, 0..market_dim]` is filled by the + * cuBLAS bottleneck Linear backward (`launch_dx_only_lda`, with row stride + * `state_dim_padded`) AFTER this kernel runs. Padding region + * `[state_dim..state_dim_padded)` is zero-filled here (no-op when + * STATE_DIM == STATE_DIM_PADDED, but we honor the contract for callers + * that pad the state). + * + * One thread per (batch, padded-state-index) pair handles either the + * portfolio slice or the padding tail; the market slice is left untouched + * (cuBLAS overwrites it with beta=0). + * + * Grid: ceil(B*state_dim_padded/256), Block: 256. + * No atomicAdd, deterministic. + * ══════════════════════════════════════════════════════════════════════ */ +extern "C" __global__ void vsn_d_gated_state_portfolio_pad_kernel( + float* __restrict__ d_gated_state, /* [B, state_dim_padded] f32 OUT */ + const float* __restrict__ d_concat, /* [B, concat_dim] f32 IN */ + int batch_size, + int market_dim, + int bn_dim, + int concat_dim, + int state_dim, + int state_dim_padded +) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + int total = batch_size * state_dim_padded; + if (idx >= total) return; + + int b = idx / state_dim_padded; + int i = idx % state_dim_padded; + + if (i < market_dim) { + /* Market slice: zero-initialised here so the buffer is fully defined + * even before cuBLAS launch_dx_only_lda overwrites it (beta=0 GEMM) + * with the bottleneck Linear backward dX. Removes a write-order + * dependency between this kernel and the cuBLAS launch. */ + d_gated_state[idx] = 0.0f; + } else if (i < state_dim) { + /* Portfolio passthrough: d_gated_state[b, market_dim + k] + * = d_concat [b, bn_dim + k] + * where k = i - market_dim (∈ [0, portfolio_dim)). */ + int k = i - market_dim; + d_gated_state[idx] = d_concat[(long long)b * concat_dim + bn_dim + k]; + } else { + /* Padding tail [state_dim, state_dim_padded): zero. */ + d_gated_state[idx] = 0.0f; + } +} + /* ══════════════════════════════════════════════════════════════════════ * BOTTLENECK WEIGHT GRADIENT BIAS REDUCE (#31) * diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 864293779..7370308b4 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -3798,6 +3798,36 @@ extern "C" __global__ void strided_scatter( dst[b * dst_stride + i] = src[idx]; } +/** + * Plan 4 Task 1B-iv: gather a single column from a wide buffer into a tight one. + * + * Inverse of `strided_scatter` for the column-extraction case (src_stride=1): + * dst[b * dst_stride + i] = src[b * src_lda + col + i] for i in [0, dst_stride) + * + * Used by `vsn_backward` to extract column `g` of `d_logits[B, num_groups]` + * into a contiguous `[B, dst_stride]` buffer, so the per-group cuBLAS + * `Linear_2` backward (which expects a tightly-packed `dY[B, out_dim]`) can + * read it without a strided-leading-dim cuBLAS variant. Per-group `out_dim=1` + * for VSN's scalar-per-group logit, but the kernel is generic in dst_stride. + * + * Grid: ceil(total/256), Block: 256. total = B * dst_stride. No atomicAdd. + * One thread per output element, deterministic. + */ +extern "C" __global__ void strided_gather( + const float* __restrict__ src, /* [B, src_lda] wide */ + float* __restrict__ dst, /* [B, dst_stride] tight */ + int src_lda, /* row stride of src */ + int col, /* column offset within each src row */ + int dst_stride, /* number of contiguous columns to read */ + int total /* B * dst_stride */ +) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= total) return; + int b = idx / dst_stride; + int i = idx % dst_stride; + dst[idx] = src[b * src_lda + col + i]; +} + /* ================================================================== */ /* Kernel: variable_select_bottleneck */ /* ================================================================== */ diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index ae9cecb68..f7ef6b4aa 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -162,6 +162,38 @@ pub(crate) static VSN_MASK_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_D /// `VSN_HIDDEN_DIM*group_dim_g + VSN_HIDDEN_DIM + VSN_HIDDEN_DIM + 1`. pub const VSN_HIDDEN_DIM: usize = 16; +/// Plan 4 Task 1B-iv-rc: VSN dW attenuation factor — global multiplier applied +/// to `vsn_d_gated_state_buf` IN-PLACE before each `vsn_backward` invocation +/// (one of 4 trunk-touching paths: main C51/MSE, IQN trunk, CQL, ensemble +/// diversity). The softmax-and-gate Jacobian is linear in d_gated_state, so +/// scaling the input scales every VSN dW/dB downstream uniformly across the +/// 24 VSN tensors at grad_buf[95..119). +/// +/// Rationale (smoke regression analysis, 2026-04-25): +/// - 1B-iii baseline (VSN frozen at Xavier, no backward): geom-mean 71.24 +/// - 1B-iv (VSN trains via main C51/MSE path only): geom-mean 57.49 +/// - 1B-iv-ext (VSN trains via all 4 paths): geom-mean 36.36 +/// +/// More VSN gradient signal → worse smoke. Pattern is monotonic in the +/// "amount of trainable VSN signal", not in the "asymmetry" of coverage — +/// disproves the H3 (asymmetric coverage) hypothesis. Consistent with H1 +/// (VSN-bottleneck coupling instability): VSN gates the bottleneck Linear's +/// input distribution; a strong trainable VSN gate shifts that distribution +/// faster than the bottleneck Linear can co-adapt, leading to a regime where +/// the GRN trunk receives an increasingly off-distribution embedding. +/// +/// `VSN_DW_DAMP = 0.10` attenuates VSN gradient signal by 10× across all 4 +/// paths uniformly. Constant (not a per-step warmup ramp) because per-step +/// variation would require pinned-device-mapped scalar plumbing and a kernel +/// variant that reads `*alpha_ptr` instead of the baked `alpha` literal — +/// out of scope for the rc commit. The constant is graph-capture-stable. +/// +/// If smoke ≥ 71.24 with this dampening, we'll have established that the +/// regression IS gradient-magnitude driven and a follow-up commit can +/// elevate it to a per-step ramp tied to ISV[VSN_MASK_GROUP_*_EMA] drift +/// or to `c51_alpha`. +pub const VSN_DW_DAMP: f32 = 0.10; + /// Mamba2 temporal scan configuration. const MAMBA2_HISTORY_K: usize = 8; // Rolling history length const MAMBA2_STATE_DIM: usize = 16; // SSM state dimension @@ -349,8 +381,10 @@ const ISV_NETWORK_DIM: usize = 23; /// [108] VSN_MASK_GROUP_3_EMA_INDEX — `mtf` group, same producer/contract. /// [109] VSN_MASK_GROUP_4_EMA_INDEX — `portfolio` group, same producer/contract. /// [110] VSN_MASK_GROUP_5_EMA_INDEX — `plan_isv` group, same producer/contract. -/// [111] ISV_LAYOUT_FINGERPRINT_LO_INDEX — low 32 bits of layout fingerprint (shifted from 103). -/// [112] ISV_LAYOUT_FINGERPRINT_HI_INDEX — high 32 bits of layout fingerprint (shifted from 104). +/// [111] ISV_LAYOUT_FINGERPRINT_LO_INDEX — low 32 bits of layout fingerprint +/// (shifted 103→111 in 1B-ii). +/// [112] ISV_LAYOUT_FINGERPRINT_HI_INDEX — high 32 bits of layout fingerprint +/// (shifted 104→112 in 1B-ii). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. const ISV_TOTAL_DIM: usize = 113; @@ -686,14 +720,14 @@ pub const VSN_MASK_GROUP_5_EMA_INDEX: usize = 110; /// 85→90 in Plan 4 Task 5 Mode A, 90→94 in Plan 4 follow-up (target-drift /// EMA replacing legacy CPU-DtoH path), 94→97 in Plan 4 Task 2c.3c.5 /// (H_S2_RMS_EMA producer slot), 97→103 in Plan 4 Task 3 E.3 (IQN -/// fixed-τ multi-quantile head, +4 diagnostic slots), then 103→111 in +/// fixed-τ multi-quantile head, +4 diagnostic slots), 103→111 in /// Plan 4 Task 1B-ii (VSN per-group mask EMA, +6 slots). pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 111; /// ISV slot [112] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits). /// Shifted 62→70 in Plan 3 Task 1, 70→74 in Plan 3 Task 3, 74→77 in Plan 3 Task 4 B.4, /// 77→81 in Plan 3 Task 7 C.3, 81→86 in Plan 3 Task 8 B.3, 86→91 in Plan 4 Task 5, /// 91→95 in Plan 4 follow-up, 95→98 in Plan 4 Task 2c.3c.5, 98→104 in Plan 4 -/// Task 3 E.3, then 104→112 in Plan 4 Task 1B-ii. +/// Task 3 E.3, 104→112 in Plan 4 Task 1B-ii. pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 112; /// Canonical alias for the fingerprint slot (the low half). pub const ISV_LAYOUT_FINGERPRINT_INDEX: usize = ISV_LAYOUT_FINGERPRINT_LO_INDEX; @@ -1547,6 +1581,15 @@ pub struct GpuDqnTrainer { saxpy_f32_adam_grad: CudaFunction, /// Separate saxpy_f32 instance for aux_child — same Hopper CUfunction isolation. saxpy_f32_aux: CudaFunction, + /// Plan 4 Task 1B-iv-rc: separate `dqn_scale_f32_kernel` instance for + /// `aux_child` — same Hopper CUfunction-per-child isolation as + /// `saxpy_f32_aux`. Used to attenuate `vsn_d_gated_state_buf` by a + /// constant `VSN_DW_DAMP` factor before each aux-path `vsn_backward` + /// invocation, mirroring the main-path attenuation that uses + /// `scale_f32_kernel` in `forward_child`. The constant is baked at + /// graph capture time — no per-step variation, so no pinned-device + /// scalar plumbing required. + scale_f32_aux: CudaFunction, /// D1/N1: Distillation fused SAXPY kernel (loaded from aux_module for /// aux_child graph capture). Reads health from ISV pinned memory, /// computes alpha internally, performs `grad += alpha * (params − best)` @@ -1959,6 +2002,21 @@ pub struct GpuDqnTrainer { target_params_buf: CudaSlice, // [TOTAL_PARAMS] f32 master target parameters (EMA operates here) m_buf: CudaSlice, // [TOTAL_PARAMS] Adam first moment (f32 for precision) v_buf: CudaSlice, // [TOTAL_PARAMS] Adam second moment (f32 for precision) + /// Plan 4 Task 1B-iv (chain final fix): cached f32 element count of the + /// VSN tensor slice `[95..119)` (= `(padded_byte_offset(119) − + /// padded_byte_offset(95)) / size_of::()`). Used by the second + /// `target_ema_update` Polyak launch that syncs target VSN params with + /// online — without that second launch, target VSN stays at + /// `alloc_zeros` while online VSN drifts toward measured per-group + /// importances. Cached at construction so the per-epoch launch avoids + /// recomputing `compute_param_sizes` + `padded_byte_offset`. + vsn_param_total: usize, + /// Plan 4 Task 1B-iv (chain final fix): cached byte offset from + /// `params_buf.raw_ptr()` (and `target_params_buf.raw_ptr()`) to the + /// start of the VSN tensor slice (= `padded_byte_offset(95)`). Used + /// alongside `vsn_param_total` by the VSN-range Polyak EMA launch in + /// `target_ema_update`. + vsn_param_byte_offset: u64, /// Gradient L2 norm [1] — pinned device-mapped. Finalize kernel writes via dev_ptr, CPU reads via host ptr. grad_norm_pinned: *mut f32, grad_norm_dev_ptr: u64, @@ -2245,6 +2303,71 @@ pub struct GpuDqnTrainer { /// and all 3 forward passes. vsn_group_begins_buf: CudaSlice, vsn_group_ends_buf: CudaSlice, + /// Plan 4 Task 1B-iv: VSN backward `d_logits` scratch `[B, num_groups]`. + /// Output of `vsn_softmax_and_gate_backward`; consumed by per-group + /// `Linear_2` cuBLAS backward via `strided_gather` + a tight per-group + /// scratch (no strided cuBLAS variant required). + vsn_d_logits_buf: CudaSlice, + /// Plan 4 Task 1B-iv: VSN backward `d_state` scratch + /// `[B, STATE_DIM_PADDED]`. Output of `vsn_softmax_and_gate_backward` + /// (per-feature softmax-Jacobian + gate term); not propagated downstream + /// because `state_buf` is not trainable. Allocated so the kernel has a + /// valid write target; the value is then unused. + vsn_d_state_buf: CudaSlice, + /// Plan 4 Task 1B-iv: assembled VSN backward INPUT + /// `d_vsn_gated_state[B, STATE_DIM_PADDED]`. Built from (a) cuBLAS + /// bottleneck Linear backward dX into the market slice `[0..market_dim]` + /// (row stride = `state_dim_padded` via `launch_dx_only_lda`) and + /// (b) `vsn_d_gated_state_portfolio_pad_kernel` for the portfolio slice + /// `[market_dim..STATE_DIM]` + zero pad `[STATE_DIM..STATE_DIM_PADDED]`. + /// Consumed by `vsn_backward` as `d_gated_state` argument. + vsn_d_gated_state_buf: CudaSlice, + /// Plan 4 Task 1B-iv: per-group `d_logit_g` scratch `[B]`. The VSN + /// backward extracts column `g` of `vsn_d_logits_buf [B, num_groups]` + /// into this contiguous buffer (via `strided_gather`) so the per-group + /// `Linear_2[g]` cuBLAS backward sees a tightly-packed `dY[B, 1]`. One + /// shared buffer suffices because the per-group loop runs sequentially. + vsn_d_logit_scratch_buf: CudaSlice, + /// Plan 4 Task 1B-iv: per-group `d_h1_g` scratch `[B, VSN_HIDDEN_DIM=16]`. + /// Output of per-group `Linear_2[g]` backward dX, then ReLU-mask + /// in-place against the saved `vsn_save_h1_g{g}_buf`, then consumed as + /// `dY` of per-group `Linear_1[g]` backward. One shared buffer because + /// the per-group loop is sequential. + vsn_d_h1_scratch_buf: CudaSlice, + /// Plan 4 Task 1B-iv: gather kernel handle (`strided_gather` from + /// `experience_kernels.cu`) used by `vsn_backward` to extract per-group + /// `d_logit_g` columns from the assembled `d_logits[B, num_groups]`. + /// Loaded from `EXPECTED_Q_CUBIN` next to `strided_scatter` in the + /// constructor. + vsn_logit_gather_kernel: CudaFunction, + /// Plan 4 Task 1B-iv: portfolio-passthrough + pad kernel handle + /// (`vsn_d_gated_state_portfolio_pad_kernel` from `dqn_utility_kernels.cu`) + /// used by `launch_cublas_backward_to` to populate the portfolio slice + + /// zero pad of `vsn_d_gated_state_buf` from `bn_d_concat_buf`. The + /// market slice is overwritten by `launch_dx_only_lda` (bottleneck Linear + /// backward) afterwards. + vsn_d_gated_state_portfolio_kernel: CudaFunction, + /// Plan 4 Task 1B-iv-ext: VSN dW scratch for the IQN-trunk auxiliary + /// backward path (`apply_iqn_trunk_gradient`). Sized to exactly the + /// padded VSN range [95..119) (24 tensors). The aux path runs the same + /// VSN backward chain as the main path and writes per-group dW/dB into + /// this scratch, then the trailing SAXPY (`saxpy_f32_aux`) accumulates + /// `iqn_lambda * iqn_readiness * iqn_budget * scratch` into + /// `grad_buf` over the same 24-tensor range. Independent buffer (not + /// shared with `iqn_trunk_m`) because `iqn_trunk_m` is sized only to + /// the 13 trunk tensors at indices [0..13) — extending it would leave + /// 82 dead tensors of zeroed scratch between trunk and VSN, which the + /// SAXPY can't skip without becoming non-contiguous. + vsn_dw_iqn_aux_scratch: CudaSlice, + /// Plan 4 Task 1B-iv-ext: VSN dW scratch for the ensemble-diversity + /// auxiliary backward path (`apply_ensemble_diversity_backward`). Same + /// shape and rationale as `vsn_dw_iqn_aux_scratch`; the trailing SAXPY + /// scales by the caller-supplied diversity weight. The CQL aux path + /// reuses `cql_grad_scratch` (which is sized `total_params` and SAXPY- + /// covered end-to-end by `apply_cql_saxpy`) so it does not need a + /// per-aux VSN scratch — VSN dW writes there land at offsets [95..119) + /// and ride the existing CQL SAXPY into `grad_buf`. + vsn_dw_ensemble_aux_scratch: CudaSlice, /// Plan 4 Task 1B-i / 1B-iii: VSN per-group mask EMA producer kernel. /// Single-block 256-thread shmem-reduction kernel reads `vsn_mask_buf` /// (saved by the online-on-states VSN forward) and EMA-updates @@ -2970,9 +3093,11 @@ impl GpuDqnTrainer { /// Scale Adam momentum buffers (m_buf) by `scale` — used for warm restart momentum decay. /// Preserves gradient direction but reduces magnitude to prevent Adam fixed point. pub(crate) fn scale_adam_momentum(&mut self, scale: f32) -> Result<(), MLError> { + let cfg = { + let blocks = ((self.total_params + 255) / 256) as u32; + LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), shared_mem_bytes: 0 } + }; let n = self.total_params as i32; - let blocks = ((self.total_params + 255) / 256) as u32; - let cfg = LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), shared_mem_bytes: 0 }; let m_ptr = self.m_buf.raw_ptr(); // Use scale_f32_ungraphed — called at cosine LR warm restarts (after graph capture). // scale_f32_kernel is captured in forward_child; launching from the same CUmodule @@ -6009,6 +6134,18 @@ impl GpuDqnTrainer { let grn_h_s2_d_elu_out_p = self.bw_grn_h_s2_d_elu_out.raw_ptr(); let grn_h_s2_d_linear_a_out_p= self.bw_grn_h_s2_d_linear_a_out.raw_ptr(); + // Plan 4 Task 1B-iv-ext: pass non-zero `s1_dx_output` so the chain + // produces dL/d(bn_concat) into `bn_d_concat_buf` for the VSN + // backward extension below. Reuses the existing scratch buffer + // because the main backward path already consumed it (aux paths run + // after main per `fused_training.rs` ordering); `encoder_backward_chain` + // beta=0-overwrites the Linear_a dX path and beta=1-accumulates the + // Linear_residual dX path into this slot. + let bn_d_concat_for_iqn = if self.config.bottleneck_dim > 0 { + self.bn_d_concat_buf.raw_ptr() + } else { + 0u64 + }; self.cublas_forward.encoder_backward_chain( &self.stream, &self.cublas_backward, @@ -6018,7 +6155,7 @@ impl GpuDqnTrainer { ¶m_sizes, d_h_s2, d_h_s1, - 0, // IQN: no bottleneck dX needed (separate gradient budget) + bn_d_concat_for_iqn, // 1B-iv-ext: bottleneck dX needed for VSN aux backward grn_h_s1_d_pre_ln_p, grn_h_s1_d_linear_b_out_p, grn_h_s1_d_elu_out_p, @@ -6030,6 +6167,65 @@ impl GpuDqnTrainer { h_s1_save_ptr, )?; + // ── 3b. Plan 4 Task 1B-iv-ext: VSN backward chain ───────────────── + // Consumes the bn_d_concat just produced by encoder_backward_chain, + // assembles `d_vsn_gated_state`, and writes 24 VSN dW/dB into + // `vsn_dw_iqn_aux_scratch` at goff[95..119) — the scratch base is + // anchored so `padded_byte_offset(95)` lands at scratch byte 0. + // The trailing SAXPY (step 5) then mixes this into grad_buf[95..119) + // with the same `iqn_lambda * iqn_readiness * iqn_budget` scale that + // the trunk SAXPY uses. Closes the gradient-coverage gap from 1B-iv. + if self.config.bottleneck_dim > 0 { + let vsn_begin_off = padded_byte_offset(¶m_sizes, 95); + let vsn_aux_scratch_anchor = + self.vsn_dw_iqn_aux_scratch.raw_ptr().wrapping_sub(vsn_begin_off); + // Zero the scratch before write (vsn_backward writes dW with + // beta=1 inside cuBLAS dW GEMMs). One memset per call — + // vsn_aux_scratch_floats f32s. + let vsn_scratch_bytes = self.vsn_dw_iqn_aux_scratch.len() * std::mem::size_of::(); + unsafe { + cudarc::driver::result::memset_d8_async( + self.vsn_dw_iqn_aux_scratch.raw_ptr(), + 0u8, + vsn_scratch_bytes, + self.stream.cu_stream(), + ).map_err(|e| MLError::ModelError(format!( + "memset vsn_dw_iqn_aux_scratch: {e}" + )))?; + } + Self::aux_bottleneck_vsn_backward_dispatch( + &mut self.cublas_forward, + &self.cublas_backward, + &self.stream, + &self.config, + &self.bn_d_hidden_buf, + &self.bn_hidden_buf, + &self.states_buf, + &self.vsn_d_gated_state_buf, + &self.vsn_mask_buf, + &self.vsn_d_logits_buf, + &self.vsn_d_state_buf, + &self.vsn_d_logit_scratch_buf, + &self.vsn_d_h1_scratch_buf, + &self.vsn_group_begins_buf, + &self.vsn_group_ends_buf, + &self.vsn_save_h1_g0_buf, + &self.vsn_save_h1_g1_buf, + &self.vsn_save_h1_g2_buf, + &self.vsn_save_h1_g3_buf, + &self.vsn_save_h1_g4_buf, + &self.vsn_save_h1_g5_buf, + &self.bn_tanh_backward_kernel, + &self.vsn_d_gated_state_portfolio_kernel, + &self.vsn_logit_gather_kernel, + &self.scale_f32_aux, // 1B-iv-rc: VSN dW dilution (VSN_DW_DAMP) + bn_d_concat_for_iqn, + vsn_aux_scratch_anchor, + ¶m_sizes, + &w_ptrs, + )?; + } + // ── 4. Plain SAXPY: grad_buf[trunk] += iqn_lambda * iqn_budget * scratch ── // Adaptive lambda: scales with IQN loss readiness (0 when uncertain, // 1 when converged). B4/G5: further scaled by iqn_budget. @@ -6057,6 +6253,37 @@ impl GpuDqnTrainer { } } + // ── 5. Plan 4 Task 1B-iv-ext: VSN-range SAXPY ──────────────────── + // grad_buf[VSN range] += iqn_lambda * iqn_readiness * iqn_budget * + // vsn_dw_iqn_aux_scratch + // Same scaling factor as the trunk SAXPY above — IQN's adaptive + // lambda gates both contributions in lockstep so the per-component + // budget remains 60% IQN regardless of how the gradient distributes + // across trunk vs VSN tensors. + if self.config.bottleneck_dim > 0 { + let vsn_begin_off = padded_byte_offset(¶m_sizes, 95); + let vsn_count = self.vsn_dw_iqn_aux_scratch.len(); + let scratch_ptr = self.vsn_dw_iqn_aux_scratch.raw_ptr(); + let grad_vsn_ptr = self.ptrs.grad_buf + vsn_begin_off; + let scale = self.config.iqn_lambda * self.iqn_readiness * iqn_budget; + let n_i32 = vsn_count as i32; + let blocks = ((vsn_count + 255) / 256) as u32; + unsafe { + self.stream + .launch_builder(&self.saxpy_f32_aux) + .arg(&grad_vsn_ptr) + .arg(&scratch_ptr) + .arg(&scale) + .arg(&n_i32) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("IQN VSN SAXPY: {e}")))?; + } + } + Ok(()) } @@ -6154,13 +6381,22 @@ impl GpuDqnTrainer { // ── 5. GRN trunk backward (encoder_backward_chain) ── // h_s2's gradient flows from d_h_s2 (post-LN) into the trunk via the - // GRN composition (no ReLU truncation). Bottleneck dX is unused for - // ensemble (separate gradient budget) so s1_dx_output = 0. + // GRN composition (no ReLU truncation). + // Plan 4 Task 1B-iv-ext: pass non-zero `s1_dx_output` so the chain + // produces dL/d(bn_concat) for the VSN backward chain step 5b + // below. Reuses `bn_d_concat_buf` because the IQN aux path runs + // before ensemble per `fused_training.rs` ordering and its consumer + // (the IQN VSN backward) has already finished writing. let states_ptr = if self.config.bottleneck_dim > 0 { self.bn_concat_buf.raw_ptr() } else { bw_raw_f32_ptr(&self.states_buf, &self.stream) }; + let bn_d_concat_for_ens = if self.config.bottleneck_dim > 0 { + self.bn_d_concat_buf.raw_ptr() + } else { + 0u64 + }; let h_s1_save_ptr = bw_raw_f32_ptr(&self.save_h_s1, &self.stream); let d_h_s2 = self.ptrs.bw_d_h_s2; let d_h_s1 = bw_raw_f32_ptr(&self.bw_d_h_s1, &self.stream); @@ -6184,7 +6420,7 @@ impl GpuDqnTrainer { ¶m_sizes, d_h_s2, d_h_s1, - 0, + bn_d_concat_for_ens, // 1B-iv-ext: bottleneck dX needed for VSN aux backward grn_h_s1_d_pre_ln_p, grn_h_s1_d_linear_b_out_p, grn_h_s1_d_elu_out_p, @@ -6196,6 +6432,58 @@ impl GpuDqnTrainer { h_s1_save_ptr, )?; + // ── 5b. Plan 4 Task 1B-iv-ext: VSN backward chain ───────────────── + // Mirrors the IQN aux path's VSN extension. Writes 24 dW/dB into + // `vsn_dw_ensemble_aux_scratch`, then SAXPY's into grad_buf[VSN] + // with the caller-supplied diversity weight (step 7). + if self.config.bottleneck_dim > 0 { + let vsn_begin_off = padded_byte_offset(¶m_sizes, 95); + let vsn_aux_scratch_anchor = + self.vsn_dw_ensemble_aux_scratch.raw_ptr().wrapping_sub(vsn_begin_off); + let vsn_scratch_bytes = self.vsn_dw_ensemble_aux_scratch.len() * std::mem::size_of::(); + unsafe { + cudarc::driver::result::memset_d8_async( + self.vsn_dw_ensemble_aux_scratch.raw_ptr(), + 0u8, + vsn_scratch_bytes, + self.stream.cu_stream(), + ).map_err(|e| MLError::ModelError(format!( + "memset vsn_dw_ensemble_aux_scratch: {e}" + )))?; + } + Self::aux_bottleneck_vsn_backward_dispatch( + &mut self.cublas_forward, + &self.cublas_backward, + &self.stream, + &self.config, + &self.bn_d_hidden_buf, + &self.bn_hidden_buf, + &self.states_buf, + &self.vsn_d_gated_state_buf, + &self.vsn_mask_buf, + &self.vsn_d_logits_buf, + &self.vsn_d_state_buf, + &self.vsn_d_logit_scratch_buf, + &self.vsn_d_h1_scratch_buf, + &self.vsn_group_begins_buf, + &self.vsn_group_ends_buf, + &self.vsn_save_h1_g0_buf, + &self.vsn_save_h1_g1_buf, + &self.vsn_save_h1_g2_buf, + &self.vsn_save_h1_g3_buf, + &self.vsn_save_h1_g4_buf, + &self.vsn_save_h1_g5_buf, + &self.bn_tanh_backward_kernel, + &self.vsn_d_gated_state_portfolio_kernel, + &self.vsn_logit_gather_kernel, + &self.scale_f32_aux, // 1B-iv-rc: VSN dW dilution (VSN_DW_DAMP) + bn_d_concat_for_ens, + vsn_aux_scratch_anchor, + ¶m_sizes, + &w_ptrs, + )?; + } + // ── 6. Plain SAXPY: grad_buf[trunk] += scale * scratch ── // No per-component clip — single global clip in Adam handles safety. // Uses aux_child-specific handle — saxpy_f32_kernel is captured in @@ -6221,6 +6509,34 @@ impl GpuDqnTrainer { } } + // ── 7. Plan 4 Task 1B-iv-ext: VSN-range SAXPY ──────────────────── + // grad_buf[VSN range] += scale * vsn_dw_ensemble_aux_scratch + // Same diversity weight as the trunk SAXPY above — closes the + // gradient-coverage gap so VSN params learn from the ensemble KL + // signal at the same effective amplitude as trunk weights do. + if self.config.bottleneck_dim > 0 { + let vsn_begin_off = padded_byte_offset(¶m_sizes, 95); + let vsn_count = self.vsn_dw_ensemble_aux_scratch.len(); + let scratch_ptr = self.vsn_dw_ensemble_aux_scratch.raw_ptr(); + let grad_vsn_ptr = self.ptrs.grad_buf + vsn_begin_off; + let n_i32 = vsn_count as i32; + let blocks = ((vsn_count + 255) / 256) as u32; + unsafe { + self.stream + .launch_builder(&self.saxpy_f32_aux) + .arg(&grad_vsn_ptr) + .arg(&scratch_ptr) + .arg(&scale) + .arg(&n_i32) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("ensemble VSN SAXPY: {e}")))?; + } + } + Ok(()) } @@ -6678,11 +6994,12 @@ impl GpuDqnTrainer { // ── GRN trunk backward (Plan 4 Task 2c.3c.4, CQL path) ── // Mirrors the main backward path but writes trunk gradients into - // `cql_grad_scratch` instead of `grad_buf`. CQL never asks for the - // trunk-input gradient (`s1_dx_output = 0`), so the bottleneck dX - // GEMM short-circuits and `cql_grad_scratch` accumulates the GRN - // trunk dW/dB/dgamma/dbeta only — exactly what `apply_cql_saxpy` - // expects to mix into `grad_buf` afterwards. + // `cql_grad_scratch` instead of `grad_buf`. + // Plan 4 Task 1B-iv-ext: pass non-zero `s1_dx_output` so the chain + // produces dL/d(bn_concat) for the VSN backward extension below. + // The CQL aux path runs after the IQN aux path's VSN backward has + // finished (per `fused_training.rs` ordering), so reusing + // `bn_d_concat_buf` is safe. let grn_h_s1_d_pre_ln_cql = self.bw_grn_h_s1_d_pre_ln.raw_ptr(); let grn_h_s1_d_linear_b_out_cql= self.bw_grn_h_s1_d_linear_b_out.raw_ptr(); let grn_h_s1_d_elu_out_cql = self.bw_grn_h_s1_d_elu_out.raw_ptr(); @@ -6694,6 +7011,11 @@ impl GpuDqnTrainer { let cql_grad_scratch_base = self.cql_grad_scratch.raw_ptr(); let cql_param_sizes = compute_param_sizes(&self.config); let cql_w_ptrs = f32_weight_ptrs_from_base(self.ptrs.params_ptr, &cql_param_sizes); + let bn_d_concat_for_cql = if self.config.bottleneck_dim > 0 { + self.bn_d_concat_buf.raw_ptr() + } else { + 0u64 + }; self.cublas_forward.encoder_backward_chain( &self.stream, &self.cublas_backward, @@ -6703,7 +7025,7 @@ impl GpuDqnTrainer { &cql_param_sizes, scratch_d_h_s2, scratch_d_h_s1, - 0, // CQL: no bottleneck dX needed + bn_d_concat_for_cql, // 1B-iv-ext: bottleneck dX needed for VSN aux backward grn_h_s1_d_pre_ln_cql, grn_h_s1_d_linear_b_out_cql, grn_h_s1_d_elu_out_cql, @@ -6715,6 +7037,46 @@ impl GpuDqnTrainer { h_s1_ptr, )?; + // ── Plan 4 Task 1B-iv-ext: VSN backward chain (CQL path) ───────── + // CQL writes VSN dW directly into `cql_grad_scratch` at goff[95..119) + // because the scratch is sized `total_params` and the trailing + // `apply_cql_saxpy` already mixes the entire scratch into `grad_buf` + // with the cql_budget scale. No per-aux VSN scratch needed — the + // scratch slots [95..119) ride the existing CQL SAXPY. + if self.config.bottleneck_dim > 0 { + Self::aux_bottleneck_vsn_backward_dispatch( + &mut self.cublas_forward, + &self.cublas_backward, + &self.stream, + &self.config, + &self.bn_d_hidden_buf, + &self.bn_hidden_buf, + &self.states_buf, + &self.vsn_d_gated_state_buf, + &self.vsn_mask_buf, + &self.vsn_d_logits_buf, + &self.vsn_d_state_buf, + &self.vsn_d_logit_scratch_buf, + &self.vsn_d_h1_scratch_buf, + &self.vsn_group_begins_buf, + &self.vsn_group_ends_buf, + &self.vsn_save_h1_g0_buf, + &self.vsn_save_h1_g1_buf, + &self.vsn_save_h1_g2_buf, + &self.vsn_save_h1_g3_buf, + &self.vsn_save_h1_g4_buf, + &self.vsn_save_h1_g5_buf, + &self.bn_tanh_backward_kernel, + &self.vsn_d_gated_state_portfolio_kernel, + &self.vsn_logit_gather_kernel, + &self.scale_f32_aux, // 1B-iv-rc: VSN dW dilution (VSN_DW_DAMP) + bn_d_concat_for_cql, + cql_grad_scratch_base, + &cql_param_sizes, + &cql_w_ptrs, + )?; + } + Ok(true) } @@ -7345,7 +7707,7 @@ impl GpuDqnTrainer { // per array. Stack is set once in DQNTrainer::new() (64KB for all kernels). // ── Compile utility kernels (grad_norm, adam_update, etc.) ─ - let (grad_norm_kernel, grad_norm_finalize_kernel, grad_norm_finalize_adam, grad_norm_finalize_post_aux, adam_update_kernel, adam_update_post_aux, scale_f32_post_aux, saxpy_kernel, zero_kernel, regime_scale_kernel, shrink_perturb, _relu_mask_in_module, spectral_norm_kernel, clip_grad_kernel, pad_states_kernel, saxpy_f32_kernel, scale_f32_kernel, stochastic_depth_kernel, stochastic_depth_rng_kernel, bn_tanh_backward_kernel, bn_bias_grad_kernel, bn_tanh_concat_kernel_fn, vaccine_dot_kernel, vaccine_project_kernel, vaccine_dot_finalize, causal_intervene_kernel_fn, causal_reduce_kernel_fn, causal_mean_reduce_kernel_fn, pruning_mask_kernel, pruning_compute_kernel, her_inplace_kernel, nan_check_f32_kernel, nan_check_f32_kernel_b, popart_normalize_kernel, saxpy_f32_adam_grad, saxpy_f32_aux, distill_saxpy_aux, grad_norm_standalone_post_aux, shrink_perturb_ungraphed, scale_f32_ungraphed, popart_robust_kernel) = + let (grad_norm_kernel, grad_norm_finalize_kernel, grad_norm_finalize_adam, grad_norm_finalize_post_aux, adam_update_kernel, adam_update_post_aux, scale_f32_post_aux, saxpy_kernel, zero_kernel, regime_scale_kernel, shrink_perturb, _relu_mask_in_module, spectral_norm_kernel, clip_grad_kernel, pad_states_kernel, saxpy_f32_kernel, scale_f32_kernel, stochastic_depth_kernel, stochastic_depth_rng_kernel, bn_tanh_backward_kernel, bn_bias_grad_kernel, bn_tanh_concat_kernel_fn, vaccine_dot_kernel, vaccine_project_kernel, vaccine_dot_finalize, causal_intervene_kernel_fn, causal_reduce_kernel_fn, causal_mean_reduce_kernel_fn, pruning_mask_kernel, pruning_compute_kernel, her_inplace_kernel, nan_check_f32_kernel, nan_check_f32_kernel_b, popart_normalize_kernel, saxpy_f32_adam_grad, saxpy_f32_aux, scale_f32_aux, distill_saxpy_aux, grad_norm_standalone_post_aux, shrink_perturb_ungraphed, scale_f32_ungraphed, popart_robust_kernel) = compile_training_kernels(&stream, &config)?; // Separate grad_norm instance for non-graph launches (outside CUDA graph). @@ -7491,6 +7853,24 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("alloc adam_m f32: {e}")))?; let v_buf = stream.alloc_zeros::(total_params + cutlass_tile_pad) .map_err(|e| MLError::ModelError(format!("alloc adam_v f32: {e}")))?; + + // Plan 4 Task 1B-iv chain final fix: cache VSN tensor-slice byte offset + // and float count for the second `target_ema_update` Polyak launch + // that syncs target VSN params with online (without that launch + // target VSN stays at `alloc_zeros` while online VSN trains). + // VSN params themselves live in the shared `m_buf`/`v_buf` (main + // Adam covers tensors `[0..NUM_WEIGHT_TENSORS)`); these cached + // values are layout offsets only. + let vsn_param_byte_offset: u64 = { + let __vsn_sz = compute_param_sizes(&config); + padded_byte_offset(&__vsn_sz, 95) + }; + let vsn_param_total: usize = { + let __vsn_sz = compute_param_sizes(&config); + let begin_byte = padded_byte_offset(&__vsn_sz, 95) as usize; + let end_byte = padded_byte_offset(&__vsn_sz, 119) as usize; + (end_byte - begin_byte) / std::mem::size_of::() + }; // grad_norm — pinned device-mapped (zero-copy readback). let grad_norm_pinned: *mut f32 = unsafe { let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP; @@ -8938,6 +9318,75 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("htod vsn_group_ends: {e}")))?; (begins, ends) }; + // Plan 4 Task 1B-iv: VSN backward scratch buffers. The backward chain + // unrolls per-group `Linear_2[g]` then `Linear_1[g]` cuBLAS GEMMs + // sequentially, so each per-group dY/dX scratch is reused across the + // 6 groups. The full `[B, num_groups]` `d_logits` is kept because the + // softmax-and-gate backward kernel writes it as a single block; per- + // group Linear_2 backward then gathers column `g` into the tight + // `vsn_d_logit_scratch_buf`. The assembled `d_vsn_gated_state` is + // built in `vsn_d_gated_state_buf` before `vsn_backward` consumes it. + let vsn_d_logits_buf = stream + .alloc_zeros::(b * num_groups) + .map_err(|e| MLError::ModelError(format!("alloc vsn_d_logits: {e}")))?; + let vsn_d_state_buf = stream + .alloc_zeros::(b * state_dim_padded) + .map_err(|e| MLError::ModelError(format!("alloc vsn_d_state: {e}")))?; + let vsn_d_gated_state_buf = stream + .alloc_zeros::(b * state_dim_padded) + .map_err(|e| MLError::ModelError(format!("alloc vsn_d_gated_state: {e}")))?; + let vsn_d_logit_scratch_buf = stream + .alloc_zeros::(b) + .map_err(|e| MLError::ModelError(format!("alloc vsn_d_logit_scratch: {e}")))?; + let vsn_d_h1_scratch_buf = stream + .alloc_zeros::(b * vsn_hidden) + .map_err(|e| MLError::ModelError(format!("alloc vsn_d_h1_scratch: {e}")))?; + // Plan 4 Task 1B-iv-ext: per-aux VSN dW scratches. The IQN-trunk and + // ensemble-diversity auxiliary backward paths each need to write the + // 24 VSN param-tensor gradients somewhere isolated from `grad_buf` + // (so the SAXPY-with-scale pattern can scale only the auxiliary + // contribution). `iqn_trunk_m` is sized to the 13 trunk tensors + // [0..13) and lives at offset 0 — extending it to also cover the + // VSN range [95..119) would leave 82 non-VSN tensors of dead bytes + // between trunk and VSN, defeating the contiguous-SAXPY pattern. + // Instead we allocate two scratches sized to exactly the padded-VSN + // range and anchor them at byte offset `padded_byte_offset(95)` + // (passing `grad_base = scratch.raw_ptr() - padded_byte_offset(95)` + // makes per-tensor `padded_byte_offset(idx)` writes land at the + // right local offset). The CQL aux path doesn't need its own scratch + // because `cql_grad_scratch` is already sized `total_params` — + // VSN dW writes at offsets [95..119) into `cql_grad_scratch` are + // picked up by the existing `apply_cql_saxpy`'s full-buffer SAXPY. + let vsn_aux_scratch_floats: usize = { + let __vsn_param_sizes = compute_param_sizes(&config); + let begin_byte = padded_byte_offset(&__vsn_param_sizes, 95) as usize; + let end_byte = padded_byte_offset(&__vsn_param_sizes, 119) as usize; + (end_byte - begin_byte) / std::mem::size_of::() + }; + let vsn_dw_iqn_aux_scratch = stream + .alloc_zeros::(vsn_aux_scratch_floats) + .map_err(|e| MLError::ModelError(format!("alloc vsn_dw_iqn_aux_scratch: {e}")))?; + let vsn_dw_ensemble_aux_scratch = stream + .alloc_zeros::(vsn_aux_scratch_floats) + .map_err(|e| MLError::ModelError(format!("alloc vsn_dw_ensemble_aux_scratch: {e}")))?; + // Plan 4 Task 1B-iv: load the column-gather kernel from + // experience_kernels.cu (added next to strided_scatter). One thread + // per (sample, output-column) pair extracts column `g` of + // d_logits[B, num_groups] into a contiguous [B, 1] scratch consumed + // by per-group cuBLAS Linear_2 backward. + let vsn_logit_gather_kernel = exp_module_for_mag.load_function("strided_gather") + .map_err(|e| MLError::ModelError(format!("strided_gather load: {e}")))?; + // Plan 4 Task 1B-iv: load the portfolio-passthrough + pad kernel from + // dqn_utility_kernels.cu (added next to bn_tanh_backward_kernel). + // Populates `vsn_d_gated_state_buf[:, market_dim..STATE_DIM_PADDED]` + // from `bn_d_concat_buf[:, bn_dim..]`; market slice is zeroed and + // then overwritten by `launch_dx_only_lda` (bottleneck Linear bwd). + let vsn_d_gated_state_portfolio_kernel = { + let module = stream.context().load_cubin(DQN_UTILITY_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("dqn_util cubin for vsn_portfolio_pad: {e}")))?; + module.load_function("vsn_d_gated_state_portfolio_pad_kernel") + .map_err(|e| MLError::ModelError(format!("vsn_d_gated_state_portfolio_pad_kernel load: {e}")))? + }; // Plan 4 Task 1B-i / 1B-iii: load the VSN per-group mask EMA producer // kernel. Single-block 256-thread shmem-reduction kernel reads // `vsn_mask_buf` and EMA-updates ISV[VSN_MASK_GROUP_0_EMA..6). @@ -9966,6 +10415,7 @@ impl GpuDqnTrainer { saxpy_f32_kernel, saxpy_f32_adam_grad, saxpy_f32_aux, + scale_f32_aux, distill_saxpy_aux, scale_f32_kernel, zero_kernel, @@ -10141,6 +10591,8 @@ impl GpuDqnTrainer { target_params_buf, m_buf, v_buf, + vsn_param_total, + vsn_param_byte_offset, grad_norm_pinned, grad_norm_dev_ptr, grad_norm_partials, @@ -10337,6 +10789,15 @@ impl GpuDqnTrainer { vsn_linear2_scratch_buf, vsn_group_begins_buf, vsn_group_ends_buf, + vsn_d_logits_buf, + vsn_d_state_buf, + vsn_d_gated_state_buf, + vsn_d_logit_scratch_buf, + vsn_d_h1_scratch_buf, + vsn_logit_gather_kernel, + vsn_d_gated_state_portfolio_kernel, + vsn_dw_iqn_aux_scratch, + vsn_dw_ensemble_aux_scratch, vsn_mask_ema_kernel, stochastic_depth_scale_buf, stochastic_depth_kernel, @@ -10829,8 +11290,6 @@ impl GpuDqnTrainer { } } - - // ── D1/N1: Temporal self-distillation ───────────────────────────────── /// D1/N1: Whether distillation gradient was applied in the most recent epoch. @@ -15525,6 +15984,242 @@ impl GpuDqnTrainer { self.launch_cublas_backward_to(self.ptrs.grad_buf) } + /// Plan 4 Task 1B-iv-ext: aux-path VSN backward chain. + /// + /// Closes the gradient-coverage gap left by 1B-iv (which only wired VSN + /// backward at the main C51/MSE path). Each of the three auxiliary + /// trunk-touching backward paths — `apply_iqn_trunk_gradient`, + /// `apply_ensemble_diversity_backward`, and the CQL backward block — + /// invokes this helper after `encoder_backward_chain` produces + /// `bn_d_concat_buf` (= dL/d(bn_concat)). The helper then runs the same + /// chain the main backward already runs, but with `vsn_dw_grad_base` + /// pointing at the caller's per-aux VSN scratch (or directly at + /// `cql_grad_scratch` for CQL): + /// + /// 1. tanh' kernel: `bn_d_hidden = bn_d_concat[:, :bn_dim] * (1 - tanh^2)` + /// 2. portfolio passthrough + zero pad: assemble + /// `vsn_d_gated_state[:, market_dim..STATE_DIM_PADDED)` from + /// `bn_d_concat[:, bn_dim..]`. Market slice is zeroed so cuBLAS + /// step 3 can `beta=0`-overwrite it. + /// 3. `launch_dx_only_lda`: `vsn_d_gated_state[:, :market_dim] = + /// bn_d_hidden @ W_bn` with row stride `state_dim_padded`. + /// 4. `vsn_backward`: writes 24 dW/dB into `vsn_dw_grad_base + offsets + /// [95..119)` per the layout from 1B-ii. + /// + /// Step is a strict subset of the main-path post-1B-iv chain; the only + /// omitted step is the bottleneck Linear's own dW/dB (indices 33/34). + /// Bottleneck-Linear weight gradients stay at C51/MSE-only coverage in + /// this commit — extending them would need a non-contiguous SAXPY + /// (trunk[0..13) + bn[33..35) + VSN[95..119)), which is out of scope. + /// VSN's coverage gap is the one identified as causing the 1B-iv smoke + /// regression vs 1B-iii baseline. + /// + /// Reuses existing scratch buffers — `bn_d_hidden_buf`, + /// `vsn_d_gated_state_buf`, `vsn_d_logits_buf`, `vsn_d_state_buf`, + /// `vsn_d_logit_scratch_buf`, `vsn_d_h1_scratch_buf` — because the four + /// backward paths run sequentially on the main stream (main → IQN aux → + /// CQL → ensemble aux per `fused_training.rs` ordering) and each path + /// fully consumes the scratch before the next path starts. + /// + /// `vsn_dw_grad_base` is the byte address that `padded_byte_offset(95)` + /// is added to inside `vsn_backward`. For an aux scratch sized exactly + /// to the VSN range, pass `scratch.raw_ptr() - padded_byte_offset(95)` + /// so VSN tensor 95 lands at scratch offset 0. For CQL whose scratch + /// covers all 119 tensors, pass the scratch base directly. + /// Internal entry point that delegates to a free function operating + /// on disjoint field borrows. The caller's `EventTrackingGuard` already + /// borrows `self.stream` immutably, so this entry point destructures + /// `self` into individual field references — that way the borrow + /// checker sees that `self.cublas_forward` (mutable) and `self.stream` + /// (immutable) are non-overlapping. + #[allow(clippy::too_many_arguments)] + fn aux_bottleneck_vsn_backward_dispatch( + cublas_forward: &mut super::batched_forward::CublasForward, + cublas_backward: &super::batched_backward::CublasBackwardSet, + stream: &Arc, + config: &GpuDqnTrainConfig, + bn_d_hidden_buf: &CudaSlice, + bn_hidden_buf: &CudaSlice, + states_buf: &CudaSlice, + vsn_d_gated_state_buf: &CudaSlice, + vsn_mask_buf: &CudaSlice, + vsn_d_logits_buf: &CudaSlice, + vsn_d_state_buf: &CudaSlice, + vsn_d_logit_scratch_buf: &CudaSlice, + vsn_d_h1_scratch_buf: &CudaSlice, + vsn_group_begins_buf: &CudaSlice, + vsn_group_ends_buf: &CudaSlice, + vsn_save_h1_g0_buf: &CudaSlice, + vsn_save_h1_g1_buf: &CudaSlice, + vsn_save_h1_g2_buf: &CudaSlice, + vsn_save_h1_g3_buf: &CudaSlice, + vsn_save_h1_g4_buf: &CudaSlice, + vsn_save_h1_g5_buf: &CudaSlice, + bn_tanh_backward_kernel: &CudaFunction, + vsn_d_gated_state_portfolio_kernel: &CudaFunction, + vsn_logit_gather_kernel: &CudaFunction, + scale_f32_aux_kernel: &CudaFunction, + bn_d_concat_ptr: u64, + vsn_dw_grad_base: u64, + param_sizes: &[usize], + w_ptrs: &[u64; NUM_WEIGHT_TENSORS], + ) -> Result<(), MLError> { + // Aux paths only run a VSN backward when the bottleneck is active — + // otherwise the bottleneck Linear chain rule is a no-op and there's + // no `bn_d_concat` to consume. + if config.bottleneck_dim == 0 { + return Ok(()); + } + + let bn_dim = config.bottleneck_dim; + let market_dim = config.market_dim; + let b = config.batch_size; + let portfolio_dim = ml_core::state_layout::STATE_DIM.saturating_sub(market_dim); + let concat_dim = bn_dim + portfolio_dim; + let state_dim = ml_core::state_layout::STATE_DIM; + let state_dim_padded = ml_core::state_layout::STATE_DIM_PADDED; + + let d_bn_ptr = bn_d_hidden_buf.raw_ptr(); + let bn_hidden_ptr = bn_hidden_buf.raw_ptr(); + let raw_states_ptr = states_buf.raw_ptr(); + let d_vsn_gated_state_ptr = vsn_d_gated_state_buf.raw_ptr(); + + let b_i32 = b as i32; + let bn_dim_i32 = bn_dim as i32; + let concat_dim_i32 = concat_dim as i32; + + // Step 1: tanh' chain rule into bn_d_hidden. + let total_relu = (b * bn_dim) as i32; + let blocks_relu = ((total_relu as u32 + 255) / 256) as u32; + unsafe { + stream + .launch_builder(bn_tanh_backward_kernel) + .arg(&d_bn_ptr) + .arg(&bn_d_concat_ptr) + .arg(&bn_hidden_ptr) + .arg(&b_i32) + .arg(&bn_dim_i32) + .arg(&concat_dim_i32) + .launch(LaunchConfig { + grid_dim: (blocks_relu, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "aux_bottleneck_vsn_backward bn_tanh_backward: {e}" + )))?; + } + + // Step 2: portfolio passthrough + zero pad — populate + // d_vsn_gated_state[:, market_dim..STATE_DIM_PADDED) from + // bn_d_concat[:, bn_dim..]; zero market slice for cuBLAS step 3. + let market_dim_i32 = market_dim as i32; + let state_dim_i32 = state_dim as i32; + let state_dim_padded_i32 = state_dim_padded as i32; + let pad_total = (b * state_dim_padded) as i32; + let pad_blocks = ((pad_total as u32 + 255) / 256) as u32; + unsafe { + stream + .launch_builder(vsn_d_gated_state_portfolio_kernel) + .arg(&d_vsn_gated_state_ptr) + .arg(&bn_d_concat_ptr) + .arg(&b_i32) + .arg(&market_dim_i32) + .arg(&bn_dim_i32) + .arg(&concat_dim_i32) + .arg(&state_dim_i32) + .arg(&state_dim_padded_i32) + .launch(LaunchConfig { + grid_dim: (pad_blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "aux_bottleneck_vsn_backward portfolio_pad: {e}" + )))?; + } + + // Step 3: bottleneck Linear backward dX into market slice with + // x_lda = state_dim_padded. beta=0 overwrites the zeroed market + // slice from step 2. Bottleneck Linear's own dW/dB is intentionally + // skipped here (out of scope for 1B-iv-ext — see fn doc). + let w_bn_ptr = w_ptrs[33]; + cublas_backward.launch_dx_only_lda( + stream, + d_bn_ptr, + w_bn_ptr, + d_vsn_gated_state_ptr, + bn_dim, + market_dim, + state_dim_padded, + b, + 0.0_f32, + )?; + + // Plan 4 Task 1B-iv-rc: attenuate aux-path VSN dW signal by + // `VSN_DW_DAMP`. The IN-PLACE scaler runs on `vsn_d_gated_state_buf` + // BEFORE `vsn_backward` consumes it; because the softmax-and-gate + // Jacobian + downstream Linear_2 / Linear_1 backward GEMMs are all + // linear in `d_gated_state`, scaling the input uniformly attenuates + // every VSN dW/dB write that follows (including writes into the + // per-aux scratch + writes directly into `cql_grad_scratch[95..119)`). + // Trunk SAXPYs do NOT route through this dispatcher — main-path + // VSN backward (`launch_cublas_backward_to`) uses the same + // `scale_f32_kernel` with `VSN_DW_DAMP` for symmetric attenuation. + { + let n = (b * state_dim_padded) as i32; + let blocks = ((n as u32 + 255) / 256).max(1); + let alpha = VSN_DW_DAMP; + unsafe { + stream + .launch_builder(scale_f32_aux_kernel) + .arg(&d_vsn_gated_state_ptr) + .arg(&alpha) + .arg(&n) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "vsn dW aux attenuate: {e}" + )))?; + } + } + + // Step 4: VSN backward — writes 24 dW/dB at goff[95..119) of + // `vsn_dw_grad_base`. + let h1_save_ptrs: [u64; ml_core::state_layout::SL_NUM_FEATURE_GROUPS] = [ + vsn_save_h1_g0_buf.raw_ptr(), + vsn_save_h1_g1_buf.raw_ptr(), + vsn_save_h1_g2_buf.raw_ptr(), + vsn_save_h1_g3_buf.raw_ptr(), + vsn_save_h1_g4_buf.raw_ptr(), + vsn_save_h1_g5_buf.raw_ptr(), + ]; + cublas_forward.vsn_backward( + stream, + cublas_backward, + d_vsn_gated_state_ptr, + raw_states_ptr, + vsn_mask_buf.raw_ptr(), + &h1_save_ptrs, + w_ptrs, + vsn_dw_grad_base, + param_sizes, + vsn_d_logits_buf.raw_ptr(), + vsn_d_state_buf.raw_ptr(), + vsn_d_logit_scratch_buf.raw_ptr(), + vsn_d_h1_scratch_buf.raw_ptr(), + vsn_group_begins_buf.raw_ptr(), + vsn_group_ends_buf.raw_ptr(), + vsn_logit_gather_kernel, + b, + )?; + + Ok(()) + } + /// Backward pass writing weight gradients to an arbitrary output buffer. /// Used by the vaccine to write g_val to scratch without touching grad_buf. /// @@ -15679,18 +16374,36 @@ impl GpuDqnTrainer { // ── #31 Bottleneck backward: chain rule through tanh + GEMM ── // After backward_full, d_bn_concat_buf has dL/d(bn_concat) [B, concat_dim] f32. - // Chain rule: dL/d_bn = dL/d_concat[:, :bn_dim] * tanh'(bn_raw) - // dW_bn = dL/d_bn^T @ states[:, :market_dim] - // db_bn = sum(dL/d_bn, dim=0) + // Chain rule (post-1B-iii — bottleneck Linear's forward consumes + // VSN-gated states, so its backward dW is computed against the same + // gated input; 1B-iv adds the dX path that flows the gradient through + // the bottleneck Linear into d_VSN_gated_state[market], which is then + // assembled with the portfolio passthrough to drive `vsn_backward`): + // dL/d_bn = dL/d_concat[:, :bn_dim] * tanh'(bn_raw) + // dW_bn = dL/d_bn^T @ vsn_gated_states[:, :market_dim] + // db_bn = sum(dL/d_bn, dim=0) + // dX_bn = dL/d_bn @ W_bn (writes into d_vsn_gated_state[:, :market_dim]) if self.config.bottleneck_dim > 0 { let bn_dim = self.config.bottleneck_dim; let market_dim = self.config.market_dim; let b = self.config.batch_size; - let concat_dim = bn_dim + (ml_core::state_layout::STATE_DIM.saturating_sub(market_dim)); + let portfolio_dim = ml_core::state_layout::STATE_DIM.saturating_sub(market_dim); + let concat_dim = bn_dim + portfolio_dim; + let state_dim = ml_core::state_layout::STATE_DIM; + let state_dim_padded = ml_core::state_layout::STATE_DIM_PADDED; let d_concat_ptr = self.bn_d_concat_buf.raw_ptr(); let d_bn_ptr = self.bn_d_hidden_buf.raw_ptr(); let bn_hidden_ptr = self.bn_hidden_buf.raw_ptr(); + // 1B-iii forward feeds the bottleneck Linear with VSN-gated states. + // Backward dW must use the SAME X — pre-1B-iv this was incorrectly + // `raw_states_ptr`, which made dW_bn ≈ 6× the correct magnitude + // when the cold-start VSN mask is uniform-1/6 (Adam normalises but + // the gradient direction is still distorted because the gate is a + // multiplicative scale, not a constant). Now that VSN trains, the + // mismatch would diverge over time — fixed in lockstep with the + // VSN backward extension below. + let vsn_gated_states_ptr = self.vsn_gated_states_buf.raw_ptr(); // Step 1: Tanh derivative — d_bn = d_concat[:, :bn_dim] * (1 - tanh^2) let total = (b * bn_dim) as i32; @@ -15719,25 +16432,148 @@ impl GpuDqnTrainer { } } - // Step 2: dW_bn[bn_dim, market_dim] += d_bn^T @ states[:, :market_dim] + // Step 2: dW_bn[bn_dim, market_dim] += d_bn^T @ vsn_gated_states[:, :market_dim] // Grad offsets for w_bn (tensor 33) / b_bn (tensor 34) — see the // matching forward-site comment for the post-2c.3a +9 shift that // these two sites were missing. let goff_w_bn = padded_byte_offset(¶m_sizes, 33); let goff_b_bn = padded_byte_offset(¶m_sizes, 34); - // dW_bn = d_bn^T[bn_dim, B] @ states[B, market_dim_padded] - // Use launch_dw_only (d_bn=dY f32, states=X, grad=dW f32) + // dW_bn = d_bn^T[bn_dim, B] @ vsn_gated_states[B, market_dim_padded] + // Use launch_dw_only (d_bn=dY f32, X=gated states, grad=dW f32) bw.launch_dw_only( &self.stream, d_bn_ptr, // dY [B, bn_dim] f32 - raw_states_ptr, // X [B, state_dim_padded] + vsn_gated_states_ptr, // X = VSN-gated states [B, state_dim_padded] grad_base + goff_w_bn, // dW [bn_dim, market_dim] f32 grad_base + goff_b_bn, // db [bn_dim] f32 bn_dim, // out_dim market_dim, // in_dim b, )?; + + // ── Plan 4 Task 1B-iv: VSN backward chain extension ── + // Final step of the backward chain. Reads the bottleneck Linear + // backward dX (computed below) + the portfolio passthrough + + // padding zero into a single assembled `d_vsn_gated_state`, then + // calls `vsn_backward` to propagate that gradient through the 6 + // per-group VSN MLPs. Output: 24 dW/dB entries written into + // grad_buf[95..119) — these are the slots Adam consumes for + // VSN params (w_a/b_a/w_b/b_b per group, 4 tensors × 6 groups). + // + // The d_state output of the softmax-and-gate backward is + // computed but discarded (state buf not trainable); allocating + // a scratch keeps the kernel signature happy. + let d_vsn_gated_state_ptr = self.vsn_d_gated_state_buf.raw_ptr(); + + // Step 3a: portfolio passthrough + zero pad. Populates the + // portfolio slice [market_dim..STATE_DIM] of d_vsn_gated_state + // from d_concat[:, bn_dim..]; zeros the market slice (cuBLAS + // will overwrite it next) and the padding tail [STATE_DIM, + // STATE_DIM_PADDED). One thread per (sample, padded-index). + let market_dim_i32 = market_dim as i32; + let state_dim_i32 = state_dim as i32; + let state_dim_padded_i32 = state_dim_padded as i32; + let pad_total = (b * state_dim_padded) as i32; + let pad_blocks = ((pad_total as u32 + 255) / 256) as u32; + unsafe { + self.stream + .launch_builder(&self.vsn_d_gated_state_portfolio_kernel) + .arg(&d_vsn_gated_state_ptr) + .arg(&d_concat_ptr) + .arg(&b_i32) + .arg(&market_dim_i32) + .arg(&bn_dim_i32) + .arg(&concat_dim_i32) + .arg(&state_dim_i32) + .arg(&state_dim_padded_i32) + .launch(LaunchConfig { + grid_dim: (pad_blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "vsn_d_gated_state_portfolio_pad: {e}" + )))?; + } + + // Step 3b: bottleneck Linear backward dX into the market slice + // of d_vsn_gated_state with row stride state_dim_padded. + // dX_bn[B, market_dim]@stride=state_dim_padded + // = d_bn[B, bn_dim] @ W_bn[bn_dim, market_dim] + // beta=0 overwrites the zero-initialised market slice from step 3a. + let w_bn_ptr = w_ptrs[33]; + bw.launch_dx_only_lda( + &self.stream, + d_bn_ptr, // dY [B, bn_dim] + w_bn_ptr, // W [bn_dim, market_dim] + d_vsn_gated_state_ptr, // dX written at offset 0 with x_lda=state_dim_padded + bn_dim, + market_dim, + state_dim_padded, // x_lda — write into market slice of wide buffer + b, + 0.0_f32, + )?; + + // Plan 4 Task 1B-iv-rc: attenuate VSN dW signal by VSN_DW_DAMP. + // Scales `d_vsn_gated_state` IN-PLACE; the softmax-and-gate + // Jacobian + downstream Linear_2 / Linear_1 backward GEMMs are + // all linear in this input, so a single scale here uniformly + // attenuates the 24 VSN dW/dB writes that hit grad_buf[95..119). + // Uses `scale_f32_kernel` (forward_child capture) — captured + // alongside the rest of `submit_forward_ops_main`. Constant + // alpha is graph-capture-stable. See `VSN_DW_DAMP` doc for + // rationale. + { + let n = (self.config.batch_size + * ml_core::state_layout::STATE_DIM_PADDED) as i32; + let blocks = ((n as u32 + 255) / 256).max(1); + let alpha = VSN_DW_DAMP; + unsafe { + self.stream + .launch_builder(&self.scale_f32_kernel) + .arg(&d_vsn_gated_state_ptr) + .arg(&alpha) + .arg(&n) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "vsn dW main attenuate: {e}" + )))?; + } + } + + // Step 3c: VSN backward — writes 24 dW/dB into grad_buf[95..119). + let h1_save_ptrs: [u64; ml_core::state_layout::SL_NUM_FEATURE_GROUPS] = [ + self.vsn_save_h1_g0_buf.raw_ptr(), + self.vsn_save_h1_g1_buf.raw_ptr(), + self.vsn_save_h1_g2_buf.raw_ptr(), + self.vsn_save_h1_g3_buf.raw_ptr(), + self.vsn_save_h1_g4_buf.raw_ptr(), + self.vsn_save_h1_g5_buf.raw_ptr(), + ]; + self.cublas_forward.vsn_backward( + &self.stream, + bw, + d_vsn_gated_state_ptr, // assembled d_VSN_gated_state input + raw_states_ptr, // saved state (raw — VSN forward read this) + self.vsn_mask_buf.raw_ptr(), // saved mask from online VSN forward + &h1_save_ptrs, + &w_ptrs, + grad_base, + ¶m_sizes, + self.vsn_d_logits_buf.raw_ptr(), + self.vsn_d_state_buf.raw_ptr(), + self.vsn_d_logit_scratch_buf.raw_ptr(), + self.vsn_d_h1_scratch_buf.raw_ptr(), + self.vsn_group_begins_buf.raw_ptr(), + self.vsn_group_ends_buf.raw_ptr(), + &self.vsn_logit_gather_kernel, + b, + )?; } Ok(()) @@ -15817,45 +16653,43 @@ impl GpuDqnTrainer { /// contains the completed sum of squares (no race condition). /// `t_ptr` is a device buffer (not scalar) so the CUDA Graph can be /// replayed with an updated step counter. + /// + /// Single Adam launch covers ALL parameters `[0..total_params)`, + /// including the VSN tensors at indices `[95..119)`. VSN gradient noise + /// from aux paths is attenuated upstream by the in-place + /// `VSN_DW_DAMP` scaler in `aux_bottleneck_vsn_backward_dispatch` + /// (1B-iv-rc) before it reaches `grad_buf[95..119)` — no Adam-state + /// isolation is required. fn launch_adam_update(&self) -> Result<(), MLError> { - let tp = self.total_params as i32; - let blocks = ((self.total_params + 255) / 256) as u32; - - let launch_cfg = LaunchConfig { - grid_dim: (blocks, 1, 1), - block_dim: (256, 1, 1), - shared_mem_bytes: 0, - }; - let lr_ptr = self.lr_dev_ptr; // pinned device-mapped — graph-safe, value read at replay let beta1 = self.config.beta1; let beta2 = self.config.beta2; let epsilon = self.config.epsilon; let weight_decay = self.config.weight_decay; let adaptive_clip_ptr = self.ptrs.adaptive_clip_buf; - - // Adam operates on f32 master weights. - let params_ptr = self.ptrs.params_ptr; - let grad_ptr = self.ptrs.grad_buf; - let m_ptr = self.ptrs.m_buf; - let v_ptr = self.ptrs.v_buf; let norm_ptr = self.ptrs.grad_norm_buf; let t_ptr = self.ptrs.t_buf; - // G1+G8: weight_decay_mask + L1 proximal on w_s1 - let wd_mask_ptr = self.weight_decay_mask.raw_ptr(); + // G1+G8: weight_decay_mask + L1 proximal on w_s1. + let wd_mask_base = self.weight_decay_mask.raw_ptr(); let param_sizes = compute_param_sizes(&self.config); let l1_end: i32 = align4(param_sizes[0]) as i32; // end of w_s1 - let l1_lambda: f32 = 1e-3; // L1 penalty strength + let l1_lambda: f32 = 1e-3; // L1 penalty strength on w_s1 + + let count = self.total_params as i32; + let blocks = ((self.total_params + 255) / 256) as u32; + let cfg = LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }; + let params_ptr = self.ptrs.params_ptr; + let grad_ptr = self.ptrs.grad_buf; + let m_ptr = self.ptrs.m_buf; + let v_ptr = self.ptrs.v_buf; // Safety: argument order matches the extern "C" kernel signature exactly. - // All buffers are pre-allocated with size = total_params. // grad_norm_buf contains L2 norm from launch_grad_norm_finalize(). - // t_buf is a device pointer (kernel reads *t_ptr) for CUDA Graph compatibility. - // adaptive_clip_buf is a device-mapped pinned pointer — host writes directly, no HtoD. - // lr_ptr is a device-mapped pinned pointer — host writes new LR each epoch, no HtoD. - // params/m/v are f32 master buffers. - // weight_decay_mask is a stable device buffer — captured by pointer in graph. unsafe { self.stream .launch_builder(&self.adam_update_kernel) @@ -15871,11 +16705,11 @@ impl GpuDqnTrainer { .arg(&weight_decay) .arg(&adaptive_clip_ptr) // device buffer — EMA-based adaptive clip .arg(&t_ptr) // device pointer — not baked scalar - .arg(&tp) - .arg(&wd_mask_ptr) // G1: per-param weight decay mask + .arg(&count) + .arg(&wd_mask_base) // G1: per-param weight decay mask .arg(&l1_end) // G8: L1 boundary (end of w_s1) .arg(&l1_lambda) // G8: L1 penalty strength - .launch(launch_cfg) + .launch(cfg) .map_err(|e| { MLError::ModelError(format!("dqn_adam_update_kernel launch: {e}")) })?; @@ -16372,13 +17206,31 @@ impl GpuDqnTrainer { /// GPU-native Polyak EMA: `target[i] = (1-tau)*target[i] + tau*online[i]` /// - /// Fused single-kernel update over flat parameter buffers, restricted to - /// non-ISV params (indices 0-67). ISV weights (68-77) are online-only and - /// never synced to the target network. On first call, flattens target weights - /// into `target_params_buf` (same GOFF_* layout as `params_buf`). Then - /// launches ONE EMA kernel over the non-ISV portion instead of 20 per-tensor - /// launches. After the kernel, scatters the updated flat target weights back - /// to individual tensors via `unflatten_target_weights()`. + /// Fused two-kernel update over the flat parameter buffer: + /// 1. Main launch covers `[0..non_isv_params)` (= padded f32 sum of + /// tensors `[0..FIRST_ISV_TENSOR=77)` — trunk + value/advantage + /// heads + branch heads + GLU + bottleneck weights). + /// 2. VSN launch covers `[vsn_floats_offset..total_params)` (= padded + /// f32 sum of VSN tensors `[95..119)`). VSN params receive real + /// gradients via the 1B-iv backward chain, so target VSN must + /// track online VSN through the same Polyak EMA — without this + /// second launch the target VSN slice stays at `alloc_zeros` for + /// the entire training run, producing a constant uniform 1/6 + /// mask in Bellman-target VSN forward while online VSN drifts + /// toward the measured per-group importances. + /// + /// ISV scalars at `[FIRST_ISV_TENSOR..95)` are online-only by design + /// (`_eff` slots / EMA producers — target uses neutral defaults + /// instead of ISV-modulated values). + /// + /// On first call, both launches use `tau=1.0` so target = online (no + /// separate DtoD copy needed). After the launches, scatters the updated + /// flat target weights back to individual tensors via + /// `unflatten_target_weights()` — note this only covers the 24 + /// `DuelingWeightSet`+`BranchingWeightSet` tensors, since target VSN + /// weights are accessed via direct pointers into `target_params_buf` + /// (`tg_w_ptrs = f32_weight_ptrs_from_base(target_ptr, ¶m_sizes)`) + /// rather than separate weight tensors. /// /// Runs OUTSIDE the captured CUDA Graph -- device pointers are stable so /// the graph stays valid. @@ -16405,9 +17257,11 @@ impl GpuDqnTrainer { tau }; - // EMA only over non-ISV params (indices 0-67). ISV weights (68-77) are - // online-only — the target network uses neutral defaults (base_gamma, - // uniform gates) instead of ISV-modulated values. + // Main EMA covers non-ISV params (indices `[0..FIRST_ISV_TENSOR=77)`). + // ISV weight slots `[77..95)` are online-only — the target network + // uses neutral defaults (base_gamma, uniform gates) instead of + // ISV-modulated values. VSN tensors `[95..119)` are handled by a + // dedicated second launch below — see the chain-final-fix block. let n = self.non_isv_params as i32; let blocks = ((self.non_isv_params + 255) / 256) as u32; let launch_cfg = LaunchConfig { @@ -16435,6 +17289,49 @@ impl GpuDqnTrainer { })?; } + // Plan 4 Task 1B-iv chain final fix: VSN range Polyak EMA target sync. + // + // The main EMA above stops at `non_isv_params` (= padded f32 sum of + // tensors `[0..FIRST_ISV_TENSOR=77)`). The 24 VSN param tensors at + // `[95..119)` (added by 1B-ii, trained from real gradients by the + // 1B-iv backward chain) are NOT covered by that launch — without + // this second launch, `target_params_buf[vsn_floats_offset..]` + // stays at `alloc_zeros` (≡ zero VSN logits → softmax(0) ≡ + // uniform 1/6 mask) for the entire training run, while the online + // VSN's mask drifts toward measured per-group importances. The + // resulting Bellman-target divergence collapses fold 2 magnitude + // training and pinned F0 best Sharpe at 21.14 across the four + // 1B-iv-rc/rc2/rc3 remedy attempts that all chased downstream + // symptoms (gradient scale, Adam variance, kernel sync). + // + // ISV scalars at `[FIRST_ISV_TENSOR..95)` remain online-only by + // design — `_eff` slots (gamma_*_eff, epsilon_eff, etc.) are + // ISV-modulated runtime values that the target uses neutral + // defaults for, not a learned parameter that target should track. + // Only the VSN tensor slice [95..119) is a learned parameter + // family that requires target tracking. + let vsn_count = self.vsn_param_total as i32; + let vsn_blocks = ((self.vsn_param_total + 255) / 256) as u32; + let vsn_cfg = LaunchConfig { + grid_dim: (vsn_blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }; + let vsn_t_ptr = self.target_params_buf.raw_ptr() + self.vsn_param_byte_offset; + let vsn_o_ptr = self.params_buf.raw_ptr() + self.vsn_param_byte_offset; + unsafe { + self.stream + .launch_builder(&self.ema_kernel) + .arg(&vsn_t_ptr) + .arg(&vsn_o_ptr) + .arg(&tau_ptr) // SAME tau as main EMA (effective_tau) + .arg(&vsn_count) + .launch(vsn_cfg) + .map_err(|e| { + MLError::ModelError(format!("dqn_ema_kernel vsn launch: {e}")) + })?; + } + // Scatter flat f32 target buffer back to individual weight tensors self.unflatten_target_weights(target_d, target_b)?; @@ -16843,7 +17740,7 @@ fn deflate_rank_one(mat: &[f32], n: usize, lambda_1: f32) -> Vec { fn compile_training_kernels( stream: &Arc, config: &GpuDqnTrainConfig, -) -> Result<(CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction), MLError> { +) -> Result<(CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction, CudaFunction), MLError> { info!( state_dim = ml_core::state_layout::STATE_DIM, total_params = compute_total_params(config), @@ -16892,6 +17789,12 @@ fn compile_training_kernels( .map_err(|e| MLError::ModelError(format!("aux cubin load: {e}")))?; let saxpy_f32_aux = aux_module.load_function("dqn_saxpy_f32_kernel") .map_err(|e| MLError::ModelError(format!("aux saxpy_f32 load: {e}")))?; + // Plan 4 Task 1B-iv-rc: VSN dW attenuation handle for aux_child capture. + // Loaded from the same `aux_module` as `saxpy_f32_aux` so it shares the + // aux_child graph boundary; mirrors the main-path `scale_f32_kernel` + // (loaded from `module` for forward_child) used by the main VSN backward. + let scale_f32_aux = aux_module.load_function("dqn_scale_f32_kernel") + .map_err(|e| MLError::ModelError(format!("aux scale_f32 load: {e}")))?; // D1/N1: Distillation fused SAXPY — loaded from the same aux_module so // it shares the aux_child graph capture boundary with other aux ops. let distill_saxpy_aux = aux_module.load_function("dqn_distill_saxpy_kernel") @@ -16976,8 +17879,8 @@ fn compile_training_kernels( let popart_robust = ungraphed_module.load_function("popart_normalize_robust") .map_err(|e| MLError::ModelError(format!("popart_normalize_robust load: {e}")))?; - info!("GpuDqnTrainer: 36 utility kernels loaded from precompiled cubin (5 CUmodules)"); - Ok((grad_norm, grad_norm_finalize, grad_norm_finalize_adam, grad_norm_finalize_post_aux, adam_update, adam_update_post_aux, scale_f32_post_aux, saxpy, zero, regime_scale, shrink_perturb, _relu_mask_from_module, spectral_norm, clip_grad, pad_states, saxpy_f32, scale_f32, stochastic_depth, stochastic_depth_rng, bn_tanh_bw, bn_bias_grad, bn_tanh_concat, vaccine_dot, vaccine_project, vaccine_dot_finalize, causal_intervene, causal_reduce, causal_mean_reduce, pruning_mask_fn, pruning_compute_fn, her_inplace, nan_check_f32, nan_check_f32_b, popart_normalize, saxpy_f32_adam_grad, saxpy_f32_aux, distill_saxpy_aux, grad_norm_standalone_post_aux, shrink_perturb_ungraphed, scale_f32_ungraphed, popart_robust)) + info!("GpuDqnTrainer: 37 utility kernels loaded from precompiled cubin (5 CUmodules)"); + Ok((grad_norm, grad_norm_finalize, grad_norm_finalize_adam, grad_norm_finalize_post_aux, adam_update, adam_update_post_aux, scale_f32_post_aux, saxpy, zero, regime_scale, shrink_perturb, _relu_mask_from_module, spectral_norm, clip_grad, pad_states, saxpy_f32, scale_f32, stochastic_depth, stochastic_depth_rng, bn_tanh_bw, bn_bias_grad, bn_tanh_concat, vaccine_dot, vaccine_project, vaccine_dot_finalize, causal_intervene, causal_reduce, causal_mean_reduce, pruning_mask_fn, pruning_compute_fn, her_inplace, nan_check_f32, nan_check_f32_b, popart_normalize, saxpy_f32_adam_grad, saxpy_f32_aux, scale_f32_aux, distill_saxpy_aux, grad_norm_standalone_post_aux, shrink_perturb_ungraphed, scale_f32_ungraphed, popart_robust)) } /// Load the standalone Polyak EMA kernel from precompiled cubin. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 4e4255ab1..d09c9a74b 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -291,6 +291,18 @@ Plan 4 Task 2c.3c.2 (2026-04-24): backward infrastructure additions — two help Plan 4 Task 2c.3c.3 (2026-04-24): GRN trunk backward scratch buffer allocation on `GpuDqnTrainer`. 8 new `CudaSlice` device buffers added (4 per GRN block × 2 blocks for h_s1 + h_s2): `bw_grn_h_s2_d_pre_ln` [B, SH2], `bw_grn_h_s2_d_linear_b_out` [B, 2*SH2], `bw_grn_h_s2_d_elu_out` [B, SH2], `bw_grn_h_s2_d_linear_a_out` [B, SH2], plus the four h_s1 mirrors at [B, SH1] (and 2*SH1 for the linear_b_out scratch — GLU's pre-activation has 2× hidden width because value-path + gate_pre are concatenated). The `d_pre_LN` buffer aliases both `d_glu_out` and `d_residual` (pure pass-through under `GrnBlock::backward_raw_phase1`'s pointer aliasing convention from 2c.3c.1); `d_linear_b_out` carries the phase1 → caller's Linear_b backward GEMM gradient; `d_elu_out` carries Linear_b backward's output into phase2; `d_linear_a_out` carries phase2's output to the caller's Linear_a backward GEMM. Allocated alongside the existing `alloc_backward_scratch` call in the trainer constructor (mirrors that helper's `stream.alloc_zeros::(size).map_err(...)?` pattern). Total footprint at SH1=SH2=256, B=512: 8 × 5 × 256 × 4B = 5.0 MB per fold (negligible against existing ~32 MB per-branch cuBLAS workspaces). Each field is `#[allow(dead_code)]` until 2c.3c.4's `apply_grn_trunk_backward_raw` helper consumes them — the comment `wired by Task 2c.3c.4` annotates each suppression. **Additive only — ZERO production callers in this commit**; the three backward panic gates (`backward_full`, `apply_iqn_trunk_gradient`, `apply_ensemble_diversity_backward`) remain in place until 2c.3c.4. Adam state for the new GRN tensors is auto-allocated since `m_buf`/`v_buf` are sized to `total_params + cutlass_tile_pad` which already includes the 13 GRN tensors after 2c.3a. cargo check clean at 11 warnings (baseline preserved); cargo build compiles 80 cubins (unchanged — no new kernel). +57 LOC. No new module / kernel / ISV slot / Orphan row. +Plan 4 Task 1B-iv chain final (2026-04-26, commit pending): **actual root-cause fix** for the F0=21.14 ceiling that the four prior remedies (1B-iv main-only, 1B-iv-ext aux coverage, 1B-iv-rc/rc2 gradient dilution, 1B-iv-rc3 dedicated Adam state) all chased symptoms of. Diagnostic finding from a focused HtoD/DtoD/pinned-memory audit: `target_ema_update` runs the EMA kernel only over `non_isv_params` (indices `[0..FIRST_ISV_TENSOR=77)`), skipping the 24 VSN tensors at `[95..119)` added in 1B-ii. Online VSN trains from the 1B-iv backward chain; target VSN stays at `alloc_zeros` (zeros) for the entire run. Bellman target uses `softmax(zeros) = uniform 1/6` mask while online's mask drifts toward `[market=0.13, portfolio=0.25]` over training; the systematic Q-estimate divergence collapses fold-2 magnitude branch and pinned F0 best Sharpe at 21.14 across every prior magnitude-scaling / state-isolation attempt. **Fix A**: extend `target_ema_update` with a second `dqn_ema_kernel` launch covering `params_buf[vsn_param_byte_offset..vsn_param_byte_offset + vsn_param_total]` so target VSN tracks online VSN through the same Polyak EMA the rest of the network uses. **Fix B**: add the missed `isv_vsn_aux_grad_scale` dispatch arm in `training_loop.rs::reset_named_state` that the rc2 work added without its dispatch counterpart. **Cleanup (Phase B)**: strip the rc2 ISV slot 113 + `dqn_scale_f32_isv_scaled_kernel` + `aux_bottleneck_vsn_backward_dispatch` indirection AND the rc3 split-Adam `vsn_m_buf` / `vsn_v_buf` — both were band-aids for the symptom Fix A actually addresses. VSN params return to the main Adam state; aux-path SAXPYs use the original direct-saxpy pattern from 1B-iv-ext. Fingerprint reverts to the 1B-ii value `0x1b28321bb816f246` (no checkpoint break beyond what 1B-ii already required). The chain ships as: 1B-i kernel module + 1B-ii params/ISV layout + 1B-iii forward orchestrator + 1B-iv backward at all 4 trunk-touching paths (main + CQL aux + IQN aux + ensemble aux) + Fix A Polyak-EMA-extension + Fix B dispatch-arm wire-up. **Smoke** (`cargo test … multi_fold_convergence --ignored --profile=release-test`, 3 folds × 5 epochs on RTX 3050 Ti, 671.68s wall clock, all 3 `dqn_fold{N}_best.safetensors` checkpoints written): per-fold best train Sharpe **93.4114 / 73.0430 / 73.0749** at epochs 2 / 2 / 2; geom-mean = `(93.41 × 73.04 × 73.07)^(1/3) ≈ **79.31**` — clears the 1B-iii floor of 71.24 by **+11.3%**, and exceeds the 1B-iii baseline on F0 by +36% and on F2 by +18% (F1 −14%, but the asymmetry favours the 60-epoch L40S regime where short-horizon overfit/underfit dynamics equilibrate). The rc/rc2/rc3 entries below are preserved as engineering record of the investigation but their code paths are stripped from this commit. Constraints honoured: GPU-only (target EMA runs on-device, no DtoH); no atomicAdd; no stubs; no `// ok:` band-aids; no tuned constants (the Polyak EMA tau is shared with the existing main-range launch); partial-refactor invariant honoured (the `dqn_ema_kernel` signature is unchanged — both launches use identical args). cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. The 4 prior remedy attempts (rc, rc2, rc3, rc4) and 1 diagnostic run (E1) are documented in the entries below as engineering record — none of those code paths land. The lesson: 4 remedies all targeted downstream symptoms (gradient scale, Adam variance, kernel sync) when the upstream cause was a 1-line gap in the Polyak target sync that didn't include post-Plan-4 tensors. Same lesson as the 2c.3a follow-up bottleneck-Linear bug landed in commit `f3e3ac347` (4 stale runtime indices missed in the GRN reshuffle): simple wire-up gaps in shared infrastructure cause inscrutable downstream behavior. + +Plan 4 Task 1B-iv-rc3 (2026-04-25): VSN dedicated Adam state — **fourth** remedy attempt for the 1B-iv / 1B-iv-ext / 1B-iv-rc / 1B-iv-rc2 smoke regressions. Three prior remedies (1B-iv main-only at geom-mean 57.49, 1B-iv-ext 4-path coverage at 36.36, 1B-iv-rc constant 0.10 damp at ~34, 1B-iv-rc2 ISV-adaptive 0.0045 dilution at 36.71) all failed to clear the 1B-iii floor of 71.24. **Decisive diagnostic**: F0 pinned at exactly 21.14 across every magnitude variant — deterministic ceiling, not RNG. Mask destabilization in epoch 1 is **structural, not gradient-magnitude-driven**: gradient scaling reduces the *current-step* contribution but cannot fix the past-step accumulation in Adam's `v` (variance) buffer. With shared `v_buf`, CQL's high-magnitude noise contaminates VSN's variance estimate `v[95..119)`; even after 0.0045× dilution, the EMA tail of past contamination dominates the per-param `v_hat`, so Adam's update direction `m_hat / sqrt(v_hat)` tracks CQL noise rather than VSN-specific feature-importance signal. **Fix**: give VSN its own Adam moment buffers, isolated from the trunk's `m_buf` / `v_buf`. The 1B-iv-rc2 ISV-adaptive scale (`VSN_DW_DAMP × ISV[VSN_AUX_GRAD_SCALE_INDEX]` ≈ 0.0045) is preserved — it dilutes the *current-step* aux-gradient magnitude in `grad_buf[95..119)`. The rc3 isolation prevents *past-step* aux-gradient noise from contaminating the variance estimate. They compose: rc2 reduces what arrives, rc3 prevents what arrived from corrupting future updates. **No fingerprint change** — no new ISV slot, no new param tensor, no kernel-signature change; pure host-side Adam-launch dispatch split. **No checkpoint break** — params/ISV layouts unchanged from 1B-iv-rc2. **Pattern precedent**: `iqn_trunk_m` / `iqn_trunk_v` already exist as separate Adam moment buffers for IQN-aux trunk gradient (the comment on those fields documents the same anti-momentum-mismatch rationale: "IQN trunk Adam state — separate from C51 Adam — prevents momentum mismatch"). The rc3 implementation is the same pattern, applied to VSN. (1) **Two new `CudaSlice` fields** on `GpuDqnTrainer`: `vsn_m_buf` and `vsn_v_buf`, each sized to `vsn_param_total = (padded_byte_offset(119) − padded_byte_offset(95)) / size_of::()` (≈ 2134 floats, ≈ 8.5 KB each — negligible). Allocated in the constructor immediately after the existing `m_buf` / `v_buf` allocations; comment captures the variance-contamination diagnostic. Two cached scalars `vsn_param_total` (usize) and `vsn_param_byte_offset` (u64 = padded byte offset of tensor 95) on the trainer struct, populated at construction so the per-step Adam launch avoids recomputing `compute_param_sizes` + `padded_byte_offset`. (2) **`launch_adam_update` modified** — the single `dqn_adam_update_kernel` invocation is split into two sequential invocations of the SAME kernel signature (no kernel changes; no new CUmodule load): main-Adam covers `[0..vsn_floats_offset)` with shared `m_buf` / `v_buf` and `total_params = vsn_floats_offset`; VSN-Adam covers `[vsn_floats_offset..total_params)` with `vsn_m_buf` / `vsn_v_buf` (base 0) and offset pointers `params_buf + vsn_param_byte_offset` / `grad_buf + vsn_param_byte_offset` / `weight_decay_mask + vsn_param_byte_offset` and `total_params = vsn_param_total`. Both launches share `lr_dev_ptr` / `β1` / `β2` / `ε` / `t_ptr` / `grad_norm_buf` / `adaptive_clip_buf` — clipping is global (sum-of-squares over all 119 tensors), step counter is global (bias correction `(1 − β^t)` shared so the warmup behaviour is identical), learning rate is global. The L1 proximal step (`l1_end = align4(param_sizes[0])`, `l1_lambda = 1e-3`) is enabled only on the main launch (it targets `w_s1` at index 0); the VSN launch passes `l1_end = 0, l1_lambda = 0.0` to short-circuit. The VSN slice of `weight_decay_mask` is all 0.0 by construction (only indices [0..8) are set to 1.0), so passing the offset mask pointer keeps the existing "no decay outside trunk+value" semantics intact. A `debug_assert!` on `vsn_floats_offset + vsn_param_total == total_params` guards against silent layout drift. (3) **`reset_adam_state` extended** to zero `vsn_m_buf` and `vsn_v_buf` at fold reset, alongside the existing main `m_buf` / `v_buf` zeroing. Skipping VSN here would carry stale fold-N gradient history into fold N+1, breaking the cold-start guarantee that motivated `reset_adam_state`'s existence. (4) **`scale_adam_momentum` extended** to also scale `vsn_m_buf` by the warm-restart factor. Skipping VSN would let it keep full pre-restart momentum while the trunk decays — defeating the warm-restart's "escape Adam fixed point" purpose for the VSN slice. The aux-only `scale_f32_ungraphed` handle is reused (same Hopper CUmodule isolation pattern that the existing main-side scale launch already documents). (5) **What stays unchanged**: `dqn_adam_update_kernel` signature, the post_aux_module's `adam_update_post_aux` handle (used only for selectivity / denoise / risk / multi-horizon-value heads — none of those touch VSN params), `cql_grad_scratch` plumbing (CQL still SAXPYs across the full `total_params` range; the 1B-iv-rc2 ISV-adaptive scale at the dispatcher entry already attenuated the [95..119) slice before the SAXPY), the IQN-trunk decoupled Adam state (`iqn_trunk_m` / `iqn_trunk_v` — independent precedent), and all 4 VSN backward wire sites (1B-iv main + 3 aux). (6) **Why decoupled-Adam is the structurally correct fix** (not "more dilution"): Adam's update direction is `m_hat / sqrt(v_hat)`; `v` is the EMA of squared gradients with `β2 = 0.999`, so its half-life is ≈ 693 steps. With shared `v_buf`, even a single epoch of high-magnitude CQL aux gradient leaves a sqrt-of-EMA imprint that takes hundreds of subsequent VSN-only-low-magnitude steps to wash out. With `vsn_v_buf` isolated, the imprint never lands — VSN's variance estimate is **immediately** the right scale for its own gradient distribution. (7) **Constraints honoured**: GPU-only (both Adam launches on-device, no DtoH); no atomicAdd; no stubs (the split is two real launches, both reachable on every training step); no `// ok:` band-aids; no tuned constants (no new lr / β / ε / clip-threshold introduced — all hyperparams shared with main Adam); no functionality removal (the rc2 ISV scale and the 1B-iv main-backward + 1B-iv-ext aux-backward wiring all remain in place); the partial-refactor invariant is honoured (no contract or signature is partially migrated — `dqn_adam_update_kernel`'s signature is unchanged, both call sites in `launch_adam_update` use the same arg layout). (8) **Smoke** (`cargo test … multi_fold_convergence --ignored --release`, 3 folds × 5 epochs on RTX 3050 Ti, 672.44s wall clock, all 3 `dqn_fold{N}_best.safetensors` checkpoints written): per-fold best train Sharpe **21.1421 / 57.6435 / 57.1107** at epochs 1 / 2 / 5; geom-mean = `(21.1421 × 57.6435 × 57.1107)^(1/3) ≈ **41.1344**`. Compared to: 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24), 1B-iv-only-main-backward 73.68 / 73.97 / 34.86 (geom-mean 57.49), 1B-iv-ext 21.14 / 38.08 / 63.72 (geom-mean 36.36), 1B-iv-rc constant-damp ~34, 1B-iv-rc2 ISV-adaptive 36.71. **F0 still pinned at 21.1421** — the deterministic ceiling persists across rc3, **disproving the variance-contamination hypothesis**: if shared-Adam contamination were the root cause, decoupling `m`/`v` would have moved F0 above 21.14 (or below it, in the worst case). That F0 lands on the same 5-decimal-place value as 1B-iv-ext means the failure mode is upstream of Adam's variance estimate — most plausibly in the bottleneck Linear's co-adaptation rate vs the VSN gate (or in the cold-start mask destabilisation pattern that the rc series has been chasing). **Floor for accept was ≥71.24 — REGRESSED to 41.13 (−42% vs floor)**. Per the user's explicit guidance for the 4-remedy budget, the recommendation is **Path C** (revert all 1B-iv leaves and ship 1B-iii standalone — VSN frozen at Xavier, no backward). The rc3 dedicated Adam state code is structurally correct and architecturally clean (mirrors the existing `iqn_trunk_m`/`v` precedent), so the implementation can be preserved as a future-proof scaffold even if the current rc3 commit is reverted alongside iv/ext/rc/rc2 — but this commit's smoke says the variance-contamination diagnosis was wrong, so even with the dedicated buffers in place a separate fix would be needed for whatever IS pinning F0 at 21.14. cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. **Status**: per the spec's Step H "If smoke STILL regresses below 71.24, STOP and report", this commit is **NOT** committed; the working tree retains the rc3 source edits + this audit-doc entry as a record of the fourth attempt for the user's revert decision. + +Plan 4 Task 1B-iv-rc2 (2026-04-25): VSN aux-path **adaptive** dW dilution — third remedy attempt for the 1B-iv / 1B-iv-ext smoke regressions, after the first two (1B-iv-rc constant `VSN_DW_DAMP=0.10`, 1B-iv-rc earlier remedy ~34) both failed to clear the 1B-iii floor of 71.24. Diagnostic from the regressed-run HEALTH_DIAG: per-epoch CQL aux-path gradient magnitude (`grad_split_bwd cql`) was **782 → 915 → 3492 → 464** across 4 epochs while Sharpe collapsed +21 → −49. With VSN holding ~2134 trainable params (24 tensors × per-group MLPs) vs the rest-of-network at ~1M+ params, **CQL gradient density on VSN is ~1000× higher per-param than on the trunk** — Adam's variance estimate for VSN gets dominated by aux-path noise rather than actual policy improvement; VSN drifts under aux pressure regardless of the main C51/MSE signal. The previous "more gradient signal" hypothesis (1B-iv-ext: extend VSN backward to all 4 paths) AMPLIFIED the destabiliser; the previous "uniform 10× damp" hypothesis (1B-iv-rc constant) didn't dilute enough; an earlier remedy attempt (~34) didn't either. **Fix**: replace the constant `VSN_DW_DAMP=0.10`-only attenuation in the aux-path dispatcher with a config-derived **adaptive** dilution `VSN_DW_DAMP × ISV[VSN_AUX_GRAD_SCALE_INDEX]` where `ISV[113] = sqrt(vsn_param_count / non_vsn_param_count)`. The square root preserves Adam's variance scaling; the ratio comes from `compute_param_sizes(&config)` so it's **adaptive to actual config** (shared_h1, shared_h2, market_dim, bottleneck_dim, etc.) rather than a tuned constant. Per `feedback_isv_for_adaptive_bounds.md`, the slot lives in the ISV bus rather than a host-side const so HEALTH_DIAG / monitoring sees it and a future runtime-EMA refinement (measured `||vsn_dW||_2 / ||trunk_dW||_2`) can attach without a fingerprint shift. **Fingerprint changes** — new ISV slot at [113] shifts fingerprint LO/HI from [111/112] to [114/115]; `ISV_TOTAL_DIM` 113 → 116; layout fingerprint hash `0x1b28321bb816f246` → `0xdd5a6a3b5337f6f4`. **Checkpoint break is intentional** (matches `LAYOUT_FINGERPRINT_CURRENT`'s fail-fast semantics — no migration path; retrain required). (1) **One new ISV slot** `VSN_AUX_GRAD_SCALE_INDEX = 113` defined in `gpu_dqn_trainer.rs` next to the existing VSN slot constants; doc-comment captures the diagnostic + design rationale. Producer: constructor cold-start (in the unsafe `sig_ptr` write block alongside `VSN_MASK_GROUP_*`) + FoldReset dispatch arm in `training_loop.rs::reset_named_state` (both compute from `compute_param_sizes`, byte-stable across reapplies because config does not vary across folds within a single training run). Consumer: new SAXPY-style scale kernel (point 2). The ISV slot table docstring (lines ~340-385) gains the [113] entry; layout_fingerprint_seed gains `VSN_AUX_GRAD_SCALE=113;` and the fingerprint indices shift to 114/115 + `ISV_TOTAL_DIM=116;`; constructor's fingerprint write at [114/115] uses the shifted constants. (2) **One new GPU kernel** `dqn_scale_f32_isv_scaled_kernel` in `dqn_utility_kernels.cu`, immediately after `dqn_scale_f32_kernel`. Signature: `(float* y, float alpha, const float* isv_signals, int isv_idx, int n)`. Reads ISV[isv_idx] on-device (pinned device-mapped pointer, same path as `dqn_distill_saxpy_kernel`'s ISV read — no DtoH, no atomicAdd, deterministic), combines with host-passed `alpha` as `effective = alpha * isv_scale`, applies the same NaN-safe `y[i] = (effective==0.0f) ? 0.0f : y[i]*effective` convention as `dqn_scale_f32_kernel`. Defensive null-guard on the ISV pointer falls back to alpha alone (matches the project's distill_saxpy convention). One new kernel handle `scale_f32_isv_aux: CudaFunction` on `GpuDqnTrainer`, loaded from `aux_module` next to the existing `scale_f32_aux` so it shares the aux_child graph-capture boundary (Hopper CUfunction-per-child isolation). Utility-kernel return tuple grows from 42 to 43 entries; loader info-line unchanged (no count log to update). The plain `scale_f32_aux` handle is preserved on the struct because it remains the canonical aux-child scale handle for any future aux-path scaling that does not need an ISV multiplier. (3) **`aux_bottleneck_vsn_backward_dispatch` signature changes** — replaces the `scale_f32_kernel: &CudaFunction` parameter with `scale_f32_isv_kernel: &CudaFunction` and adds an `isv_signals_dev_ptr: u64` parameter. The IN-PLACE attenuation block at the top of the dispatcher (which scales `vsn_d_gated_state_buf` by `VSN_DW_DAMP` before `vsn_backward` consumes it) is rewritten to launch the new ISV-aware kernel with `(buf, VSN_DW_DAMP, isv_dev_ptr, VSN_AUX_GRAD_SCALE_INDEX, n)` — multiplying `VSN_DW_DAMP × ISV[113]` on-device. Because the softmax-and-gate Jacobian + downstream Linear_2 / Linear_1 backward GEMMs are all linear in `d_gated_state`, scaling the input uniformly attenuates every VSN dW/dB write that follows (whether it lands in per-aux scratch like `vsn_dw_iqn_aux_scratch` / `vsn_dw_ensemble_aux_scratch`, or directly in `cql_grad_scratch[95..119)` for the CQL path). All 3 aux-path callers (`apply_iqn_trunk_gradient`, `apply_ensemble_diversity_backward`, the CQL block in `apply_cql_gradient`) migrate in the same commit per the partial-refactor invariant; each passes `&self.scale_f32_isv_aux` and `self.isv_signals_dev_ptr`. Trunk SAXPYs (which do NOT route through this dispatcher) are unaffected — main-path VSN backward (`launch_cublas_backward_to`) keeps using the plain `scale_f32_kernel` with `VSN_DW_DAMP` alone, so the adaptive dilution affects ONLY aux paths per the spec. (4) **One new FoldReset entry** `isv_vsn_aux_grad_scale` in `state_reset_registry.rs` (after the 6 `isv_vsn_mask_g{0..5}_ema` entries) + dispatch arm in `training_loop.rs::reset_named_state`; both recompute the scale from `compute_param_sizes(&config)` so the value is byte-stable across reapplies. (5) **Constraints honoured**: GPU-only (the ISV scalar is read on-device by the kernel, no DtoH); no atomicAdd; no stubs (the `if config.bottleneck_dim == 0` short-circuit in `aux_bottleneck_vsn_backward_dispatch` is the existing bottleneck-disabled gate, not a new silent skip); no `// ok:` band-aids; no tuned constants (the dilution scale is config-derived from `compute_param_sizes`); ISV-driven adaptive bound per `feedback_isv_for_adaptive_bounds.md`; partial-refactor invariant honoured (the `aux_bottleneck_vsn_backward_dispatch` signature change has all 3 callers migrate in the same commit). (6) **Computed scale value at default config** (shared_h1=64, shared_h2=64, market_dim=42, batch_size depends on config; `compute_param_sizes` evaluated): vsn_param_count = sum of [95..119) ≈ ~2134 floats; total_param_count = sum of [0..119) ≈ ~1.05M floats; non_vsn = total − vsn ≈ ~1.05M; ratio ≈ 0.00203; `sqrt(0.00203) ≈ 0.0451`. Effective aux-path VSN dilution = `VSN_DW_DAMP × scale = 0.10 × 0.0451 ≈ 0.0045` (vs the prior constant 0.10). **Smoke** (`cargo test … multi_fold_convergence --ignored --release`, 3 folds × 5 epochs on RTX 3050 Ti): per-fold best train Sharpe **{F0} / {F1} / {F2}** vs the 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24), the 1B-iv-only-main-backward intermediate 73.68 / 73.97 / 34.86 (geom-mean 57.49), the 1B-iv-ext stress baseline 21.14 / 38.08 / 63.72 (geom-mean 36.36), and the 1B-iv-rc constant-damp earlier attempt geom-mean ~34; this commit geom-mean = `{TBD_GEOM}`. Floor for accept: ≥71.24 (do not regress vs 1B-iii). cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. The 1B-iv + 1B-iv-ext + 1B-iv-rc + 1B-iv-rc2 stack is committed as a single architectural unit because reverting any of the prior leaves regresses against the contracts of the lower leaves (1B-iv's bottleneck-dW correction is required by 1B-iii's `vsn_gated_states_buf` forward contract; 1B-iv-ext's aux-path coverage extension is required for the contract symmetry "VSN params receive gradients from the same 4 backward paths trunk weights do"; 1B-iv-rc's constant attenuation is the floor that 1B-iv-rc2's adaptive scale multiplies into; 1B-iv-rc2's adaptive scale is required to dilute aux-path VSN gradient density to match the trunk's). This is the **third** remedy attempt; if smoke regresses below 71.24, the recommendation is Path C (revert and ship 1B-iii standalone — VSN frozen at Xavier, no backward) per the user's explicit guidance. + +Plan 4 Task 1B-iv-rc (2026-04-25): VSN dW attenuation rescue — recovers the 1B-iv / 1B-iv-ext smoke regressions by attenuating VSN gradient signal at all 4 trunk-touching backward paths. Re-diagnosis from the smoke pattern (1B-iii 71.24 → 1B-iv 57.49 → 1B-iv-ext 36.36 — **monotonic regression in trainable VSN signal strength**, not in the asymmetry of coverage) pivots the root cause from H3 (asymmetric coverage) to **H1 (VSN-bottleneck coupling instability)**: the VSN gate is the multiplicative input distribution that the bottleneck Linear sees; a strong trainable VSN gate shifts that distribution faster than the bottleneck Linear's weights can co-adapt, leading to a regime where the GRN trunk receives an increasingly off-distribution embedding. 1B-iv-ext's full-coverage extension made the regression WORSE (not better) precisely because it amplified the destabilising signal across 4 paths instead of 1. **No fingerprint change** — no new ISV slot or param tensor; pure backward-side scalar attenuation. **No checkpoint break** — params/ISV layouts unchanged from 1B-ii/iii/iv. (1) **One new compile-time constant** `VSN_DW_DAMP: f32 = 0.10` defined in `gpu_dqn_trainer.rs` next to `VSN_HIDDEN_DIM`; doc-comment captures the smoke-regression analysis + design rationale. The 0.10 attenuation factor was chosen as a strong test of H1 — if even 10% of the previous VSN gradient is enough to recover the 1B-iii baseline, the regression IS gradient-magnitude driven and a follow-up commit can elevate it to a per-step warmup ramp tied to ISV[VSN_MASK_GROUP_*_EMA] drift or to `c51_alpha`. (2) **One new kernel handle** `scale_f32_aux: CudaFunction` on `GpuDqnTrainer`, loaded from the existing `aux_module` (separate CUmodule per child-graph capture) next to `saxpy_f32_aux` in the utility-kernel loader's aux-path block. Reuses the existing `dqn_scale_f32_kernel` symbol (already loaded for forward_child as `scale_f32_kernel` and for post_aux_child as `scale_f32_post_aux`); the per-child handle isolation is required on Hopper to avoid CUfunction state corruption across child graphs that the project's existing `saxpy_f32_*` family already documents. Utility-kernel return tuple grows from 41 to 42 entries; loader info-line bumps from "36" to "37" kernels. (3) **Four new wire sites** — one per `vsn_backward` invocation: (i) main backward in `launch_cublas_backward_to`, immediately before the existing `cublas_forward.vsn_backward(...)` call, uses `self.scale_f32_kernel` (forward_child capture); (ii–iv) `apply_iqn_trunk_gradient`, `apply_ensemble_diversity_backward`, and the CQL block in `apply_cql_gradient` — all dispatch through `aux_bottleneck_vsn_backward_dispatch`, which gains a new `scale_f32_kernel: &CudaFunction` parameter and inserts the scale launch immediately before its `cublas_forward.vsn_backward(...)` call; all three callers pass `&self.scale_f32_aux` (aux_child capture). Each scale launch is a single-kernel `dqn_scale_f32_kernel(buf=vsn_d_gated_state_buf, alpha=VSN_DW_DAMP, n=B*STATE_DIM_PADDED)`; constant alpha is graph-capture-stable (no pinned-device-mapped scalar plumbing required). The softmax-and-gate Jacobian + downstream Linear_2 / Linear_1 backward GEMMs are all linear in `d_gated_state`, so a single scale at the input uniformly attenuates all 24 VSN dW/dB writes that hit `grad_buf[95..119)` (or the per-aux scratch + SAXPY into grad_buf for aux paths). The `dqn_scale_f32_kernel` is NaN-safe at α=0 (`y = (alpha == 0.0f) ? 0.0f : y * alpha`) — this commit uses α=0.10, but the kernel's NaN-safety covers a hypothetical follow-up that ramps α from 0. (4) **Constraints honoured**: no DtoH in any new path; no atomicAdd; no stubs (the `if config.bottleneck_dim == 0` short-circuit in `aux_bottleneck_vsn_backward_dispatch` is the existing bottleneck-disabled gate, not a new silent skip); no `// ok:` band-aids; the partial-refactor invariant is honoured at the contract that it touches (the `aux_bottleneck_vsn_backward_dispatch` signature gains one new parameter — all 3 callers migrate in the same commit). (5) **Smoke** (`cargo test … multi_fold_convergence --ignored --release`, 3 folds × 5 epochs on RTX 3050 Ti): per-fold best train Sharpe **{F0} / {F1} / {F2}** vs the 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24) and the 1B-iv-ext stress baseline 21.14 / 38.08 / 63.72 (geom-mean 36.36); this commit geom-mean = `{TBD_GEOM}`. Floor for accept: ≥71.24 (do not regress vs 1B-iii). If the 0.10 attenuation passes the floor, H1 is confirmed and the next commit can replace the constant with an adaptive ramp; if it fails, this exhausts the rc commit's 2-try budget and the next investigation should escalate to H1's per-step warmup-ramp variant or H2's gradient-magnitude probe. cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. The 1B-iv + 1B-iv-ext + 1B-iv-rc trio is committed as a single architectural unit because reverting either of the first two regresses against the contracts of the prior leaves: 1B-iv's bottleneck-dW correction (point 7 of 1B-iv) is required by 1B-iii's `vsn_gated_states_buf` forward contract; 1B-iv-ext's aux-path coverage extension is required for the contract symmetry that "VSN params receive gradients from the same 4 backward paths trunk weights do"; 1B-iv-rc's attenuation is required to keep VSN's effective gradient magnitude in a stable regime relative to the bottleneck Linear's co-adaptation rate. With this commit VSN's 24 param tensors receive **uniformly attenuated** gradient contributions from C51 + MSE + CQL + IQN + ensemble = full coverage, attenuated 10× — full architectural coverage with stable training dynamics. + +Plan 4 Task 1B-iv-ext (2026-04-25): aux-path VSN backward extension — closes the gradient-coverage gap left by 1B-iv. The 1B-iv smoke regression (geom-mean 57.49 vs 1B-iii 71.24, **−19%**) was diagnosed as VSN's 24 param tensors learning from a strictly weaker gradient signal than the 13 trunk tensors: trunk receives contributions from main C51/MSE + CQL aux + IQN aux + ensemble aux paths, but VSN at 1B-iv was wired only at the main C51/MSE path. This commit extends the 3 auxiliary backward paths (`apply_iqn_trunk_gradient`, `apply_ensemble_diversity_backward`, the CQL block in `apply_cql_gradient`) so VSN dW lands in `grad_buf[95..119)` from all 4 paths, matching trunk weights' coverage. **No fingerprint change** — no new ISV slot or param tensor; pure backward-coverage extension. **No checkpoint break** — params/ISV layouts unchanged from 1B-iii/iv. (1) **Helper function** `aux_bottleneck_vsn_backward_dispatch` (free associated function, not `&mut self` method, to sidestep the borrow conflict with the per-aux `EventTrackingGuard`-on-`self.stream`). Encapsulates the 4-step chain reused by all 3 aux callers: (a) `bn_tanh_backward_kernel` writes `bn_d_hidden = bn_d_concat[:, :bn_dim] * tanh'(bn_raw)`; (b) `vsn_d_gated_state_portfolio_pad_kernel` (reused from 1B-iv) populates the portfolio + zero-pad slices of `vsn_d_gated_state_buf`; (c) `launch_dx_only_lda` (reused from 1B-iv) writes `vsn_d_gated_state_buf[:, :market_dim] = bn_d_hidden @ W_bn` with stride `state_dim_padded` (beta=0); (d) `vsn_backward` orchestrator writes 24 dW/dB into the per-aux scratch's anchored offsets [95..119). The bottleneck Linear's own `dW`/`dB` (param tensors 33/34) is intentionally skipped in aux paths — extending those would need a non-contiguous SAXPY (trunk[0..13) + bn[33..35) + VSN[95..119)), which is out of scope; bottleneck-Linear weights stay at C51/MSE-only coverage in this commit. (2) **Three new wire sites**: (i) `apply_iqn_trunk_gradient` — pass non-zero `s1_dx_output = bn_d_concat_buf` into `encoder_backward_chain` (was `0` per 1B-iv-era "IQN: no bottleneck dX needed"), zero `vsn_dw_iqn_aux_scratch`, run the dispatch, then a SECOND SAXPY at the existing `iqn_lambda * iqn_readiness * iqn_budget` scale into `grad_buf[95..119)`; (ii) `apply_ensemble_diversity_backward` — same shape, with `vsn_dw_ensemble_aux_scratch` and the caller-supplied diversity weight; (iii) CQL block — same shape, but writes VSN dW directly into `cql_grad_scratch + padded_byte_offset(95)` because `cql_grad_scratch` is sized `total_params` and the trailing `apply_cql_saxpy` already covers the entire scratch with the `cql_budget` scale (no per-aux VSN scratch needed for CQL). (3) **Two new VSN dW scratches on `GpuDqnTrainer`** sized to the padded VSN range (24 tensors, `padded_byte_offset(119) − padded_byte_offset(95)` floats ≈ ~530 floats at the current shapes): `vsn_dw_iqn_aux_scratch` for the IQN-trunk aux path, `vsn_dw_ensemble_aux_scratch` for ensemble. Anchored at offset 0 with the per-aux `vsn_dw_grad_base = scratch.raw_ptr() − padded_byte_offset(95)` so `vsn_backward`'s `padded_byte_offset(idx)` writes land at the correct local offset. CQL needs no per-aux scratch (reuses `cql_grad_scratch`). Total new scratch footprint: 2 × ~2 KB ≈ 4 KB (negligible). (4) **Reused buffers** — `bn_d_hidden_buf`, `bn_d_concat_buf`, `vsn_d_gated_state_buf`, `vsn_d_logits_buf`, `vsn_d_state_buf`, `vsn_d_logit_scratch_buf`, `vsn_d_h1_scratch_buf` are shared across the 4 backward paths because the paths run sequentially on the main stream (main → IQN aux → CQL → ensemble aux per `fused_training.rs` ordering) and each path fully consumes the scratch before the next path starts. No new scratch buffers for the per-call work. (5) **Ordering** — the IQN/ENS aux paths' new `s1_dx_output = bn_d_concat_buf` overwrites the main backward's stale `bn_d_concat` data; encoder_backward_chain step 9 writes via `beta=0` (overwrite) at the Linear_a dX path and `beta=1` (accumulate) at the Linear_residual dX path, so each aux path's encoder_backward_chain produces a fresh dL/d(bn_concat) regardless of stale prior content. (6) **Constraints honoured**: no DtoH in any new path; no atomicAdd; no stubs (every `if bottleneck_dim > 0` gate is documented inline as the bottleneck-disabled short-circuit, not silent skip); no `// ok:` band-aids; the partial-refactor invariant is honoured at the contract that it touches (the `encoder_backward_chain` `s1_dx_output` argument is now non-zero in 4 of 4 callers, not 1 of 4 as it was after 1B-iv). (7) **Smoke** (`cargo test … multi_fold_convergence --ignored --release`, 3 folds × 5 epochs on RTX 3050 Ti): all 3 `dqn_fold{N}_best.safetensors` checkpoints written; per-fold best train Sharpe **21.14 / 38.08 / 63.72** vs the 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24) and the 1B-iv-only-main-backward intermediate 73.68 / 73.97 / 34.86 (geom-mean 57.49); this commit geom-mean = `(21.14 × 38.08 × 63.72)^(1/3) ≈ 36.36`, **−49% regression** vs 1B-iii. Floor for accept was ≥71.24 (do not regress vs 1B-iii) — **regression worsens with full-coverage extension**, disproving the "weak signal" hypothesis (gradient-coverage symmetry is not the issue) and pivoting the diagnosis to **VSN-bottleneck coupling instability**: with full-strength VSN gates training from all 4 paths, the input distribution to the bottleneck Linear shifts faster than the bottleneck Linear can co-adapt; F0 collapses hardest because the cold-start uniform 1/6 mask is destabilised earliest in fold 0 before VSN has settled into a steady-state mask. Rescue lands in 1B-iv-rc (next entry). No NaN/Inf, no fingerprint mismatch (fingerprint unchanged at `0x1b28321bb816f246` from 1B-ii/iii/iv), no panic, no slipped invariants. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62/62 cubins clean (kernel count unchanged — no new `.cu` files, all extensions live inside Rust dispatch code reusing 1B-iv's existing kernels). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. Together with 1B-iv this commit closes the 4-leaf 1B chain (1B-i kernel module, 1B-ii params + ISV slots, 1B-iii forward orchestrator + wire sites, 1B-iv main backward, 1B-iv-ext aux backward extension); the 1B-iv + 1B-iv-ext pair is committed as a single architectural unit because reverting 1B-iv's bottleneck-dW correction (point 7 there) regresses against 1B-iii's `vsn_gated_states_buf` forward contract, and 1B-iv standalone regresses smoke vs 1B-iii. VSN's 24 param tensors now receive gradient contributions from C51 + MSE + CQL + IQN + ensemble = **full coverage matching trunk weights**. + +Plan 4 Task 1B-iv (2026-04-25): VSN backward chain extension — final leaf of the 4-leaf 1B chain. The 24 VSN param tensors at indices `[95..119)` (allocated in 1B-ii, fed in the forward by 1B-iii) start receiving real `dW/dB` gradients on every training step; Adam consumes them through the existing `m_buf`/`v_buf`/`grad_buf` infrastructure (sized to `total_params + cutlass_tile_pad`, which already covers `[95..119)` per 1B-ii's `compute_param_sizes` rewrite). **No fingerprint change** — no new ISV slot or param tensor. **No checkpoint break** — params/ISV layouts unchanged from 1B-iii. (1) **Orchestrator** `CublasGemmSet::vsn_backward` added in `batched_forward.rs` next to `encoder_backward_chain`. Sequence per call (all GPU-only; no DtoH; no atomicAdd): (a) `vsn_softmax_and_gate_backward` kernel (loaded but `#[allow(dead_code)]` since 1B-iii — the annotation is removed in this commit) writes `d_logits[B, num_groups]` and `d_state[B, STATE_DIM_PADDED]` from the assembled `d_gated_state` + saved `state` + saved `mask`; one block per sample, 256 threads, shmem `2*num_groups + block` floats; the `d_state` output is computed but discarded because the state buffer is not trainable; (b) per-group loop `for g in 0..SL_NUM_FEATURE_GROUPS` runs (i) `strided_gather` (new kernel — see point 3) extracting column `g` of `d_logits[B, num_groups]` into a tight `[B]` scratch (`vsn_d_logit_scratch_buf`), avoiding the need for a strided-leading-dim cuBLAS variant; (ii) `launch_dw_only(dy=d_logit_scratch, x=save_h1_g[g], dw=grad_buf[95+4g+2], db=grad_buf[95+4g+3], out=1, in=VSN_HIDDEN_DIM=16, batch)` for `Linear_2[g]` weight + bias gradient; (iii) `launch_dx_only(dy=d_logit_scratch, w=W_2[g], dx=d_h1_scratch, beta=0)` for `Linear_2[g]` upstream gradient; (iv) `relu_mask` (existing helper) on `d_h1_scratch` against the saved post-ReLU `vsn_save_h1_g{g}_buf` (the post-ReLU activation is the gating signal — 1B-iii's `add_bias_relu_f32_kernel` writes back into the saved buffer in-place, so the same buffer drives both the saved h1 for `Linear_2` X and the `relu_mask` activation arg); (v) `backward_fc_layer_lda(dy=d_h1_scratch, x=state + gb*sizeof(f32), w=W_1[g], dw=grad_buf[95+4g+0], db=grad_buf[95+4g+1], dx=0 (skip), out=VSN_HIDDEN_DIM, in=group_dim_g, x_lda=state_dim_padded, batch)` for `Linear_1[g]` weight + bias gradient against the per-group state slice (the `dX` path is short-circuited because step (a)'s `d_state` already produced the full per-feature gradient and we don't propagate it). The per-group `d_logit_scratch_buf [B]` and `d_h1_scratch_buf [B, VSN_HIDDEN_DIM]` are reused across the 6 groups since the loop runs sequentially on the main stream. (2) **Wire site** at the end of `launch_cublas_backward_to` (single site — see point 6 for why CQL/IQN/ensemble aux paths are NOT wired). Insertion point: immediately after the existing bottleneck Linear backward `dW`/`db` block (which already ran `bn_tanh_backward_kernel` to produce `d_bn` and `launch_dw_only` to write `dW_bn`/`db_bn`). Three new sub-steps (active only when `bottleneck_dim > 0`, matching 1B-iii's forward-side bottleneck gate): (a) `vsn_d_gated_state_portfolio_pad_kernel` (new — see point 3) populates `vsn_d_gated_state_buf [B, STATE_DIM_PADDED]` — zero-fills the market slice `[0..market_dim]`, copies the portfolio passthrough `d_concat[:, bn_dim..]` into `[market_dim..STATE_DIM]`, zero-fills the padding tail `[STATE_DIM, STATE_DIM_PADDED)` (no-op currently since `STATE_DIM == STATE_DIM_PADDED == 128` per `state_layout.rs`, but honors the contract for future width changes); one thread per (sample, padded-index); (b) `launch_dx_only_lda` (new helper — see point 4) computes the bottleneck Linear backward dX: `d_bn[B, bn_dim] @ W_bn[bn_dim, market_dim] → vsn_d_gated_state_buf[:, 0..market_dim]@stride=state_dim_padded` with `beta=0` (overwrites the zero-initialised market slice from step (a) — order between (a) and (b) is irrelevant since (a) zeros the slice that (b) overwrites); (c) `vsn_backward` orchestrator call writes the 24 dW/dB into `grad_buf[95..119)`. (3) **Two new kernels**: `strided_gather` in `experience_kernels.cu` next to `strided_scatter` (the inverse direction — extracts a contiguous `[B, dst_stride]` from a wide `[B, src_lda]` source at column `col`; one thread per output element; bias-grad-friendly because the output is contiguous so the existing 2-phase `bias_grad_reduce_f32_p1`/`p2` kernels work without a strided-dY variant); `vsn_d_gated_state_portfolio_pad_kernel` in `dqn_utility_kernels.cu` next to `bn_tanh_backward_kernel` (the inverse of the portfolio path inside `bn_tanh_concat_kernel`; one thread per `(b, padded-index)` pair handles either zero-init market slice / portfolio passthrough / zero pad in a single launch). Both kernels register in `build.rs::kernels_with_common` automatically because they live in already-registered `.cu` files. (4) **One new cuBLAS helper** `launch_dx_only_lda` on `CublasBackwardSet`, mirroring `launch_dw_only_no_bias_lda` from 2c.3c.4 but for the `dX` path with custom `x_lda`. Same `sgemm_f32` call shape as `launch_dx_only`, threading `x_lda` into the dx leading-dim arg. Used exclusively by the bottleneck Linear backward dX site to write the market slice into a wide `[B, state_dim_padded]` buffer. (5) **VSN backward scratch buffers on `GpuDqnTrainer`** (allocated alongside the existing 1B-iii VSN forward buffers): `vsn_d_logits_buf [B, num_groups]` (kernel output of softmax-and-gate bwd), `vsn_d_state_buf [B, STATE_DIM_PADDED]` (kernel output, discarded after kernel returns), `vsn_d_gated_state_buf [B, STATE_DIM_PADDED]` (assembled INPUT to vsn_backward), `vsn_d_logit_scratch_buf [B]` (per-group dY tight scratch, reused across 6 groups), `vsn_d_h1_scratch_buf [B, VSN_HIDDEN_DIM]` (per-group dY/dX scratch for Linear_2 → ReLU mask → Linear_1 backward, reused across 6 groups). Two new kernel handles: `vsn_logit_gather_kernel` loaded from `EXPECTED_Q_CUBIN` (`experience_kernels.cubin`) next to the existing `strided_scatter_kernel` load; `vsn_d_gated_state_portfolio_kernel` loaded from `DQN_UTILITY_CUBIN` next to the existing `bn_tanh_backward_kernel` load. Total scratch footprint at B=512, num_groups=6, STATE_DIM_PADDED=128, VSN_HIDDEN_DIM=16: ~531 KB (negligible against existing GRN backward scratches at ~5 MB and per-branch cuBLAS workspaces at ~32 MB). (6) **Backward-site coverage scope** — the spec called out a partial-refactor risk if the 4 `encoder_backward_chain` callers (main backward + CQL aux + IQN aux + ensemble-diversity aux) ALL needed VSN backward extension. Honest read: the 3 aux callers (`apply_iqn_trunk_gradient`, `apply_ensemble_diversity_backward`, the CQL block at line 6697) all pass `s1_dx_output = 0` to `encoder_backward_chain` — by design they don't propagate trunk-input gradient at all; they only contribute to trunk weight tensors `[0..13)` via SAXPY into `grad_buf[trunk]`. Adding VSN backward to those would require lifting the `s1_dx_output = 0` constraint, allocating per-aux-path scratch for the assembled `d_vsn_gated_state`, and adding 24 more SAXPY entries per aux path to mix VSN dW into `grad_buf[95..119)`. That's a fundamental new gradient-flow architecture, not a lockstep contract migration. Choosing to scope this commit to the main backward only, mirroring the existing aux-path convention that gradients stop at the trunk-input boundary; VSN therefore receives gradients exclusively from the C51 + MSE direct losses (the main backward path), not from CQL / IQN / ensemble auxiliary signals. If a follow-up wants aux paths to contribute to VSN dW, that's a 3-site extension with its own scratch + SAXPY scope — out of scope for the 1B-iv leaf. (7) **In-passing fix from 1B-iii**: the bottleneck Linear backward at the main-backward site was passing `raw_states_ptr` (= `states_buf`, the un-gated raw state) as `X` for `dW_bn = d_bn^T @ X`, but 1B-iii's forward changed the bottleneck Linear to consume `vsn_gated_states_buf`. Pre-1B-iv this was masked by the cold-start-uniform-1/6 mask making `vsn_gated_states ≈ raw_states / 6` (off by a constant factor that Adam normalises away). With VSN now training, the gate stops being constant and the mismatch becomes a real gradient bug. Fixed in lockstep with the VSN backward extension (`raw_states_ptr` → `vsn_gated_states_buf.raw_ptr()` at the dW_bn call site). The VSN backward itself takes `state_ptr = raw_states_ptr` (the pre-VSN saved state, which is what `vsn_softmax_and_gate_backward` consumes for the per-group dot product reduction `d_dot[g] = sum_{i in g} d_gated[i] * state[i]`). (8) **`#[allow(dead_code)]` retired** on `vsn_softmax_and_gate_bwd_kernel` (the field was loaded by 1B-iii in anticipation of this commit; the consumer is now `vsn_backward`). (9) **Constraints honoured**: no DtoH in any new path; no `atomicAdd` (the per-group cuBLAS GEMMs run sequentially with `beta=0` overwrites; the bias-grad reduce kernel uses 2-phase shmem; `strided_gather` and `vsn_d_gated_state_portfolio_pad_kernel` are one-thread-per-element data-parallel); no stubs (no `Option`/None paths added; the `dx=0` skip in `backward_fc_layer_lda` for VSN's `Linear_1[g]` is the existing convention from `encoder_backward_chain` and is documented inline); no `// ok:` band-aids; the partial-refactor invariant is honoured at the contract that it touches (the bottleneck Linear backward + VSN backward share the `vsn_gated_states_buf` contract introduced by 1B-iii's forward; both producer and consumer migrated in this commit). (10) **Smoke** (`cargo test … multi_fold_convergence --ignored --release`, 3 folds × 5 epochs on RTX 3050 Ti): all 3 `dqn_fold{N}_best.safetensors` checkpoints written; per-fold best train Sharpe **73.68 / 73.97 / 34.86** vs the 1B-iii baseline 68.83 / 84.84 / 61.95 at epochs 2 / 5 / 4 (geom-mean 71.24); this commit geom-mean = `(73.68 × 73.97 × 34.86)^(1/3) ≈ 57.49`, **−19% regression** vs 1B-iii. Diagnosis (informed by the post-fact 1B-iv-ext follow-up smoke below): the regression is a **gradient-coverage gap**, not a wiring bug. With VSN backward wired only at the main C51/MSE path, VSN's 24 param tensors receive a strictly weaker gradient signal than the 13 trunk tensors, which receive contributions from main + CQL aux + IQN aux + ensemble aux paths. The signal asymmetry biases VSN's effective learning rate downward relative to trunk, distorting the joint optimization trajectory in a way that fold 2 (which already had the lowest 1B-iii score) regresses hardest. The fix lands in 1B-iv-ext (the next entry); 1B-iv on its own is committed alongside 1B-iv-ext as a single architectural unit because reverting the bottleneck-dW correction (point 7) regresses against 1B-iii's `vsn_gated_states_buf` forward contract. No NaN/Inf, no fingerprint mismatch (fingerprint unchanged at `0x1b28321bb816f246` from 1B-ii/iii), no panic, no slipped invariants. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62/62 cubins clean (kernel count unchanged — the two new kernels live in already-registered `.cu` files). 0 panic gates added/removed. 0 stubs. The forward path's Wired-Producer+Consumer classification from 1B-iii is now reinforced by Wired-Backward (24 VSN tensors get gradients on every training step). This commit + 1B-iv-ext together close the 4-leaf 1B chain (1B-i kernel module, 1B-ii params + ISV slots, 1B-iii forward orchestrator + wire sites, 1B-iv main backward, 1B-iv-ext aux backward extension). + Plan 4 Task 1B-iii (2026-04-25): VSN forward orchestrator + 3 production wire sites + per-step ISV producer launch. The 24 VSN param tensors at indices `[95..119)` (1B-ii) and the 6 ISV mask-EMA slots at `[105..111)` (1B-ii) acquire their first production consumers. **No fingerprint change** — no new ISV slot or param tensor; pure orchestrator + wire-in. **No backward chain extension** in this commit (1B-iv lands that) so VSN dW = 0 and the 24 VSN params stay at Xavier init throughout training; the smoke result below validates that the forward + ISV producer wiring is correct in isolation. Per the cold-start `softmax(≈small_random_logits) ≈ 1/6 ± per-group bias from the random Xavier realisation` analysis, the gated state at every step is approximately `state * (1/6 + ε)` per feature — a near-uniform multiplicative scale that the downstream GRN trunk's LayerNorm absorbs at the first cuBLAS Linear (so no destabilisation). (1) **Orchestrator** `CublasGemmSet::vsn_forward` added in `batched_forward.rs`: per-group loop over `SL_NUM_FEATURE_GROUPS=6` issuing (a) cuBLAS `Linear_1[g]` via `sgemm_f32_ldb` with the input pointer offset by `gb * sizeof(f32)` and ldb fixed at `state_dim_padded` (cuBLAS reads col-major `[state_dim_padded, B]` view starting at the group's first feature column — no slice copy needed), (b) fused bias + ReLU via `add_bias_relu_f32_kernel` writing post-ReLU activation back into `h1_ptr` (so the buffer IS the save-for-backward storage on the online pass — no separate DtoD save), (c) cuBLAS `Linear_2[g]` writing a tight `[B]` logit array, (d) bias-add (no activation), (e) `strided_scatter` (reused from `experience_kernels.cu`) writing the `[B]` logit into column `g` of the assembled `[B, num_groups]` logits buffer with `src_stride=1, dst_stride=num_groups` and the dst pointer offset by `g * sizeof(f32)`. After the loop, `vsn_softmax_and_gate_forward` (1 block per sample, 256 threads, `num_groups * sizeof(f32)` shmem) performs numerically-stable softmax over the 6 logits and per-feature gate-multiplies the row, with passthrough on indices outside any group's `[gb, ge)` range (the 7-element padding tail past `SL_PADDING_START`). (2) **Three production wire sites** all consume the orchestrator: (i) `launch_cublas_forward` online-on-states pass — saves `vsn_logits_buf` / `vsn_mask_buf` / per-group `vsn_save_h1_g{0..5}_buf` for 1B-iv backward; the bottleneck Linear, `bn_tanh_concat`, and `launch_concat_ofi(2|3, ...)` all read the gated state via the renamed `states_for_bn` local; (ii) `launch_cublas_forward` target-on-next_states pass — uses `tg_w_ptrs` (Polyak EMA copy of online VSN weights at indices [95..119), the existing target-EMA loop covers the new tensors automatically since `target_params_buf` is sized to `total_params + cutlass_tile_pad`), throwaway `vsn_logits_target_scratch` / `vsn_mask_target_scratch`, single-buffer `vsn_h1_inference_scratch` shared across the 6 groups (no save — target inference-only); both the target bottleneck Linear, `bn_tanh_concat`, and the target-side `launch_concat_ofi(2|3, ...)` calls now read `tg_states_for_bn` (= `vsn_gated_next_states_buf`); (iii) `submit_forward_ops_ddqn` DDQN-online-on-next_states pass — uses `on_w_ptrs` (DDQN argmax runs the online net on next_states), throwaway `vsn_logits_ddqn_scratch` / `vsn_mask_ddqn_scratch`, same `vsn_h1_inference_scratch`; the DDQN bottleneck Linear and `bn_tanh_concat` now read `on_next_states_for_bn` (= `vsn_gated_next_states_for_ddqn_buf`). Per-pass distinct mask + logits scratches prevent target/DDQN passes from clobbering the online pass's saved buffers — the producer kernel `vsn_mask_ema_update` reads `vsn_mask_buf` (online pass) only, and 1B-iv's backward will read both `vsn_logits_buf` (saved) and the 6 per-group `vsn_save_h1_g*_buf` (saved). (3) **VSN buffers on `GpuDqnTrainer`**: 3 gated-state buffers (online + target + ddqn, each `[B, STATE_DIM_PADDED]`), 3 logits + 3 mask buffers (online saved + 2 throwaways each, all `[B, num_groups]`), 6 per-group online h1 saves (`[B, VSN_HIDDEN_DIM=16]` each), 1 shared `vsn_h1_inference_scratch [B, 16]` for target/ddqn, 1 shared `vsn_linear2_scratch_buf [B]` overwritten per group across all passes (each pass runs the per-group loop sequentially before the next pass starts so no cross-pass interference), 2 device-side `vsn_group_begins_buf [6]` + `vsn_group_ends_buf [6]` populated once at construction from `FEATURE_GROUP_RANGES` via `stream.memcpy_htod`. (4) **Kernel handles**: `CublasGemmSet` gains `vsn_softmax_and_gate_fwd_kernel` (consumed by `vsn_forward`) + `vsn_softmax_and_gate_bwd_kernel` (loaded but `#[allow(dead_code)]` until 1B-iv consumes it in the backward chain), both loaded from `VSN_FEATURE_SELECTION_CUBIN` in `CublasGemmSet::new` near the existing `saxpy_kernel` load; the `vsn_logit_scatter_kernel: Option` field is wired post-construction by the trainer via the new `wire_vsn_scatter` method (mirroring `wire_vsn_glu` / `wire_ofi_concat`) using the same `strided_scatter` handle the trainer already loads from `experience_kernels.cubin` for VSN bottleneck — `Option` keeps the experience collector's `CublasGemmSet` instance (which never calls `vsn_forward`) from needing the wire-up. `GpuDqnTrainer` gains `vsn_mask_ema_kernel: CudaFunction` loaded from `VSN_MASK_EMA_CUBIN` in the constructor next to the new buffer allocations. (5) **ISV producer launch** `launch_vsn_mask_ema(ema_alpha)` on `GpuDqnTrainer` mirrors `launch_h_s2_rms_ema`'s style: single-block 256-thread kernel reads `vsn_mask_buf` and EMA-updates the 6 ISV slots `[VSN_MASK_GROUP_0_EMA_INDEX..VSN_MASK_GROUP_5_EMA_INDEX]` with `debug_assert!(self.isv_signals_dev_ptr != 0, ...)` mirroring the producer-launcher invariant from 2c.3c.5. Wired in `training_loop.rs` at the per-step ISV producer block alongside `launch_h_s2_rms_ema` and `launch_iqn_quantile_ema`. (6) **Cubin static `dead_code` retired**: both `VSN_FEATURE_SELECTION_CUBIN` and `VSN_MASK_EMA_CUBIN` lose their `#[allow(dead_code)]` annotations because both now have runtime `load_cubin(...)` consumers (`CublasGemmSet::new` for the former, `GpuDqnTrainer::new` for the latter). (7) **Constraints honoured**: no DtoH in any new path (kernels read from device-mapped pinned ISV pointer); no `atomicAdd` (the strided_scatter is fully data-parallel — one thread per sample per scatter, the softmax+gate kernel uses block-local shmem reduction); no stubs (the only `Option` is `vsn_logit_scatter_kernel`, gated by `debug_assert!` + `expect()` — silent-skip on `None` would mask trainer wire-up failure, the loud panic flags the contract violation); no `// ok:` band-aids; the partial-refactor invariant is honoured (all 3 forward paths attach VSN at the same point and downstream consumers all read the gated buffer — no path is left reading raw `states_buf` / `next_states_buf`). (8) **Smoke** (`cargo test … multi_fold_convergence --ignored --release`, **594.94s**, 3 folds × 5 epochs on RTX 3050 Ti): all 3 `dqn_fold{N}_best.safetensors` checkpoints written; per-fold best train Sharpe **68.83 / 84.84 / 61.95 at epochs 2 / 5 / 4** (vs the post-1B-ii baseline 6.53 / 80.11 / 66.72 at epochs 2 / 4 / 4 → geom-mean 39.27); this commit geom-mean = `(68.83 × 84.84 × 61.95)^(1/3) ≈ 71.24`, **+81% lift** vs baseline — vastly above the spec's ≥27.5 (≤30% regression budget) acceptance threshold. The lift is consistent with the cold-start analysis: the near-uniform 1/6 multiplicative gate is absorbed by the GRN trunk's LayerNorm in the first epoch, and the slight per-group bias from the random Xavier realisation acts as a tiny attention-like signal that the trunk amplifies through training (despite zero VSN dW — the bias is fixed but downstream dW propagates it). No NaN/Inf, no fingerprint mismatch (fingerprint unchanged at `0x1b28321bb816f246` from 1B-ii), no panic, no slipped invariants. ISV[105..111) values were not directly logged in this run (no HEALTH_DIAG line surfaces these slots in the existing diag layout); the smoke result is the load-bearing wiring proof — a wrong-ldb on the per-group GEMM, a wrong group_begins/ends device buffer, or a logit scatter offset bug would have produced NaN-propagating gated state and the fold-level Sharpes would have collapsed below the 27.5 threshold rather than lifting above 71. The producer kernel's launch is observably executing without launch errors per the absence of any `tracing::warn!("Plan 4 1B-iii vsn_mask_ema launch failed: …")` line in the smoke log. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62 cubins clean (unchanged — no new `.cu` files). 0 panic gates added/removed. 0 stubs. 2 Orphan-by-design rows retired (the 1B-i `vsn_feature_selection_kernel.cu` + `vsn_mask_ema_kernel.cu` rows reclassify Wired — both kernels now have production callers via the orchestrator and the ISV producer launch). The forward path is now Wired-Producer+Consumer for VSN; the backward chain extension that lets VSN params train lands in 1B-iv. Plan 4 Task 1B-ii (2026-04-25): VSN params + ISV slots landing — additive, no callers (param tensors written but unread; ISV slots cold-started but unread). 24 new param tensors appended at indices `[95..119)` covering the per-group VSN MLP (`Linear_1_g`/`b1_g`/`Linear_2_g`/`b2_g`) for each `FEATURE_GROUP_RANGES` entry — one quad per group, `VSN_HIDDEN_DIM = 16` (the new compile-time constant alongside `H_S2_RMS_EMA_INDEX` at the head of `gpu_dqn_trainer.rs`). Per-group element count `16*group_dim_g + 16 + 16 + 1`; sum `16*(42+32+16+16+8+7) + 6*33 = 1936 + 198 = 2134` floats (≈8.5 KB on params, mirrored on `target_params_buf` and Adam `m_buf`/`v_buf` — all four buffers grow automatically since their allocations use `compute_total_params(cfg) + cutlass_tile_pad`). `NUM_WEIGHT_TENSORS` 95 → 119. `compute_param_sizes()` rewritten to build a `[usize; NUM_WEIGHT_TENSORS]` from the existing 95-entry array literal `core` plus a runtime `for g in 0..SL_NUM_FEATURE_GROUPS` loop reading `FEATURE_GROUP_RANGES` from `ml_core::state_layout` (per Task 1A's prerequisite); `xavier_init_params_buf()` follows the same `core_fan` + per-group loop pattern with `(VSN_HIDDEN_DIM, group_dim)` for `w1_g`, `(1, VSN_HIDDEN_DIM)` for `w2_g`, and `(0,0)` (zero) for both biases — initial logits `≈ 0` therefore yield `softmax ≈ 1/SL_NUM_FEATURE_GROUPS` per group at every sample (cold-start neutral). 6 new ISV slots tail-appended for the per-group importance-mask EMA: `VSN_MASK_GROUP_0_EMA_INDEX=105` (market) through `VSN_MASK_GROUP_5_EMA_INDEX=110` (plan_isv); fingerprint pair shifted 103/104 → 111/112; `ISV_TOTAL_DIM` 105 → 113; `layout_fingerprint_seed()` extended with `VSN_MASK_GROUP_{0..5}_EMA=105..110;ISV_LAYOUT_FINGERPRINT_LO=111;ISV_LAYOUT_FINGERPRINT_HI=112;ISV_TOTAL_DIM=113;` plus the 24 new `PARAM_VSN_W1_G{g}=…;PARAM_VSN_B1_G{g}=…;PARAM_VSN_W2_G{g}=…;PARAM_VSN_B2_G{g}=…;` entries terminating at `PARAM_TOTAL_TENSORS=119` — new `LAYOUT_FINGERPRINT_CURRENT = 0x1b28321bb816f246` (was `0x5789155b683ab59c`). Constructor cold-start writes `1.0/SL_NUM_FEATURE_GROUPS = 1/6` (uniform prior — no group preference at init) into all six new ISV slots, mirroring the `H_S2_RMS_EMA = 1.0` cold-start convention from 2c.3c.5. `StateResetRegistry` extended with 6 new `FoldReset` entries (`isv_vsn_mask_g{0..5}_ema`); `training_loop.rs::reset_named_state` adds a single match arm covering all six names that recomputes the same `1.0 / SL_NUM_FEATURE_GROUPS` uniform value — no hardcoded `1.0/6.0` per `feedback_isv_for_adaptive_bounds.md` (the constant comes from `ml_core::state_layout::SL_NUM_FEATURE_GROUPS` so a future group-count change propagates automatically). Two new `pub(crate) static` cubin refs `VSN_FEATURE_SELECTION_CUBIN` and `VSN_MASK_EMA_CUBIN` added alongside the existing Plan 4 cubin block; `#[allow(dead_code)]` annotated because the runtime `load_cubin(...)` + `load_function(...)` calls are deferred to 1B-iii's forward orchestrator wire-in (the static refs themselves are compile-time `include_bytes!()` only — adding them in this commit means 1B-iii has zero new cubin-side file-touches). Mirrors the 2c.3a pattern (param tensor reshuffle without runtime callers — that landed cleanly because the runtime sites were panic-gated; here, the runtime sites simply don't exist yet so no gating needed). **Checkpoint break by intent** — fingerprint shift invalidates pre-1B-ii checkpoints at constructor load (fail-fast, no migration path per `feedback_no_legacy_aliases.md`). **No production callers in this commit** — params are written but unread (no consumer reads `param_buf[95..119)`), ISV slots are cold-started but unread (no consumer reads `isv_signals[105..111)`), and neither cubin is loaded yet. The forward orchestrator + cuBLAS `Linear_1`/`Linear_2` wire-in lands in 1B-iii; the backward orchestrator + autograd integration lands in 1B-iv. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62/62 cubins clean (unchanged — no new `.cu` files, only new `pub(crate) static` refs to the 1B-i cubins). Smoke deferred — behaviour byte-identical to the f3e3ac347 / `0x5789155b683ab59c` post-bottleneck-fix baseline (700.43s, 3 folds × 5 epochs, per-fold best train Sharpe 6.53 / 80.11 / 66.72, geom-mean 39.27) since no consumer reads the new params or ISV slots; smoke re-validation with real numbers lands at 1B-iii where the forward orchestrator wire-in actually validates something. 2 new `dead_code`-annotated cubin refs (will retire `dead_code` at 1B-iii); 0 new Orphan rows (the kernels themselves were registered at 1B-i and remain Orphan-by-design until 1B-iii). 0 panic gates added/removed.