diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index 2d66dc74a..3fe4192c1 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -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
diff --git a/crates/ml/src/cuda_pipeline/bw_d_h_s2_p99_kernel.cu b/crates/ml/src/cuda_pipeline/bw_d_h_s2_p99_kernel.cu
new file mode 100644
index 000000000..e09fea902
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/bw_d_h_s2_p99_kernel.cu
@@ -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
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index d57ac7ba8..a574ba1bb 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -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,
diff --git a/crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu b/crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu
new file mode 100644
index 000000000..b30165cd9
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu
@@ -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
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
index bc553b218..b52b97253 100644
--- a/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
+++ b/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
@@ -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);
diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs
index e33725ce0..44f8ace61 100644
--- a/crates/ml/src/trainers/dqn/state_reset_registry.rs
+++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs
@@ -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",
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index dd2bfb007..997a3324c 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -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();
                 }
diff --git a/crates/ml/tests/sp4_producer_unit_tests.rs b/crates/ml/tests/sp4_producer_unit_tests.rs
index 1f61f9d99..5dd0c72bd 100644
--- a/crates/ml/tests/sp4_producer_unit_tests.rs
+++ b/crates/ml/tests/sp4_producer_unit_tests.rs
@@ -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);
+}
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 0f854eb89..e6585dbfe 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -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).