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

fix(sp4): #260 — eliminate SP1-era 1e6 × isv multipliers via ISV-driven producers

Browse Source 
Two SP1-Phase-C cuBLAS backward sanitisers at gpu_dqn_trainer.rs:7615
(bw_d_h_s2 input) and :20679 (d_value_logits + d_adv_logits) used
hardcoded `1e6 × signal.max(1.0)` formulas — outside SP4 mechanisms
1/2/5/6/9/10 but flagged by feedback_isv_for_adaptive_bounds review as
the last remaining hardcoded multipliers in the cuBLAS-backward
sanitiser consumer path after Layer A/B closed out the rest of SP4.

Migrated to the proper SP4 pattern:
  - 2 new SP4 ISV slots [171, 172) extending [131..171) → [131..173):
    BW_D_H_S2_BOUND_INDEX=171 (single-buffer p99) and
    Q_DIR_GRAD_BOUND_INDEX=172 (multi-sub-buffer p99 over
    d_value_logits ∪ d_adv_logits, n_sub=2).
  - 2 new producer kernels:
    bw_d_h_s2_p99_kernel.cu (mirrors h_s2_p99 single-buffer pattern)
    q_dir_grad_p99_kernel.cu (mirrors param_group_oracle multi-sub
    convention via shared sp4_histogram_p99_multi<256> template;
    reuses oracle_subbuf_table_buf + oracle_subbuf_counts_buf).
  - Pearls A+D wired via existing apply_pearls_ad_kernel chained on
    same stream (GPU-only, graph-capture-compatible).
  - Consumers read ISV[slot].max(EPS_CLAMP_FLOOR) directly.

Buffer growth (single source of truth in gpu_dqn_trainer.rs):
  SP4_PRODUCER_COUNT 69 → 71, SP4_WIENER_TOTAL_FLOATS 207 → 213.
  ISV_TOTAL_DIM 171 → 173. LAYOUT_FINGERPRINT_SEED extended (existing
  checkpoints will fail-fast at load — re-train required).

StateResetRegistry: 2 new FoldReset entries (sp4_bw_d_h_s2_bound,
sp4_q_dir_grad_bound) + 2 matching reset_named_state dispatch arms;
both halves of Pearl A's sentinel contract reset together per
feedback_no_partial_refactor.

Unit tests: 2 new GPU-gated tests in sp4_producer_unit_tests.rs
exercising (single-buffer Pearl-A→Pearl-D convergence, multi-sub-
buffer p99 vs analytical |N(0,1)|). All 16 GPU tests pass on local
RTX 3050 Ti (rel-err 0.00000 / 0.00526).

Behaviour change: bound is now Pearls-smoothed p99 of the actual
gradient distribution, tighter than the pre-existing 1e6 ceiling but
appropriate-scale for observed values. Smoke validation must verify
F0/F1/F2 still converge.

Refs: task #260, baseline commit 1c150d190 (Layer A + B + fix-up +
L40S smoke pass).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 12:41:29 +02:00
parent 1c150d1901
commit 88ae74ca73
9 changed files with 1078 additions and 101 deletions
Download Patch File Download Diff File Expand all files Collapse all files

30
crates/ml/build.rs
View File

@@ -306,6 +306,36 @@ fn main() {
// (slot 96) which tracks RMS — different signal. Task A13 will
// retrofit that existing kernel to also use Pearls A+D, separately.
"h_s2_p99_kernel.cu",
// SP4 Layer C #260 (2026-05-01): bw_d_h_s2_p99 producer for
// ISV[BW_D_H_S2_BOUND_INDEX=171]. Single-block 256-thread kernel
// reading `bw_d_h_s2 [B × SH2]` (the IQN trunk-backward d_h_s2 buffer
// post-DtoD copy from `iqn_d_h_s2_ptr` and post slot-35 NaN check)
// via `sp4_histogram_p99<256>`; output → `producer_step_scratch_buf[69]`
// with `__threadfence_system()`. The host launcher
// (`GpuDqnTrainer::launch_sp4_bw_d_h_s2_p99`) chains the GPU
// `apply_pearls_ad_kernel` on the same stream — writes new x_mean to
// ISV[171] + Wiener state at slot 69 (offset 120). Migrated from
// SP1-Phase-C `1e6 × ISV[H_S2_RMS_EMA].max(1.0)` per
// `feedback_isv_for_adaptive_bounds`. Mirrors `h_s2_p99_kernel.cu`'s
// shape exactly — only differs in the consumed buffer + ISV slot.
"bw_d_h_s2_p99_kernel.cu",
// SP4 Layer C #260 (2026-05-01): q_dir_grad_p99 producer for
// ISV[Q_DIR_GRAD_BOUND_INDEX=172]. Single-block 256-thread kernel
// reading the union of `d_value_logits` ∪ `d_adv_logits` (both
// dueling-head logits gradient buffers) via `sp4_histogram_p99_multi<256>`
// with `n_sub=2`; output → `producer_step_scratch_buf[70]` with
// `__threadfence_system()`. The host launcher
// (`GpuDqnTrainer::launch_sp4_q_dir_grad_p99`) builds a 2-element
// device-pointer table + 2-element count table in mapped-pinned memory
// (reusing the same `oracle_subbuf_table_buf` allocated for
// `param_group_oracle_kernel.cu`), then chains the GPU
// `apply_pearls_ad_kernel` on the same stream — writes new x_mean to
// ISV[172] + Wiener state at slot 70 (offset 123). Migrated from
// SP1-Phase-C `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)` per
// `feedback_isv_for_adaptive_bounds`. The same bound was previously
// applied to BOTH d_value_logits and d_adv_logits, so the union p99
// is the natural ISV-driven replacement.
"q_dir_grad_p99_kernel.cu",
// SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread
// device-side loop applies Pearls A+D to N consecutive ISV slots
// (each with its own Wiener-state triple) on the producer's

44
crates/ml/src/cuda_pipeline/bw_d_h_s2_p99_kernel.cu Normal file
View File

@@ -0,0 +1,44 @@
// crates/ml/src/cuda_pipeline/bw_d_h_s2_p99_kernel.cu
//
// SP4 Layer C #260 (2026-05-01) — bw_d_h_s2_p99 producer for
// ISV[BW_D_H_S2_BOUND_INDEX=171].
//
// Single-block, 256-thread kernel that `#include "sp4_histogram_p99.cuh"`
// directly, reads `bw_d_h_s2 [B × SH2]` (the IQN trunk-backward d_h_s2 buffer
// post-DtoD copy from `iqn_d_h_s2_ptr` and post slot-35 NaN check), computes
// p99(|bw_d_h_s2|) via `sp4_histogram_p99<256>`, and writes the per-step
// observation to `producer_step_scratch_buf[scratch_idx]` with
// `__threadfence_system()` so the chained `apply_pearls_ad_kernel` reads the
// fresh value.
//
// Migrated from SP1-Phase-C hardcoded formula
// `1e6 × ISV[H_S2_RMS_EMA_INDEX=96].max(1.0)` per
// `feedback_isv_for_adaptive_bounds`. The hardcoded multiplier was the last
// remaining `1e6 × isv` constant in the cuBLAS backward-grad sanitiser path
// after Layer A/B closed out the rest of SP4.
//
// The host launcher (`GpuDqnTrainer::launch_sp4_bw_d_h_s2_p99`) follows up
// with the GPU `apply_pearls_ad_kernel` on the same stream which writes the
// new x_mean to ISV[BW_D_H_S2_BOUND_INDEX=171] + Wiener state at slot 69
// (offset `(171 - SP4_SLOT_BASE) * 3 = 120`). Graph-capture-compatible (no
// host sync; same-stream ordering with the producer kernel).
//
// Mirrors `h_s2_p99_kernel.cu` exactly — single-buffer histogram p99
// producer; the only difference is the input pointer (`bw_d_h_s2` vs
// `save_h_s2`) and the consumer ISV slot/scratch slot.
#include "sp4_histogram_p99.cuh"
extern "C" __global__ void bw_d_h_s2_p99_update(
const float* __restrict__ buf,
int count,
float* __restrict__ scratch_buf,
int scratch_idx
) {
if (blockIdx.x != 0) return;
float p99 = sp4_histogram_p99<256>(buf, count);
if (threadIdx.x == 0) {
scratch_buf[scratch_idx] = p99;
__threadfence_system(); // make host-mapped write visible
}
}

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

@@ -225,6 +225,40 @@ static SP4_GRAD_NORM_P99_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"),
/// migrates in Layer B.
static SP4_H_S2_P99_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/h_s2_p99_kernel.cubin"));
/// SP4 Layer C #260 (2026-05-01): bw_d_h_s2_p99 producer cubin for
/// ISV[BW_D_H_S2_BOUND_INDEX=171]. Single-block, 256-thread kernel reading
/// `bw_d_h_s2 [B × SH2]` (the IQN trunk-backward d_h_s2 buffer post-DtoD
/// copy and post slot-35 NaN check) via `sp4_histogram_p99<256>`; output →
/// `producer_step_scratch_buf[69]`. The launcher
/// (`GpuDqnTrainer::launch_sp4_bw_d_h_s2_p99`) chains the GPU
/// `apply_pearls_ad_kernel` on the same stream which writes the new x_mean
/// to ISV[BW_D_H_S2_BOUND_INDEX=171] + Wiener state at offset
/// `(171 - SP4_SLOT_BASE) * 3 = 120`. Migrated from SP1-Phase-C
/// `1e6 × ISV[H_S2_RMS_EMA].max(1.0)` hardcoded multiplier per
/// `feedback_isv_for_adaptive_bounds`. Consumer at
/// `apply_iqn_trunk_gradient` (gpu_dqn_trainer.rs:~7615) reads
/// `ISV[171].max(EPS_CLAMP_FLOOR)` directly.
static SP4_BW_D_H_S2_P99_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/bw_d_h_s2_p99_kernel.cubin"));
/// SP4 Layer C #260 (2026-05-01): q_dir_grad_p99 producer cubin for
/// ISV[Q_DIR_GRAD_BOUND_INDEX=172]. Single-block, 256-thread kernel reading
/// the union of `d_value_logits` ∪ `d_adv_logits` via
/// `sp4_histogram_p99_multi<256>` with n_sub=2; output →
/// `producer_step_scratch_buf[70]`. The launcher
/// (`GpuDqnTrainer::launch_sp4_q_dir_grad_p99`) builds a 2-element
/// device-pointer table + 2-element count table in mapped-pinned memory
/// (reusing `oracle_subbuf_table_buf`), then chains the GPU
/// `apply_pearls_ad_kernel` on the same stream which writes the new x_mean
/// to ISV[Q_DIR_GRAD_BOUND_INDEX=172] + Wiener state at offset
/// `(172 - SP4_SLOT_BASE) * 3 = 123`. Migrated from SP1-Phase-C
/// `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)` hardcoded multiplier per
/// `feedback_isv_for_adaptive_bounds`. Consumer at
/// `launch_cublas_backward_to` (gpu_dqn_trainer.rs:~20679) reads
/// `ISV[172].max(EPS_CLAMP_FLOOR)` directly.
static SP4_Q_DIR_GRAD_P99_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/q_dir_grad_p99_kernel.cubin"));
/// SP4 GPU-only Pearls A+D applicator cubin (2026-05-01). Loaded once by
/// `GpuDqnTrainer::new` and shared with the experience collector via the
/// `apply_pearls_ad_kernel_handle` accessor. Replaces the host-side
@@ -555,7 +589,7 @@ const ISV_NETWORK_DIM: usize = 23;
/// (shifted 112→116 in Plan 4 Task 6 Commit A).
/// 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 = 171; // SP4: 131 + 40 magnitude/EMA-rate/regularization slots
const ISV_TOTAL_DIM: usize = 173; // SP4: 131 + 42 magnitude/EMA-rate/regularization slots (40 Layer A/B + 2 Layer C #260)
/// Legacy alias preserved for call sites that haven't been audited for the
/// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight
/// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer).
@@ -742,19 +776,22 @@ pub const REWARD_COMPONENT_COUNT: usize = 6;
/// - 40 SP4 Layer A producers (Tasks A5-A9): TARGET_Q_BOUND, ATOM_POS×4,
/// WEIGHT/ADAM_M/ADAM_V/WD_RATE×8, L1_LAMBDA, GRAD_CLIP, H_S2 — slots 0..40
/// - 29 Task-A13 retrofit producers — slots 40..69
/// - 2 Layer C #260 producers (2026-05-01) — slots 69..71:
/// 69 = `bw_d_h_s2_p99` (single buffer) → ISV[171]
/// 70 = `q_dir_grad_p99` (multi-sub-buffer) → ISV[172]
///
/// Co-located with `SP4_WIENER_FLOATS_PER_SLOT` and `SP4_WIENER_TOTAL_FLOATS`
/// so the relationship `wiener_total = count × floats_per_slot` is visible at
/// declaration time. Re-exported via `cuda_pipeline::SP4_PRODUCER_COUNT`
/// for consumption by `crates/ml/tests/sp4_producer_unit_tests.rs`.
pub const SP4_PRODUCER_COUNT: usize = 69;
pub const SP4_PRODUCER_COUNT: usize = 71;
/// Per-slot Wiener-state float count (`sample_var`, `diff_var`, `x_lag`).
/// Matches `WienerState`'s flat layout in `sp4_wiener_ema.rs`.
pub const SP4_WIENER_FLOATS_PER_SLOT: usize = 3;
/// Total floats in the SP4 Wiener-state buffer (one triple per producer slot).
/// Equals `SP4_PRODUCER_COUNT × SP4_WIENER_FLOATS_PER_SLOT = 207`.
/// Equals `SP4_PRODUCER_COUNT × SP4_WIENER_FLOATS_PER_SLOT = 213`.
pub const SP4_WIENER_TOTAL_FLOATS: usize =
SP4_PRODUCER_COUNT * SP4_WIENER_FLOATS_PER_SLOT;
@@ -1179,7 +1216,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
WD_RATE_GROUP_0=160;WD_RATE_GROUP_1=161;WD_RATE_GROUP_2=162;WD_RATE_GROUP_3=163;\
WD_RATE_GROUP_4=164;WD_RATE_GROUP_5=165;WD_RATE_GROUP_6=166;WD_RATE_GROUP_7=167;\
GRAD_CLIP_BOUND=168;H_S2_BOUND=169;L1_LAMBDA_TRUNK=170;\
ISV_TOTAL_DIM=171;\
BW_D_H_S2_BOUND=171;Q_DIR_GRAD_BOUND=172;\
ISV_TOTAL_DIM=173;\
PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\
PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\
PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\
@@ -2937,9 +2975,13 @@ pub struct GpuDqnTrainer {
/// SP1 Phase C surgical fix: in-place finite clamp kernel
/// (`dqn_clamp_finite_f32_kernel` in `dqn_utility_kernels.cu`). Replaces
/// NaN/Inf with 0 and clamps finite |v| to ±max_abs. max_abs is sourced
/// from existing ISV slots (1e6 × ISV[H_S2_RMS_EMA_INDEX=96] for the
/// IQN trunk path; 1e6 × ISV[Q_DIR_ABS_REF_INDEX=21] for the
/// Bottleneck Linear path) with a 1e3 floor for uninitialised state.
/// from dedicated SP4 ISV bound slots:
/// - IQN trunk path (`bw_d_h_s2`): ISV[BW_D_H_S2_BOUND_INDEX=171]
/// (Layer C #260, producer `bw_d_h_s2_p99`)
/// - Bottleneck Linear path (`d_value_logits` ∪ `d_adv_logits`):
/// ISV[Q_DIR_GRAD_BOUND_INDEX=172] (Layer C #260, producer
/// `q_dir_grad_p99` multi-sub-buffer n_sub=2)
/// with `EPS_CLAMP_FLOOR=1.0` ε on cold-start (Invariant 1 carve-out).
/// Used as pre-call sanitisation on cuBLAS GEMM backward inputs and
/// post-call defence-in-depth on outputs at slots 26/32 per
/// `docs/dqn-backward-nan-audit.md`.
@@ -2951,13 +2993,16 @@ pub struct GpuDqnTrainer {
/// boundary or on training-guard halts.
pub(crate) nan_flags_buf: CudaSlice<i32>,
/// SP4 Pearl D Wiener-EMA state for all 69 ISV-bound EMAs (40 SP4 +
/// 29 Task-A13 retrofit existing producers). Per slot, 3 floats laid
/// out as `[sample_var, diff_var, x_lag]`. Mapped-pinned per
/// SP4 Pearl D Wiener-EMA state for all 71 ISV-bound EMAs (40 SP4 +
/// 29 Task-A13 retrofit + 2 Layer-C-#260 producers). Per slot, 3 floats
/// laid out as `[sample_var, diff_var, x_lag]`. Mapped-pinned per
/// `feedback_no_htod_htoh_only_mapped_pinned`. Reset to all-zeros at
/// fold boundary (Pearl A sentinel via state-reset registry, Task A12;
/// buffer grew 141→207 in Task A13.0 to cover the 29 retrofit producers).
/// Indexed by ISV slot's offset from `SP4_SLOT_BASE` × 3.
/// buffer grew 141→207 in Task A13.0 to cover the 29 retrofit producers,
/// then 207→213 in Layer C #260 to cover bw_d_h_s2_p99 +
/// q_dir_grad_p99). Indexed by ISV slot's offset from `SP4_SLOT_BASE`
/// × 3 for SP4-block slots [131..173); retrofit producers (slots
/// 40..69 in scratch) use scratch_idx × 3 directly.
pub(crate) wiener_state_buf: MappedF32Buffer,
/// SP4 Pearl C engagement counter — per (param-group, Adam-launch-block),
@@ -2989,17 +3034,18 @@ pub struct GpuDqnTrainer {
/// at fold boundary via the state-reset registry (Task A12).
pub(crate) pearl_c_rate_deficit_ema: [f32; SP4_PARAM_GROUP_COUNT],
/// SP4 producer step-observation scratch — each of the 69 producers
/// (40 SP4 + 29 retrofit) writes its per-step `step_observation` to
/// its dedicated slot here. The chained GPU `apply_pearls_ad_kernel`
/// (2026-05-01 GPU-Pearls refactor) reads these step observations
/// directly via `dev_ptr`, computes the new `x_mean` per slot, and
/// writes to the corresponding ISV bound slot — same-stream, no host
/// sync. Mapped-pinned f32 retained for cross-boundary debug
/// inspection; the host-side `read_volatile` Phase 2 pattern was
/// eliminated per `feedback_no_cpu_compute_strict.md`.
/// SP4 producer step-observation scratch — each of the 71 producers
/// (40 SP4 + 29 retrofit + 2 Layer-C-#260) writes its per-step
/// `step_observation` to its dedicated slot here. The chained GPU
/// `apply_pearls_ad_kernel` (2026-05-01 GPU-Pearls refactor) reads
/// these step observations directly via `dev_ptr`, computes the new
/// `x_mean` per slot, and writes to the corresponding ISV bound
/// slot — same-stream, no host sync. Mapped-pinned f32 retained for
/// cross-boundary debug inspection; the host-side `read_volatile`
/// Phase 2 pattern was eliminated per
/// `feedback_no_cpu_compute_strict.md`.
///
/// Layout (post Task A13):
/// Layout (post Layer C #260):
/// [0..40) — SP4 producers (Tasks A5-A11)
/// [40] — h_s2_rms_ema → ISV[96] (A13.0)
/// [41..43) — aux_heads_loss_ema → ISV[113],ISV[114] (A13.1)
@@ -3009,6 +3055,8 @@ pub struct GpuDqnTrainer {
/// [53..59) — vsn_mask_ema → ISV[105..111) (A13.3)
/// [59..63) — iqn_quantile_ema → ISV[99..103]\med (A13.4)
/// [63..69) — reward_component_ema → ISV[63..69) (A13.5)
/// [69] — bw_d_h_s2_p99 → ISV[171] (Layer C #260)
/// [70] — q_dir_grad_p99 → ISV[172] (Layer C #260)
pub(crate) producer_step_scratch_buf: MappedF32Buffer,
/// SP4 Task A7 fix-up #2: oracle sub-buffer device-pointer table.
@@ -3707,6 +3755,33 @@ pub struct GpuDqnTrainer {
/// from `SP4_H_S2_P99_CUBIN`.
h_s2_p99_update: CudaFunction,
/// SP4 Layer C #260 (2026-05-01): bw_d_h_s2_p99 producer kernel for
/// ISV[BW_D_H_S2_BOUND_INDEX=171]. Single-block 256-thread histogram
/// p99 kernel reading `bw_d_h_s2 [B × SH2]` (IQN trunk-backward d_h_s2
/// buffer); writes step_p99 to `producer_step_scratch_buf[69]`. The
/// launcher (`launch_sp4_bw_d_h_s2_p99`) chains the GPU
/// `apply_pearls_ad_kernel` on the same stream — graph-capture-
/// compatible by construction. Migrated from SP1-Phase-C
/// `1e6 × ISV[H_S2_RMS_EMA].max(1.0)` per
/// `feedback_isv_for_adaptive_bounds`. Loaded from
/// `SP4_BW_D_H_S2_P99_CUBIN`.
bw_d_h_s2_p99_update: CudaFunction,
/// SP4 Layer C #260 (2026-05-01): q_dir_grad_p99 producer kernel for
/// ISV[Q_DIR_GRAD_BOUND_INDEX=172]. Single-block 256-thread multi-sub-
/// buffer histogram p99 kernel reading the union of `d_value_logits`
/// ∪ `d_adv_logits` (n_sub=2); writes step_p99 to
/// `producer_step_scratch_buf[70]`. The launcher
/// (`launch_sp4_q_dir_grad_p99`) populates the 2-element pointer +
/// count tables in `oracle_subbuf_table_buf` / `oracle_subbuf_counts_buf`
/// (reusing the existing param_group_oracle infra), then chains the GPU
/// `apply_pearls_ad_kernel` on the same stream — graph-capture-
/// compatible by construction. Migrated from SP1-Phase-C
/// `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)` per
/// `feedback_isv_for_adaptive_bounds`. Loaded from
/// `SP4_Q_DIR_GRAD_P99_CUBIN`.
q_dir_grad_p99_update: CudaFunction,
/// SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread
/// device-side loop kernel that applies Pearls A+D to N consecutive
/// ISV slots on the producer's stream. Consumed by every SP4 producer
@@ -7584,35 +7659,34 @@ impl GpuDqnTrainer {
// `docs/dqn-backward-nan-audit.md` (slot 35 row).
self.check_nan_f32(self.ptrs.bw_d_h_s2, b * sh2, 35)?;
// SP1 Phase C (slot 35 buffer / GEMM input sanitise): clamp NaN/Inf
// to 0 + bound magnitude to 1e6 × ISV[H_S2_RMS_EMA_INDEX=96]. The
// buffer here is `bw_d_h_s2` post-DtoD copy from `iqn_d_h_s2_ptr`
// (slot 27 source). Pre-call defence — bounds the cuBLAS GEMM
// accumulator's intermediate products so the reductions inside
// `encoder_backward_chain` (Linear_a/Linear_b dW/dX/dB GEMMs) cannot
// produce f32-overflow NaN from the kind of pathological inputs
// observed at smoke-test-xvzgk (commit f139a63ee). 1e3 ε floor
// covers the uninitialised-ISV case (Invariant 1 carve-out per
// `feedback_isv_for_adaptive_bounds`). F0 paper-review: F0
// ISV[96] ≈ 1.0 (LayerNorm output RMS), max_abs = 1e6 — F0-typical
// |iqn_d_h_s2| ≤ ~10² so guard is a no-op on F0; F1 overflows
// (≥ 1e6) trigger clamp. The slot 35 NaN check at line 6949 fires
// BEFORE this clamp, so any NaN entry is captured in the regression
// sentinel before sanitisation — see Issue 1 fix in commit fix-up.
// SP1 Phase C fix-up #2 (cold-start ε floor): smoke-test-dr2bn
// (commit 19b008e1c) FAILED at F1 step 240 with the ORIGINAL formula
// `(1e6 × isv).max(1e3)` because fold-boundary ISV reset puts ISV[96]
// near 0, giving max_abs = 1e3 — too tight; clipped legitimate F1
// startup gradients to ±1e3, destabilizing Adam → cuBLAS overflow.
// New formula `1e6 × isv.max(1.0)` guarantees max_abs ≥ 1e6 in all
// ISV states (cold-start = 1e6, warm = 1e6 × isv). F0 no-op intent
// preserved (F0 inputs ≪ 1e6); F1 startup gradients ≤ 1e6 are
// un-clipped. Per `feedback_isv_for_adaptive_bounds`, the 1.0 ε
// floor is a numerical-stability bound (Invariant 1 carve-out) on
// the ISV multiplier, not a hardcoded tuning constant on the bound
// itself.
let h_s2_rms_ema = self.read_isv_signal_at(H_S2_RMS_EMA_INDEX);
let max_abs_input = 1e6_f32 * h_s2_rms_ema.max(1.0_f32);
// SP4 Layer C #260 (2026-05-01): backward-grad sanitiser bound for
// `bw_d_h_s2` is now ISV-driven via the dedicated
// `BW_D_H_S2_BOUND_INDEX=171` slot. Producer
// (`launch_sp4_bw_d_h_s2_p99`) measures p99(|bw_d_h_s2|) post-DtoD-
// copy + post-NaN-check, applies Pearls A+D on the same stream, and
// writes the smoothed bound to ISV[171]. Consumer reads the slot
// directly with `EPS_CLAMP_FLOOR=1.0` ε on cold-start (Invariant 1
// carve-out per `feedback_isv_for_adaptive_bounds`).
//
// Migrated from SP1-Phase-C `1e6 × ISV[H_S2_RMS_EMA].max(1.0)`
// hardcoded multiplier — the last remaining `1e6 × isv` constant
// in the cuBLAS backward-grad sanitiser path after Layer A/B
// closed out the rest of SP4. The new bound is Pearls-smoothed
// p99 of the actual gradient distribution, tighter than the
// pre-existing 1e6 ceiling but appropriate-scale for the observed
// values.
//
// Ordering invariant: the slot 35 NaN check at line 6949 (and the
// companion check just above at line 7585) fires BEFORE this
// clamp, so any NaN entry is captured in the regression sentinel
// before the producer measures the clean distribution. Producer
// launches BEFORE the consumer at this site.
use crate::cuda_pipeline::sp4_isv_slots::BW_D_H_S2_BOUND_INDEX;
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
self.launch_sp4_bw_d_h_s2_p99()?;
let max_abs_input = self
.read_isv_signal_at(BW_D_H_S2_BOUND_INDEX)
.max(EPS_CLAMP_FLOOR);
self.launch_clamp_finite_f32(self.ptrs.bw_d_h_s2, b * sh2, max_abs_input)?;
// ── 3. GRN trunk backward (encoder_backward_chain) ──
@@ -9652,6 +9726,243 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP4 Layer C #260 (2026-05-01): producer for ISV[BW_D_H_S2_BOUND_INDEX=171].
///
/// Reads `bw_d_h_s2 [B × SH2]` (the IQN trunk-backward d_h_s2 buffer
/// post-DtoD copy from `iqn_d_h_s2_ptr` and post slot-35 NaN check),
/// launches the single-block 256-thread `bw_d_h_s2_p99_update` kernel
/// which computes p99(|bw_d_h_s2|) via `sp4_histogram_p99<256>` and
/// writes the step observation to `producer_step_scratch_buf[SCRATCH_IDX=69]`
/// with `__threadfence_system()`. Then chains the GPU
/// `apply_pearls_ad_kernel` on the same stream — applies Pearls A+D
/// to ISV[BW_D_H_S2_BOUND_INDEX=171] using
/// `wiener_state_buf[(171 - SP4_SLOT_BASE) * 3 = 120 .. +3]`.
/// Graph-capture-compatible by construction (no host sync; same-stream
/// ordering with the producer kernel).
///
/// **Migrated from SP1-Phase-C `1e6 × ISV[H_S2_RMS_EMA].max(1.0)`**
/// hardcoded multiplier per `feedback_isv_for_adaptive_bounds`.
/// Consumer at `apply_iqn_trunk_gradient` (~line 7615) reads
/// `ISV[BW_D_H_S2_BOUND_INDEX].max(EPS_CLAMP_FLOOR)` directly.
///
/// Caller contract: invoke BEFORE the consumer's
/// `launch_clamp_finite_f32(ptrs.bw_d_h_s2, ...)` so the bound is
/// fresh; the producer launches AFTER the slot-35 NaN check at
/// `apply_iqn_trunk_gradient`'s line 7585 (the NaN check captures
/// the regression sentinel before the producer measures the clean
/// distribution — same ordering invariant as the inline slot 24/25
/// checks in `launch_cublas_backward_to`).
pub fn launch_sp4_bw_d_h_s2_p99(&self) -> Result<(), MLError> {
use crate::cuda_pipeline::sp4_isv_slots::{BW_D_H_S2_BOUND_INDEX, SP4_SLOT_BASE};
use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
// Layer C #260: scratch slot 69 — first Layer C slot post-Layer-A
// retrofits (which occupy [40..69)).
const SCRATCH_IDX: usize = 69;
// Shared memory: 8 warps × 256 bins × sizeof(int) = 8192 bytes.
// Same contract as the other histogram-based SP4 producers.
const SHARED_BYTES: u32 = (256 / 32) * 256 * 4;
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: SHARED_BYTES,
};
let b = self.config.batch_size;
let sh2 = self.config.shared_h2;
let count: i32 = (b * sh2) as i32;
let buf_ptr = self.ptrs.bw_d_h_s2;
let scratch_dev = self.producer_step_scratch_buf.dev_ptr;
let scratch_idx_arg: i32 = SCRATCH_IDX as i32;
// Safety: kernel signature `(const float*, int, float*, int)`
// matches the four args below; both device pointers are valid for
// the lifetime of `self`.
unsafe {
self.stream
.launch_builder(&self.bw_d_h_s2_p99_update)
.arg(&buf_ptr)
.arg(&count)
.arg(&scratch_dev)
.arg(&scratch_idx_arg)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("bw_d_h_s2_p99_update launch: {e}")))?;
}
// GPU Pearls A+D: stream-ordered with the producer kernel, no host
// sync (graph-capture-compatible). Wiener offset for SP4 slot 171
// → (171-131)*3 = 120. Kernel-side `step_obs == 0` short-circuit
// handles degenerate-zero cases (e.g. fresh memset of bw_d_h_s2
// before the IQN DtoD copy fires the very first step).
let isv_idx = BW_D_H_S2_BOUND_INDEX;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev,
SCRATCH_IDX as i32,
self.isv_signals_dev_ptr,
isv_idx as i32,
self.wiener_state_buf.dev_ptr,
((isv_idx - SP4_SLOT_BASE) * 3) as i32,
1,
)?;
}
Ok(())
}
/// SP4 Layer C #260 (2026-05-01): producer for ISV[Q_DIR_GRAD_BOUND_INDEX=172].
///
/// Reads the union of `d_value_logits` ∪ `d_adv_logits` (both dueling-
/// head logits gradient buffers) as one logical distribution, launches
/// the single-block 256-thread `q_dir_grad_p99_update` kernel which
/// computes p99(|union|) via `sp4_histogram_p99_multi<256>` with
/// n_sub=2, and writes the step observation to
/// `producer_step_scratch_buf[SCRATCH_IDX=70]` with
/// `__threadfence_system()`. Then chains the GPU
/// `apply_pearls_ad_kernel` on the same stream — applies Pearls A+D
/// to ISV[Q_DIR_GRAD_BOUND_INDEX=172] using
/// `wiener_state_buf[(172 - SP4_SLOT_BASE) * 3 = 123 .. +3]`.
/// Graph-capture-compatible by construction (no host sync; same-stream
/// ordering with the producer kernel).
///
/// **Migrated from SP1-Phase-C `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)`**
/// hardcoded multiplier per `feedback_isv_for_adaptive_bounds`.
/// Consumer at `launch_cublas_backward_to` (~line 20679) reads
/// `ISV[Q_DIR_GRAD_BOUND_INDEX].max(EPS_CLAMP_FLOOR)` directly. The
/// same bound was previously applied to BOTH `d_value_logits` AND
/// `d_adv_logits`, so the union p99 is the natural ISV-driven
/// replacement (one bound, one ISV slot, one producer).
///
/// Sub-buffer pointer table reuses `oracle_subbuf_table_buf` (the
/// 4 × K_MAX = 16 u64 mapped-pinned table allocated for
/// `param_group_oracle_kernel.cu`'s n_sub multi-buffer launch).
/// We use the first ptr-array slot (offset 0) and the first 2 entries
/// of the count table — the param_group_oracle launcher syncs the
/// stream after each per-group launch so there's no concurrent reader
/// of the table when this launcher writes (and our launch site is
/// strictly before the param_group_oracle invocation in the per-step
/// trainer flow if both are ever co-scheduled — currently we run
/// inside `launch_cublas_backward_to`, well before the once-per-step
/// SP4 producer block in `training_loop.rs`).
///
/// Caller contract: invoke BEFORE the consumer's
/// `launch_clamp_finite_f32(d_value_logits/d_adv_logits, ...)` so the
/// bound is fresh; the producer launches AFTER the inline slot 24/25
/// NaN checks (~line 20682 / 20684) — same ordering invariant as
/// site 1.
pub fn launch_sp4_q_dir_grad_p99(&self) -> Result<(), MLError> {
use crate::cuda_pipeline::sp4_isv_slots::{Q_DIR_GRAD_BOUND_INDEX, SP4_SLOT_BASE};
use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
// Layer C #260: scratch slot 70 — second Layer C slot.
const SCRATCH_IDX: usize = 70;
// Shared memory: 8 warps × 256 bins × sizeof(int) = 8192 bytes.
// Same contract as the other histogram-based SP4 producers.
const SHARED_BYTES: u32 = (256 / 32) * 256 * 4;
// Sub-buffer table layout: reuse the param_group_oracle
// `oracle_subbuf_table_buf` (4 × K_MAX = 16 u64 entries; we only
// use the first ptr-array slot, entries [0..2)) + the first 2
// entries of `oracle_subbuf_counts_buf`. K_MAX is a private const
// in the param_group_oracle launcher; for q_dir_grad we always
// have n_sub=2 so we don't need the full 4-slot capacity.
const K_MAX: usize = 4;
const N_SUB: usize = 2;
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: SHARED_BYTES,
};
// Populate the host-visible side of the mapped-pinned tables.
// First ptr-array slot (offset 0) holds the 2 sub-buffer device
// pointers; the count table holds the 2 element counts.
//
// Safety: `oracle_subbuf_table_buf.host_ptr` is valid for
// `4 * K_MAX = 16` u64s; `oracle_subbuf_counts_buf.host_ptr` is
// valid for K_MAX = 4 i32s. Volatile writes ensure the kernel-
// visible mapping observes the freshest values once we hit the
// launch (mapped-pinned memory has device-visible coherence after
// a `__threadfence_system()`-paired write/launch ordering, which
// holds since the next op on the stream is the kernel launch).
let n_val = self.d_value_logits_buf.len();
let n_adv = self.d_adv_logits_buf.len();
let p_val = self.d_value_logits_buf.raw_ptr();
let p_adv = self.d_adv_logits_buf.raw_ptr();
unsafe {
let table_host = self.oracle_subbuf_table_buf.host_ptr;
let counts_host = self.oracle_subbuf_counts_buf.host_ptr;
// Ptr-array slot 0 (offset 0 within the 4 × K_MAX layout):
// entries [0..N_SUB) = sub-buffer device pointers.
std::ptr::write_volatile(table_host.add(0 * K_MAX + 0), p_val);
std::ptr::write_volatile(table_host.add(0 * K_MAX + 1), p_adv);
// Zero unused tail entries so a bug that reads past N_SUB
// lands in known-safe territory.
for s in N_SUB..K_MAX {
std::ptr::write_volatile(table_host.add(0 * K_MAX + s), 0u64);
}
// Counts table: entries [0..N_SUB) = sub-buffer element counts.
std::ptr::write_volatile(counts_host.add(0), n_val as i32);
std::ptr::write_volatile(counts_host.add(1), n_adv as i32);
for s in N_SUB..K_MAX {
std::ptr::write_volatile(counts_host.add(s), 0i32);
}
}
// Device-pointer offsets into `oracle_subbuf_table_buf` /
// `oracle_subbuf_counts_buf`. We use the FIRST ptr-array slot
// (offset 0); the counts table is a flat i32 array.
let sub_buf_ptrs_dev = self.oracle_subbuf_table_buf.dev_ptr;
let sub_counts_dev = self.oracle_subbuf_counts_buf.dev_ptr;
let n_sub_i32: i32 = N_SUB as i32;
let total_count_i32: i32 = (n_val + n_adv) as i32;
let scratch_dev = self.producer_step_scratch_buf.dev_ptr;
let scratch_idx_arg: i32 = SCRATCH_IDX as i32;
// Safety: kernel signature
// `(const u64*, const i32*, i32, i32, float*, i32)`
// matches the six args below; both device pointers (table buf,
// counts buf) come from constructor-allocated mapped-pinned
// buffers and are valid for the lifetime of `self`.
unsafe {
self.stream
.launch_builder(&self.q_dir_grad_p99_update)
.arg(&sub_buf_ptrs_dev)
.arg(&sub_counts_dev)
.arg(&n_sub_i32)
.arg(&total_count_i32)
.arg(&scratch_dev)
.arg(&scratch_idx_arg)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("q_dir_grad_p99_update launch: {e}")))?;
}
// GPU Pearls A+D: stream-ordered with the producer kernel, no host
// sync (graph-capture-compatible). Wiener offset for SP4 slot 172
// → (172-131)*3 = 123. Kernel-side `step_obs == 0` short-circuit
// handles degenerate-zero cases (uninitialised buffers).
let isv_idx = Q_DIR_GRAD_BOUND_INDEX;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev,
SCRATCH_IDX as i32,
self.isv_signals_dev_ptr,
isv_idx as i32,
self.wiener_state_buf.dev_ptr,
((isv_idx - SP4_SLOT_BASE) * 3) as i32,
1,
)?;
}
Ok(())
}
/// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update`
/// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129].
///
@@ -11521,6 +11832,35 @@ impl GpuDqnTrainer {
.map_err(|e| MLError::ModelError(format!("h_s2_p99_update load: {e}")))?
};
// SP4 Layer C #260: load bw_d_h_s2_p99 producer kernel (cold-path).
// Single-block 256-thread kernel reading `bw_d_h_s2 [B × SH2]` via
// `sp4_histogram_p99<256>`, writing step_p99 to
// `producer_step_scratch_buf[69]`. Pearls A+D applied via the GPU
// `apply_pearls_ad_kernel`. Replaces SP1-Phase-C
// `1e6 × ISV[H_S2_RMS_EMA].max(1.0)` hardcoded multiplier in
// `apply_iqn_trunk_gradient` (~line 7615).
let bw_d_h_s2_p99_update = {
let module = stream.context().load_cubin(SP4_BW_D_H_S2_P99_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp4 bw_d_h_s2_p99 cubin load: {e}")))?;
module.load_function("bw_d_h_s2_p99_update")
.map_err(|e| MLError::ModelError(format!("bw_d_h_s2_p99_update load: {e}")))?
};
// SP4 Layer C #260: load q_dir_grad_p99 producer kernel (cold-path).
// Single-block 256-thread kernel reading the union of
// `d_value_logits` ∪ `d_adv_logits` (n_sub=2) via
// `sp4_histogram_p99_multi<256>`, writing step_p99 to
// `producer_step_scratch_buf[70]`. Pearls A+D applied via the GPU
// `apply_pearls_ad_kernel`. Replaces SP1-Phase-C
// `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)` hardcoded multiplier in
// `launch_cublas_backward_to` (~line 20679).
let q_dir_grad_p99_update = {
let module = stream.context().load_cubin(SP4_Q_DIR_GRAD_P99_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp4 q_dir_grad_p99 cubin load: {e}")))?;
module.load_function("q_dir_grad_p99_update")
.map_err(|e| MLError::ModelError(format!("q_dir_grad_p99_update load: {e}")))?
};
// SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread
// device-side loop kernel that consumes scratch_buf step
// observations and applies Pearls A+D in-place to ISV + Wiener
@@ -14538,6 +14878,8 @@ impl GpuDqnTrainer {
param_group_oracle_update,
grad_norm_p99_update,
h_s2_p99_update,
bw_d_h_s2_p99_update,
q_dir_grad_p99_update,
apply_pearls_ad_kernel,
q_drift_rate_ema_kernel,
fold_warmup_factor_kernel,
@@ -20638,50 +20980,53 @@ impl GpuDqnTrainer {
0u64
};
// SP1 Phase C surgical fix (slots 24/25 input sanitise): pre-call
// clamp on the cuBLAS GEMM inputs that feed `backward_full`. NaN
// or extreme finite values in `d_value_logits_buf` /
// `d_adv_logits_buf` propagate through the dueling backward GEMMs
// into `d_h_s2`/`bn_d_concat_buf` and produce the slot 32 NaN
// observed at smoke-test-xvzgk (commit f139a63ee). Bound by 1e6 ×
// ISV[Q_DIR_ABS_REF_INDEX=21] (the direction-branch |Q| EMA — same
// reference scale used elsewhere for q-magnitude bounds), with a
// 1e3 ε floor for uninitialised state (Invariant 1 carve-out per
// `feedback_isv_for_adaptive_bounds`). F0 paper-review:
// F0 |d_value_logits|, |d_adv_logits| are bounded post the existing
// 5.0 L2 norm-clip (gpu_dqn_trainer.rs:16827-16868, max_d_logit_norm
// = 5.0_f32 + clip_grad_kernel). L2 norm ≤ 5.0 across N elements
// bounds per-element |x| ≤ 5.0 absolute worst-case (single-element
// edge case) and ≤ 5.0/sqrt(N) typical-case. Both are several orders
// of magnitude below the 1e6 max_abs guard threshold, so guards are
// no-ops on F0 dynamics. F1 overflows enter via post-norm-clip GEMM
// accumulator overflow; the fix targets the GEMM input/output sites
// not the clip itself. NB: the existing 5.0 L2 norm-clip is a
// per-buffer-norm operation (aggregate L2 over the entire buffer);
// this kernel is a per-element magnitude clamp — they catch
// different failure modes (norm = aggregate; clamp_finite =
// per-element NaN/Inf/extreme), so both stay.
// SP4 Layer C #260 (2026-05-01): backward-grad sanitiser bound for
// `d_value_logits ∪ d_adv_logits` is now ISV-driven via the
// dedicated `Q_DIR_GRAD_BOUND_INDEX=172` slot. Producer
// (`launch_sp4_q_dir_grad_p99`) measures p99(|union of both
// buffers|) via `sp4_histogram_p99_multi<256>` with n_sub=2,
// applies Pearls A+D on the same stream, and writes the smoothed
// bound to ISV[172]. Consumer reads the slot directly with
// `EPS_CLAMP_FLOOR=1.0` ε on cold-start (Invariant 1 carve-out per
// `feedback_isv_for_adaptive_bounds`).
//
// Migrated from SP1-Phase-C `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)`
// hardcoded multiplier. The same bound was previously applied to
// BOTH d_value_logits AND d_adv_logits, so the union p99 is the
// natural ISV-driven replacement (one bound, one ISV slot, one
// producer). The new bound is Pearls-smoothed p99 of the actual
// joint gradient distribution, tighter than the pre-existing 1e6
// ceiling but appropriate-scale for the observed values.
//
// Inline slot 24/25 NaN checks IMMEDIATELY BEFORE each clamp
// preserve the regression sentinel: `run_nan_checks_post_backward`
// (fused_training.rs:1540) runs AFTER these clamps, so without
// these pre-clamp checks the slot 24/25 flags would observe
// sanitised zeros and never fire. Mirrors the slot 35 check at
// line 6949 pattern.
// SP1 Phase C fix-up #2 (cold-start ε floor): same pathology as
// line 6967 — fold-boundary ISV reset puts Q_DIR_ABS_REF near 0
// under original `(1e6 × isv).max(1e3)`, clipping legitimate F1
// c51_grad outputs to ±1e3 and destabilizing Bottleneck Linear bwd.
// Formula `1e6 × isv.max(1.0)` guarantees max_abs ≥ 1e6 (cold-start
// = 1e6, warm = 1e6 × isv). F0 no-op intent preserved; F1 startup
// gradients ≤ 1e6 are un-clipped.
let q_dir_abs_ref = self.read_isv_signal_at(Q_DIR_ABS_REF_INDEX);
let max_abs_q = 1e6_f32 * q_dir_abs_ref.max(1.0_f32);
// sanitised zeros and never fire. Producer launches BEFORE the
// consumer's clamp_finite calls, AFTER the slot 24/25 NaN checks
// (same ordering invariant as site 1: NaN-check → producer →
// consumer-clamp).
//
// Why the existing 5.0 L2 norm-clip is still in play: the L2
// norm-clip (gpu_dqn_trainer.rs:16827-16868, max_d_logit_norm =
// 5.0_f32 + clip_grad_kernel) is a per-buffer-norm aggregate
// operation; this kernel is a per-element magnitude clamp — they
// catch different failure modes (norm = aggregate; clamp_finite =
// per-element NaN/Inf/extreme), so both stay.
use crate::cuda_pipeline::sp4_isv_slots::Q_DIR_GRAD_BOUND_INDEX;
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
let n_val_clamp = self.d_value_logits_buf.len();
let n_adv_clamp = self.d_adv_logits_buf.len();
self.check_nan_f32(d_value_logits_ptr, n_val_clamp, 24)?;
self.launch_clamp_finite_f32(d_value_logits_ptr, n_val_clamp, max_abs_q)?;
self.check_nan_f32(d_adv_logits_ptr, n_adv_clamp, 25)?;
// Producer reads BOTH sub-buffers via the union p99; launch AFTER
// both NaN checks so the regression sentinel captures any NaN entry
// before the producer measures the clean distribution.
self.launch_sp4_q_dir_grad_p99()?;
let max_abs_q = self
.read_isv_signal_at(Q_DIR_GRAD_BOUND_INDEX)
.max(EPS_CLAMP_FLOOR);
self.launch_clamp_finite_f32(d_value_logits_ptr, n_val_clamp, max_abs_q)?;
self.launch_clamp_finite_f32(d_adv_logits_ptr, n_adv_clamp, max_abs_q)?;
bw.backward_full(
@@ -22471,9 +22816,10 @@ impl GpuDqnTrainer {
/// SP4 Layer A Task A12 (2026-04-30): bulk-zero the Pearl D Wiener-EMA
/// state at fold boundary.
///
/// `wiener_state_buf` is a 207-float mapped-pinned buffer holding all
/// 69 producers' Wiener triples `[sample_var, diff_var, x_lag]`. Pearl
/// A's first-observation sentinel requires both `prev_x_mean` (the ISV
/// `wiener_state_buf` is a 213-float mapped-pinned buffer holding all
/// 71 producers' Wiener triples `[sample_var, diff_var, x_lag]` (40 SP4
/// + 29 Task-A13 retrofit + 2 Layer-C-#260 producers). Pearl A's
/// first-observation sentinel requires both `prev_x_mean` (the ISV
/// bound slot) AND `state.x_lag` (Wiener triple offset 2) to be exactly
/// 0 when a producer first fires in a new fold. Without this reset the
/// new fold's first producer launch would EMA against fold-N's stale
@@ -22490,7 +22836,7 @@ impl GpuDqnTrainer {
// Safety: `host_ptr` is a mapped-pinned `*mut f32` of length
// `self.wiener_state_buf.len` returned by `cuMemHostAlloc`.
// f32 zero == bytewise-zero per IEEE-754 +0.0 representation, so
// `write_bytes(.., 0, ..)` produces 207 valid f32 zeros.
// `write_bytes(.., 0, ..)` produces 213 valid f32 zeros.
unsafe {
std::ptr::write_bytes(
self.wiener_state_buf.host_ptr as *mut u8,

50
crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu Normal file
View File

@@ -0,0 +1,50 @@
// crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu
//
// SP4 Layer C #260 (2026-05-01) — q_dir_grad_p99 producer for
// ISV[Q_DIR_GRAD_BOUND_INDEX=172].
//
// Single-block, 256-thread kernel reading the union of two sub-buffers
// (`d_value_logits` ∪ `d_adv_logits`) as one logical distribution, computes
// p99(|union|) via `sp4_histogram_p99_multi<256>`, and writes the per-step
// observation to `producer_step_scratch_buf[scratch_idx]` with
// `__threadfence_system()` so the chained `apply_pearls_ad_kernel` reads the
// fresh value.
//
// Migrated from SP1-Phase-C hardcoded formula
// `1e6 × ISV[Q_DIR_ABS_REF_INDEX=21].max(1.0)` per
// `feedback_isv_for_adaptive_bounds`. The same `max_abs_q` bound was applied
// to BOTH `d_value_logits` AND `d_adv_logits` at the SP1 site, so the
// migrated bound is necessarily a single ISV slot whose p99 reflects the
// joint distribution of both buffers.
//
// The host launcher (`GpuDqnTrainer::launch_sp4_q_dir_grad_p99`) builds a
// 2-element device-pointer table + 2-element count table in mapped-pinned
// memory (reusing the same 4 × K_MAX oracle table buffer used by
// `param_group_oracle_kernel.cu`'s n_sub multi-buffer launch — adequate for
// n_sub=2). Then chains the GPU `apply_pearls_ad_kernel` on the same stream,
// which writes the new x_mean to ISV[Q_DIR_GRAD_BOUND_INDEX=172] + Wiener
// state at slot 70 (offset `(172 - SP4_SLOT_BASE) * 3 = 123`). Graph-
// capture-compatible (no host sync; same-stream ordering with the producer
// kernel).
//
// Mirrors `param_group_oracle_kernel.cu`'s multi-sub-buffer signature
// convention (`unsigned long long*` ptr arrays + `int*` count array), so
// the shared `sp4_histogram_p99_multi<256>` device template can be reused.
#include "sp4_histogram_p99.cuh"
extern "C" __global__ void q_dir_grad_p99_update(
const unsigned long long* __restrict__ sub_buf_ptrs, /* [n_sub] device pointers */
const int* __restrict__ sub_counts, /* [n_sub] per-sub-buffer element counts */
int n_sub, /* always 2 in this caller */
int total_count, /* sum of sub_counts (used for p99 normalisation) */
float* __restrict__ scratch_buf,
int scratch_idx
) {
if (blockIdx.x != 0) return;
float p99 = sp4_histogram_p99_multi<256>(sub_buf_ptrs, sub_counts, n_sub, total_count);
if (threadIdx.x == 0) {
scratch_buf[scratch_idx] = p99;
__threadfence_system(); // make host-mapped write visible
}
}

57
crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
View File

@@ -1,7 +1,14 @@
//! SP4 ISV slot indices for signal-driven magnitude bounds.
//!
//! Allocated at indices 131..171 (extending ISV_TOTAL_DIM 131 → 171).
//! Allocated at indices 131..173 (extending ISV_TOTAL_DIM 131 → 173).
//! Per param-group families use 8 contiguous slots.
//!
//! Layer C #260 (2026-05-01) added two new SP4 slots:
//! - 171 = `BW_D_H_S2_BOUND_INDEX` (cuBLAS backward-grad sanitiser bound for
//! `bw_d_h_s2`; replaces SP1-Phase-C `1e6 × ISV[H_S2_RMS_EMA].max(1.0)`).
//! - 172 = `Q_DIR_GRAD_BOUND_INDEX` (cuBLAS backward-grad sanitiser bound for
//! `d_value_logits ∪ d_adv_logits`; replaces SP1-Phase-C
//! `1e6 × ISV[Q_DIR_ABS_REF].max(1.0)`).
/// Number of distinct param groups (DQN trunk/value/branch/IQN/IQL-hi/IQL-lo/attn/curiosity).
pub const SP4_PARAM_GROUP_COUNT: usize = 8;
@@ -42,11 +49,44 @@ pub const H_S2_BOUND_INDEX: usize = 169;
/// (mean|g|/mean|w|) × entropy_deficit — trunk L1 lambda. Group 0 only.
pub const L1_LAMBDA_TRUNK_INDEX: usize = 170;
/// SP4 Layer C #260 (2026-05-01): backward-grad sanitiser bound for
/// `bw_d_h_s2` buffer (slot 35 NaN-check accomplice).
///
/// Migrated from SP1-Phase-C hardcoded formula
/// `1e6 × ISV[H_S2_RMS_EMA_INDEX=96].max(1.0)` per
/// `feedback_isv_for_adaptive_bounds`. Producer: `bw_d_h_s2_p99_kernel`
/// (single-buffer 256-thread histogram p99 over `bw_d_h_s2 [B × SH2]`)
/// — see `GpuDqnTrainer::launch_sp4_bw_d_h_s2_p99`. Pearls A+D applied
/// via `apply_pearls_ad_kernel` chained on the same stream. Consumer:
/// `gpu_dqn_trainer.rs::apply_iqn_trunk_gradient` (~line 7615) reads
/// `ISV[BW_D_H_S2_BOUND_INDEX].max(EPS_CLAMP_FLOOR)` directly.
pub const BW_D_H_S2_BOUND_INDEX: usize = 171;
/// SP4 Layer C #260 (2026-05-01): backward-grad sanitiser bound for
/// `d_value_logits ∪ d_adv_logits` buffers (slot 24/25 NaN-check
/// accomplices).
///
/// Migrated from SP1-Phase-C hardcoded formula
/// `1e6 × ISV[Q_DIR_ABS_REF_INDEX=21].max(1.0)` per
/// `feedback_isv_for_adaptive_bounds`. Producer: `q_dir_grad_p99_kernel`
/// (multi-sub-buffer 256-thread histogram p99 with `n_sub=2`, treating
/// the union of `d_value_logits` + `d_adv_logits` as one logical
/// distribution) — see `GpuDqnTrainer::launch_sp4_q_dir_grad_p99`.
/// Pearls A+D applied via `apply_pearls_ad_kernel` chained on the same
/// stream. Consumer: `gpu_dqn_trainer.rs::launch_cublas_backward_to`
/// (~line 20679) reads `ISV[Q_DIR_GRAD_BOUND_INDEX].max(EPS_CLAMP_FLOOR)`
/// directly.
pub const Q_DIR_GRAD_BOUND_INDEX: usize = 172;
/// First slot used by SP4 (for sanity-check assertions).
pub const SP4_SLOT_BASE: usize = 131;
/// One past last slot used by SP4. Must equal `ISV_TOTAL_DIM` post-bump.
pub const SP4_SLOT_END: usize = 171;
///
/// Layer C #260 (2026-05-01) extended this from 171 → 173 to add
/// `BW_D_H_S2_BOUND_INDEX=171` + `Q_DIR_GRAD_BOUND_INDEX=172` for the
/// SP1-Phase-C `1e6 × isv` migration.
pub const SP4_SLOT_END: usize = 173;
// ── SP4 Task A14: Pearl C engagement-counter buffer layout ───────────
@@ -187,11 +227,11 @@ mod tests {
use super::*;
#[test]
fn slot_layout_is_contiguous_and_total_40() {
// SP4 occupies [131..171), 40 slots total.
fn slot_layout_is_contiguous_and_total_42() {
// SP4 occupies [131..173), 42 slots total post-Layer-C #260.
assert_eq!(SP4_SLOT_BASE, 131);
assert_eq!(SP4_SLOT_END, 171);
assert_eq!(SP4_SLOT_END - SP4_SLOT_BASE, 40);
assert_eq!(SP4_SLOT_END, 173);
assert_eq!(SP4_SLOT_END - SP4_SLOT_BASE, 42);
// Per-group families are contiguous and disjoint.
assert_eq!(WEIGHT_BOUND_BASE - ATOM_POS_BOUND_BASE, 4);
@@ -200,10 +240,13 @@ mod tests {
assert_eq!(WD_RATE_BASE - ADAM_V_BOUND_BASE, 8);
assert_eq!(GRAD_CLIP_BOUND_INDEX - WD_RATE_BASE, 8);
// Final 3 single-output slots: GRAD_CLIP, H_S2, L1_LAMBDA.
// Final 5 single-output slots: GRAD_CLIP, H_S2, L1_LAMBDA,
// BW_D_H_S2_BOUND, Q_DIR_GRAD_BOUND.
assert_eq!(GRAD_CLIP_BOUND_INDEX, 168);
assert_eq!(H_S2_BOUND_INDEX, 169);
assert_eq!(L1_LAMBDA_TRUNK_INDEX, 170);
assert_eq!(BW_D_H_S2_BOUND_INDEX, 171);
assert_eq!(Q_DIR_GRAD_BOUND_INDEX, 172);
// Convenience accessors.
assert_eq!(weight_bound(0), 136);

12
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -521,10 +521,20 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "ISV[L1_LAMBDA_TRUNK_INDEX=170] — trunk L1 lambda (group-0-only); SP4 Pearl A sentinel 0 at fold boundary (Task A12)",
},
RegistryEntry {
name: "sp4_bw_d_h_s2_bound",
category: ResetCategory::FoldReset,
description: "ISV[BW_D_H_S2_BOUND_INDEX=171] — Layer C #260 backward-grad sanitiser bound for `bw_d_h_s2`. SP4 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_sp4_bw_d_h_s2_p99` observation replaces directly (no stale prev_x_mean). Companion to sp4_wiener_state's slot-69 triple (offset 120..123) — both halves reset together per `feedback_no_partial_refactor.md`",
},
RegistryEntry {
name: "sp4_q_dir_grad_bound",
category: ResetCategory::FoldReset,
description: "ISV[Q_DIR_GRAD_BOUND_INDEX=172] — Layer C #260 backward-grad sanitiser bound for `d_value_logits ∪ d_adv_logits` (multi-sub-buffer p99). SP4 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_sp4_q_dir_grad_p99` observation replaces directly. Companion to sp4_wiener_state's slot-70 triple (offset 123..126) — both halves reset together per `feedback_no_partial_refactor.md`",
},
RegistryEntry {
name: "sp4_wiener_state",
category: ResetCategory::FoldReset,
description: "GpuDqnTrainer.wiener_state_buf [207 floats = 69 producers × {sample_var, diff_var, x_lag}] — SP4 Pearl D Wiener-EMA state (40 SP4 + 29 Task-A13 retrofit producers). Mapped-pinned f32; bulk write_bytes(0) at fold boundary so Pearl A's sentinel branch fires on the new fold's first producer launch (`prev_x_mean == 0 AND state.x_lag == 0` → first-observation replacement). Companion to the 9 SP4 ISV bound entries above — both halves of the sentinel contract reset together per `feedback_no_partial_refactor.md` (Task A12; buffer grew 141→207 in Task A13.0 to cover the 29 retrofit producers — h_s2_rms, aux_heads_loss×2, aux_label_scale, moe_expert_util×8, moe_gate_entropy, vsn_mask×6, iqn_quantile×4, reward_component×6)",
description: "GpuDqnTrainer.wiener_state_buf [213 floats = 71 producers × {sample_var, diff_var, x_lag}] — SP4 Pearl D Wiener-EMA state (40 SP4 + 29 Task-A13 retrofit + 2 Layer-C-#260 producers). Mapped-pinned f32; bulk write_bytes(0) at fold boundary so Pearl A's sentinel branch fires on the new fold's first producer launch (`prev_x_mean == 0 AND state.x_lag == 0` → first-observation replacement). Companion to the 11 SP4 ISV bound entries above — both halves of the sentinel contract reset together per `feedback_no_partial_refactor.md` (Task A12; buffer grew 141→207 in Task A13.0 to cover the 29 retrofit producers, then 207→213 in Layer C #260 to cover bw_d_h_s2_p99 + q_dir_grad_p99)",
},
RegistryEntry {
name: "sp4_clamp_engage_counters",

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

@@ -5884,13 +5884,43 @@ impl DQNTrainer {
);
}
}
"sp4_bw_d_h_s2_bound" => {
// Layer C #260 (2026-05-01): backward-grad sanitiser bound
// for `bw_d_h_s2`. Pearl A sentinel 0 at fold boundary so the
// new fold's first `launch_sp4_bw_d_h_s2_p99` observation
// replaces directly. Companion `sp4_wiener_state` bulk reset
// zeros the slot-69 wiener triple in lockstep so both halves
// of the sentinel contract reset together.
if let Some(ref fused) = self.fused_ctx {
fused.trainer().write_isv_signal_at(
crate::cuda_pipeline::sp4_isv_slots::BW_D_H_S2_BOUND_INDEX,
0.0_f32,
);
}
}
"sp4_q_dir_grad_bound" => {
// Layer C #260 (2026-05-01): backward-grad sanitiser bound
// for `d_value_logits ∪ d_adv_logits` (multi-sub-buffer p99).
// Pearl A sentinel 0 at fold boundary so the new fold's first
// `launch_sp4_q_dir_grad_p99` observation replaces directly.
// Companion `sp4_wiener_state` bulk reset zeros the slot-70
// wiener triple in lockstep so both halves of the sentinel
// contract reset together.
if let Some(ref fused) = self.fused_ctx {
fused.trainer().write_isv_signal_at(
crate::cuda_pipeline::sp4_isv_slots::Q_DIR_GRAD_BOUND_INDEX,
0.0_f32,
);
}
}
"sp4_wiener_state" => {
// Bulk write_bytes(0) on `wiener_state_buf` (207 mapped-
// pinned f32 = 69 producers × 3-tuple). Pearl A's first-
// Bulk write_bytes(0) on `wiener_state_buf` (213 mapped-
// pinned f32 = 71 producers × 3-tuple). Pearl A's first-
// observation sentinel requires `state.x_lag == 0` on the
// new fold's first producer launch, in lockstep with the
// SP4 ISV bound slots above (and the 29 Task-A13 retrofit
// ISV slots, all of which now also reset to 0.0).
// ISV slots + 2 Layer-C-#260 SP1-multiplier-replacement
// slots, all of which now also reset to 0.0).
if let Some(ref mut fused) = self.fused_ctx {
fused.trainer_mut().reset_sp4_wiener_state();
}

332
crates/ml/tests/sp4_producer_unit_tests.rs
View File

@@ -2398,3 +2398,335 @@ fn sp4_apply_pearls_ad_kernel_matches_host_reference_within_fp32_ulp() {
assert_eq!(wiener_gpu[8], 50.0);
}
}
// ── SP4 Layer C #260: bw_d_h_s2_p99 producer (single-buffer histogram) ──────
/// Test-only cubin for the SP4 Layer C #260 `bw_d_h_s2_p99_update` kernel.
/// Same cubin loaded by `GpuDqnTrainer::new` in production via
/// `SP4_BW_D_H_S2_P99_CUBIN`; this kernel-direct test bypasses the full
/// trainer harness and exercises the kernel's scratch-slot write contract on
/// a controlled stationary signal.
const SP4_BW_D_H_S2_P99_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/bw_d_h_s2_p99_kernel.cubin"));
fn load_sp4_bw_d_h_s2_p99_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP4_BW_D_H_S2_P99_CUBIN.to_vec())
.expect("load bw_d_h_s2_p99_kernel cubin");
module
.load_function("bw_d_h_s2_p99_update")
.expect("load bw_d_h_s2_p99_update function")
}
/// SP4 Layer C #260: validate `bw_d_h_s2_p99_update`'s per-step write
/// contract on a stationary controlled signal. Mirrors
/// `sp4_h_s2_p99_writes_step_p99_to_scratch_via_pearl_a` (the Layer A
/// cousin producer for slot 169) — same shape, different ISV slot (171)
/// and different scratch slot (69).
///
/// Inputs: 1024 constant samples = 5.0 → analytical p99 = 5.0 (constant
/// distribution). Stationary signal makes Pearl D's variance-ratio
/// converge α* → 0 after a Pearl A bootstrap, so the EMA anchors at 5.0.
///
/// Pearl A semantics tested explicitly:
/// - `prev_x_mean = 0.0` + `WienerState::ZERO` → first observation
/// replaces sentinel directly; `state.x_lag` = 5.0.
///
/// Pearl D convergence tested explicitly:
/// - 1000 stationary observations after the bootstrap → x_mean stays
/// within 1% of 5.0.
///
/// Production launcher contract validated:
/// - Kernel writes step_p99 to `scratch[SCRATCH_IDX=69]`.
/// - All other scratch slots remain zero (cascade guard).
/// - `BW_D_H_S2_BOUND_INDEX == 171`, `SP4_SLOT_BASE == 131`,
/// `(171 - SP4_SLOT_BASE) * 3 == 120` (Wiener state offset).
#[test]
#[ignore = "requires GPU"]
fn sp4_bw_d_h_s2_p99_writes_step_obs_via_pearl_a_then_converges_pearl_d() {
use ml::cuda_pipeline::sp4_isv_slots::{BW_D_H_S2_BOUND_INDEX, SP4_SLOT_BASE};
use ml::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState};
// 1024 constant samples = 5.0. Analytical p99 of a constant
// distribution is the constant itself; the kernel's histogram should
// bin all samples into the top bin and the cumulative-from-top scan
// should yield bin upper edge ≈ 5.0 (within 1/256 quantization).
const N: usize = 1024;
let samples: Vec<f32> = vec![5.0_f32; N];
let true_p99: f32 = 5.0;
// Layer C #260 scratch slot — first slot post-Layer-A retrofits
// (which occupy [40..69)).
const SCRATCH_IDX: usize = 69;
const SHARED_BYTES: u32 = (256 / 32) * 256 * 4;
let stream = make_test_stream();
let kernel = load_sp4_bw_d_h_s2_p99_kernel(&stream);
// Safety: CUDA context active on this thread (resolved via stream).
let in_buf = unsafe { MappedF32Buffer::new(samples.len()) }
.expect("alloc MappedF32Buffer for SP4 bw_d_h_s2_p99 input");
in_buf.write_from_slice(&samples);
let scratch_buf = unsafe { MappedF32Buffer::new(SP4_PRODUCER_COUNT) }
.expect("alloc MappedF32Buffer for SP4 producer scratch");
let count_i32 = i32::try_from(samples.len()).expect("sample count fits in i32");
let scratch_idx_arg: i32 = SCRATCH_IDX as i32;
let in_dev_ptr = in_buf.dev_ptr;
let scratch_dev_ptr = scratch_buf.dev_ptr;
unsafe {
stream
.launch_builder(&kernel)
.arg(&in_dev_ptr)
.arg(&count_i32)
.arg(&scratch_dev_ptr)
.arg(&scratch_idx_arg)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: SHARED_BYTES,
})
.expect("launch bw_d_h_s2_p99_update");
}
stream
.synchronize()
.expect("synchronize after bw_d_h_s2_p99_update launch");
// Verify ALL non-target slots remained zero — cascade-write guard.
let host = scratch_buf.read_all();
for (i, &v) in host.iter().enumerate() {
if i != SCRATCH_IDX {
assert_eq!(
v, 0.0,
"bw_d_h_s2_p99 producer wrote to scratch[{i}] outside target slot {SCRATCH_IDX}: {v}",
);
}
}
let step_p99 = host[SCRATCH_IDX];
let rel_err = ((step_p99 - true_p99) / true_p99).abs();
println!(
"SP4 bw_d_h_s2_p99 — true_p99={true_p99:.5}, step_p99 (kernel)={step_p99:.5}, rel_err={rel_err:.5}",
);
assert!(
step_p99 > 0.0,
"kernel step_p99 (slot {SCRATCH_IDX}) should be positive; got {step_p99}",
);
assert!(
rel_err < 0.05,
"kernel step_p99 ({step_p99}) vs true p99 ({true_p99}) rel_err {rel_err} > 5%",
);
// ── Pearl A bootstrap (host-side, mirrors `launch_sp4_bw_d_h_s2_p99`) ──
let prev_x_mean: f32 = 0.0;
let mut state = WienerState::ZERO;
let new_x_mean = pearls_ad_update(prev_x_mean, &mut state, step_p99);
assert_eq!(
new_x_mean, step_p99,
"Pearl A: first observation must replace sentinel directly; got {new_x_mean} vs step_p99={step_p99}",
);
assert_eq!(
state.x_lag, step_p99,
"Pearl A: x_lag must seed to first observation",
);
assert_eq!(state.sample_var, 0.0, "Pearl A: sample_var stays zero");
assert_eq!(state.diff_var, 0.0, "Pearl A: diff_var stays zero");
// ── Pearl D stationary convergence: 1000 observations of 5.0 ──
let mut x_mean = new_x_mean;
for _ in 0..1000 {
x_mean = pearls_ad_update(x_mean, &mut state, step_p99);
}
assert!(
((x_mean - true_p99) / true_p99).abs() < 0.01,
"Pearl D: stationary signal converged to {x_mean} (expected {true_p99}, 1% tolerance)",
);
// Slot-layout invariants — guard against off-by-one drift in
// BW_D_H_S2_BOUND_INDEX or SP4_SLOT_BASE.
assert_eq!(BW_D_H_S2_BOUND_INDEX, 171);
assert_eq!(SP4_SLOT_BASE, 131);
assert_eq!((BW_D_H_S2_BOUND_INDEX - SP4_SLOT_BASE) * 3, 120);
}
// ── SP4 Layer C #260: q_dir_grad_p99 producer (multi-sub-buffer histogram) ──
/// Test-only cubin for the SP4 Layer C #260 `q_dir_grad_p99_update`
/// kernel. Same cubin loaded by `GpuDqnTrainer::new` in production via
/// `SP4_Q_DIR_GRAD_P99_CUBIN`; this kernel-direct test bypasses the full
/// trainer harness and exercises the multi-sub-buffer path on a known
/// analytical distribution.
const SP4_Q_DIR_GRAD_P99_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/q_dir_grad_p99_kernel.cubin"));
fn load_sp4_q_dir_grad_p99_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP4_Q_DIR_GRAD_P99_CUBIN.to_vec())
.expect("load q_dir_grad_p99_kernel cubin");
module
.load_function("q_dir_grad_p99_update")
.expect("load q_dir_grad_p99_update function")
}
/// SP4 Layer C #260: validate `q_dir_grad_p99_update`'s multi-sub-buffer
/// p99 contract on a controlled distribution. The kernel reads the union
/// of two sub-buffers (`d_value_logits` ∪ `d_adv_logits` in production)
/// as one logical distribution.
///
/// Test surface: two sub-buffers each with `n=2048` samples drawn from
/// |N(0,1)| via Box-Muller (deterministic seed). Combined union has
/// `total=4096` samples; analytical p99 of |N(0,1)| ≈ 2.576.
///
/// Pearl A semantics tested explicitly (one bootstrap iteration on the
/// host-side reference). Production launcher contract validated:
/// - Kernel writes step_p99 to `scratch[SCRATCH_IDX=70]`.
/// - All other scratch slots remain zero (cascade guard).
/// - `Q_DIR_GRAD_BOUND_INDEX == 172`, `SP4_SLOT_BASE == 131`,
/// `(172 - SP4_SLOT_BASE) * 3 == 123` (Wiener state offset).
#[test]
#[ignore = "requires GPU"]
fn sp4_q_dir_grad_p99_multi_sub_buffer_oracle() {
use ml::cuda_pipeline::sp4_isv_slots::{Q_DIR_GRAD_BOUND_INDEX, SP4_SLOT_BASE};
use ml::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState};
use rand::rngs::StdRng;
use rand::{Rng, SeedableRng};
// Two sub-buffers each with 2048 |N(0,1)| samples; union p99 ≈ 2.576.
const N_SUB0: usize = 2048;
const N_SUB1: usize = 2048;
const TOTAL: usize = N_SUB0 + N_SUB1;
let mut rng = StdRng::seed_from_u64(0xC9_2D_60_F3);
let make_box_muller = |rng: &mut StdRng, n: usize| -> Vec<f32> {
(0..n)
.map(|_| {
let u1: f32 = rng.gen_range(1e-8_f32..1.0_f32);
let u2: f32 = rng.gen_range(0.0_f32..1.0_f32);
let z = (-2.0_f32 * u1.ln()).sqrt() * (2.0_f32 * std::f32::consts::PI * u2).cos();
z.abs()
})
.collect()
};
let samples0: Vec<f32> = make_box_muller(&mut rng, N_SUB0);
let samples1: Vec<f32> = make_box_muller(&mut rng, N_SUB1);
// Compute analytical p99 over the union via sort.
let mut all: Vec<f32> = samples0.iter().chain(samples1.iter()).copied().collect();
all.sort_by(|a, b| a.partial_cmp(b).expect("|N(0,1)| samples are finite"));
let true_p99 = all[(TOTAL * 99) / 100];
// Layer C #260 scratch slot — second Layer C slot.
const SCRATCH_IDX: usize = 70;
const N_SUB: usize = 2;
const K_MAX: usize = 4;
const SHARED_BYTES: u32 = (256 / 32) * 256 * 4;
let stream = make_test_stream();
let kernel = load_sp4_q_dir_grad_p99_kernel(&stream);
// Safety: CUDA context active on this thread (resolved via stream).
let in_buf0 = unsafe { MappedF32Buffer::new(N_SUB0) }
.expect("alloc MappedF32Buffer for SP4 q_dir_grad sub-buffer 0");
in_buf0.write_from_slice(&samples0);
let in_buf1 = unsafe { MappedF32Buffer::new(N_SUB1) }
.expect("alloc MappedF32Buffer for SP4 q_dir_grad sub-buffer 1");
in_buf1.write_from_slice(&samples1);
let scratch_buf = unsafe { MappedF32Buffer::new(SP4_PRODUCER_COUNT) }
.expect("alloc MappedF32Buffer for SP4 producer scratch");
// Build the device-readable sub-buffer pointer + count tables. Mirrors
// the production launcher's reuse of `oracle_subbuf_table_buf` (which
// is `4 × K_MAX = 16` u64 entries; we use the first ptr-array slot
// [0..N_SUB)) and `oracle_subbuf_counts_buf` (`K_MAX = 4` i32 entries;
// we use the first N_SUB).
let table_buf = unsafe { MappedU64Buffer::new(4 * K_MAX) }
.expect("alloc oracle_subbuf_table_buf for q_dir_grad test");
let counts_buf = unsafe { MappedI32Buffer::new(K_MAX) }
.expect("alloc oracle_subbuf_counts_buf for q_dir_grad test");
unsafe {
let table_host = table_buf.host_ptr;
let counts_host = counts_buf.host_ptr;
std::ptr::write_volatile(table_host.add(0 * K_MAX + 0), in_buf0.dev_ptr);
std::ptr::write_volatile(table_host.add(0 * K_MAX + 1), in_buf1.dev_ptr);
for s in N_SUB..K_MAX {
std::ptr::write_volatile(table_host.add(0 * K_MAX + s), 0u64);
}
std::ptr::write_volatile(counts_host.add(0), N_SUB0 as i32);
std::ptr::write_volatile(counts_host.add(1), N_SUB1 as i32);
for s in N_SUB..K_MAX {
std::ptr::write_volatile(counts_host.add(s), 0i32);
}
}
let sub_buf_ptrs_dev = table_buf.dev_ptr;
let sub_counts_dev = counts_buf.dev_ptr;
let n_sub_i32: i32 = N_SUB as i32;
let total_count_i32: i32 = TOTAL as i32;
let scratch_dev_ptr = scratch_buf.dev_ptr;
let scratch_idx_arg: i32 = SCRATCH_IDX as i32;
unsafe {
stream
.launch_builder(&kernel)
.arg(&sub_buf_ptrs_dev)
.arg(&sub_counts_dev)
.arg(&n_sub_i32)
.arg(&total_count_i32)
.arg(&scratch_dev_ptr)
.arg(&scratch_idx_arg)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: SHARED_BYTES,
})
.expect("launch q_dir_grad_p99_update");
}
stream
.synchronize()
.expect("synchronize after q_dir_grad_p99_update launch");
// Verify ALL non-target slots remained zero — cascade-write guard.
let host = scratch_buf.read_all();
for (i, &v) in host.iter().enumerate() {
if i != SCRATCH_IDX {
assert_eq!(
v, 0.0,
"q_dir_grad_p99 producer wrote to scratch[{i}] outside target slot {SCRATCH_IDX}: {v}",
);
}
}
let step_p99 = host[SCRATCH_IDX];
let rel_err = ((step_p99 - true_p99) / true_p99).abs();
println!(
"SP4 q_dir_grad_p99 — true_p99={true_p99:.5}, step_p99 (kernel)={step_p99:.5}, rel_err={rel_err:.5}",
);
assert!(
step_p99 > 0.0,
"kernel step_p99 (slot {SCRATCH_IDX}) should be positive; got {step_p99}",
);
assert!(
rel_err < 0.05,
"kernel step_p99 ({step_p99}) vs true p99 ({true_p99}) rel_err {rel_err} > 5% (multi-sub-buffer p99)",
);
// ── Pearl A bootstrap (host-side, mirrors `launch_sp4_q_dir_grad_p99`) ──
let prev_x_mean: f32 = 0.0;
let mut state = WienerState::ZERO;
let new_x_mean = pearls_ad_update(prev_x_mean, &mut state, step_p99);
assert_eq!(
new_x_mean, step_p99,
"Pearl A: first observation must replace sentinel directly; got {new_x_mean} vs step_p99={step_p99}",
);
assert_eq!(state.x_lag, step_p99, "Pearl A: x_lag must seed to first observation");
assert_eq!(state.sample_var, 0.0, "Pearl A: sample_var stays zero");
assert_eq!(state.diff_var, 0.0, "Pearl A: diff_var stays zero");
// Slot-layout invariants — guard against off-by-one drift.
assert_eq!(Q_DIR_GRAD_BOUND_INDEX, 172);
assert_eq!(SP4_SLOT_BASE, 131);
assert_eq!((Q_DIR_GRAD_BOUND_INDEX - SP4_SLOT_BASE) * 3, 123);
}

92
docs/dqn-wire-up-audit.md
View File

@@ -2834,3 +2834,95 @@ Verification (post-fix-up):
- `git grep -nE "100 × Q_ABS_REF\\.max\\(1\\.0\\)|100\\.0f \\*
q_abs_ref|100 \\* Q_ABS_REF" crates/ml/src/cuda_pipeline/
crates/ml/src/trainers/`: ZERO matches.
### Layer C #260 — SP1-era 1e6 multiplier elimination (cuBLAS backward sanitisers, 2026-05-01)
Two SP1-Phase-C cuBLAS backward sanitisers at
`gpu_dqn_trainer.rs:7615` (bw_d_h_s2 input) and `:20679`
(d_value_logits + d_adv_logits) used hardcoded
`1e6 × signal.max(1.0)` formulas — outside SP4 mechanisms 1/2/5/6/9/10
but flagged by `feedback_isv_for_adaptive_bounds` review as the last
remaining hardcoded multipliers in the cuBLAS-backward sanitiser
consumer path after Layer A/B closed out the rest of SP4.
Migrated to the proper SP4 pattern:
- 2 new SP4 ISV slots [171, 172) extending [131..171) → [131..173)
- `BW_D_H_S2_BOUND_INDEX = 171` (single-buffer histogram p99)
- `Q_DIR_GRAD_BOUND_INDEX = 172` (multi-sub-buffer histogram p99,
n_sub=2 over `d_value_logits` ∪ `d_adv_logits`)
- 2 new producer kernels:
- `bw_d_h_s2_p99_kernel.cu` (single buffer; mirrors `h_s2_p99_kernel.cu`)
- `q_dir_grad_p99_kernel.cu` (multi-sub-buffer; mirrors the
`param_group_oracle_kernel.cu` n_sub launch convention via the
shared `sp4_histogram_p99_multi<256>` template + reuses
`oracle_subbuf_table_buf` / `oracle_subbuf_counts_buf`)
- Pearls A+D wired via existing `apply_pearls_ad_kernel` (chained on
same stream as producer; no host sync; graph-capture-compatible)
- Consumer migration:
- `apply_iqn_trunk_gradient` (line ~7615): replaces
`1e6 × ISV[H_S2_RMS_EMA].max(1.0)` with
`ISV[BW_D_H_S2_BOUND_INDEX].max(EPS_CLAMP_FLOOR)`
- `launch_cublas_backward_to` (line ~20679): replaces
`1e6 × ISV[Q_DIR_ABS_REF].max(1.0)` with
`ISV[Q_DIR_GRAD_BOUND_INDEX].max(EPS_CLAMP_FLOOR)`
- Buffer growth (single source of truth in `gpu_dqn_trainer.rs`):
- `SP4_PRODUCER_COUNT`: 69 → 71
- `SP4_WIENER_TOTAL_FLOATS`: 207 → 213 (= 71 × 3)
- `producer_step_scratch_buf`: 69 → 71 floats
- `wiener_state_buf`: 207 → 213 floats
- StateResetRegistry: 2 new `FoldReset` entries
(`sp4_bw_d_h_s2_bound`, `sp4_q_dir_grad_bound`); bulk
`sp4_wiener_state` reset now spans 213 floats; existing
`reset_named_state` dispatch arms in `training_loop.rs` extended
with 2 matching arms that write 0.0 to ISV[171] / ISV[172] at fold
boundary — both halves of Pearl A's sentinel contract reset
together per `feedback_no_partial_refactor`.
- LAYOUT_FINGERPRINT_SEED bumped (added `BW_D_H_S2_BOUND=171;` +
`Q_DIR_GRAD_BOUND=172;` + `ISV_TOTAL_DIM=173`); existing
checkpoints will fail-fast at load (intentional — re-train).
- Unit tests: 2 new GPU-gated tests in
`crates/ml/tests/sp4_producer_unit_tests.rs`:
- `sp4_bw_d_h_s2_p99_writes_step_obs_via_pearl_a_then_converges_pearl_d`
— stationary 5.0 signal, 1000-iter Pearl D convergence (1% rel-err)
- `sp4_q_dir_grad_p99_multi_sub_buffer_oracle` — 2-sub-buffer union
of |N(0,1)| samples, p99 vs analytical (5% rel-err).
Behaviour change: bound is now Pearls-smoothed p99 of the actual
gradient distribution, tighter than the pre-existing 1e6 ceiling but
appropriate-scale for the observed values. Smoke validation must
verify F0/F1/F2 still converge.
Verification (Layer C #260):
- `cargo check -p ml --offline`: clean, same 11 pre-existing warnings.
- `cargo test -p ml --lib --offline -- sp4_wiener_ema sp4_isv_slots
state_reset_registry`: 11 passed (incl. renamed
`slot_layout_is_contiguous_and_total_42`).
- `cargo test -p ml --lib --offline`: 928 passed, 14 failed (no new
failures vs. Layer B fix-up baseline).
- `cargo test -p ml --test sp4_producer_unit_tests --release --
--ignored --nocapture --test-threads=1`: 16 GPU tests passed (14
pre-existing + 2 new) on local RTX 3050 Ti.
- `git grep -nE "1e6_?f32 \\* h_s2_rms_ema|1e6_?f32 \\* q_dir_abs_ref|
1e6_?f32 \\*.*\\.max\\(1\\.0_?f32\\)" crates/ml/src/cuda_pipeline/`:
ZERO matches.
Files touched (Layer C #260):
- `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs` — 2 new slot constants,
test rename + asserts updated.
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — `SP4_PRODUCER_COUNT`
bumped, `ISV_TOTAL_DIM` bumped, `LAYOUT_FINGERPRINT_SEED` extended,
2 new cubin statics + 2 new struct fields + 2 new constructor
loaders + 2 new launchers (`launch_sp4_bw_d_h_s2_p99`,
`launch_sp4_q_dir_grad_p99`), 2 consumer sites migrated.
- `crates/ml/src/cuda_pipeline/bw_d_h_s2_p99_kernel.cu` — new kernel.
- `crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu` — new kernel.
- `crates/ml/build.rs` — 2 new kernel registrations.
- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — 2 new entries +
wiener_state count update.
- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — 2 new
`reset_named_state` dispatch arms + count update.
- `crates/ml/tests/sp4_producer_unit_tests.rs` — 2 new GPU-gated tests.
- `docs/dqn-wire-up-audit.md` — this entry.
Refs: task #260, baseline commit `1c150d190` (Layer A + B + fix-up
+ L40S smoke pass).