refactor(sp4): #260 follow-ups — symbolic doc anchors, dedicated q_dir_grad table, p99 helper

3 IMPORTANT items from #260 code-quality review (commit 88ae74ca7), addressed
post-smoke-validation (smoke-test-tkkx6 Succeeded):

1. Stale numeric line-references in 8+ doc-comments replaced with symbolic
   code anchors per feedback_trust_code_not_docs. Pre-existing stale
   `ISV_TOTAL_DIM = 60` comment also corrected.

2. q_dir_grad launcher allocated dedicated `q_dir_grad_subbuf_table_buf` +
   `q_dir_grad_subbuf_counts_buf` (2 entries each) instead of reusing the
   shared `oracle_subbuf_table_buf`. Eliminates implicit "must-run-before-
   oracle" temporal coupling; removes `K_MAX=4` local redefinition.

3. `launch_sp4_p99_producer_single_buf` helper extracted on GpuDqnTrainer.
   Collapses ~30-line boilerplate x 3 call sites (target_q, h_s2,
   bw_d_h_s2) into single-line calls. Multi-sub-buffer launchers
   (q_dir_grad, param_group_oracle) and shape-distinct producers
   (grad_norm 1-thread, atom_pos 4-iter+batched-Pearls) keep their
   bespoke shape.

Build clean, 11 SP4 lib tests pass, 16 SP4 GPU tests pass on RTX 3050 Ti.

Refs: #260 code-quality review, smoke-test-tkkx6 (commit 88ae74ca7).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 13:22:01 +02:00
parent 88ae74ca73
commit 45da3eac69
9 changed files with 274 additions and 282 deletions

View File

@@ -12,16 +12,17 @@
// fresh value.
//
// Migrated from SP1-Phase-C hardcoded formula
// `1e6 × ISV[H_S2_RMS_EMA_INDEX=96].max(1.0)` per
// `1e6 × ISV[H_S2_RMS_EMA_INDEX].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).
// new x_mean to ISV[BW_D_H_S2_BOUND_INDEX] + Wiener state at the slot's
// scratch index plus a wiener offset of
// `(BW_D_H_S2_BOUND_INDEX - SP4_SLOT_BASE) * 3` floats. 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

View File

@@ -978,7 +978,7 @@ impl DecisionTransformer {
let l1_end_dt: i32 = 0; // no L1 for DT
let l1_lambda_dt: f32 = 0.0; // disabled
// SP3 Mech 9: post-Adam |p_val| clamp is keyed on the main DQN
// trainer's ISV[Q_ABS_REF_INDEX=16]; DecisionTransformer is a
// trainer's ISV[Q_ABS_REF_INDEX]; DecisionTransformer is a
// separate offline-RL trainer with no ISV bus access and is
// outside the SP3 cuBLAS-overflow scope. Pass 0.0 to disable
// the clamp here — the kernel branch is `if (bound > 0.0f)`,

View File

@@ -99,7 +99,7 @@ extern "C" __global__ void dqn_grad_norm_finalize(
* SP3 Mech 9 (2026-04-30): post-Adam ISV-driven weight clamp. After the
* Adam + AdamW + L1-proximal step writes p_val, |p_val| is clamped to
* `weight_clamp_max_abs` (computed host-side from
* ISV[Q_ABS_REF_INDEX=16].max(1.0) × 100 — see `update_adaptive_clip` and
* ISV[Q_ABS_REF_INDEX].max(1.0) × 100 — see `update_adaptive_clip` and
* the Adam launch sites in `gpu_dqn_trainer.rs`). Targets cuBLAS sgemm
* f32-accumulator overflow at `iqn_trunk_m` / `bn_d_concat_buf` GEMMs
* with finite inputs but extreme-magnitude weights (root cause of the
@@ -498,8 +498,8 @@ extern "C" __global__ void dqn_zero_kernel(
* depth on outputs (slots 26 + 32) where extreme intermediate-product
* accumulator overflow has been observed (smoke-test-xvzgk @ commit
* f139a63ee). 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
* - 1e6 × ISV[H_S2_RMS_EMA_INDEX] for the IQN trunk path
* - 1e6 × ISV[Q_DIR_ABS_REF_INDEX] for the Bottleneck Linear path
* with a 1e3 floor applied host-side when the ISV slot is uninitialised
* (Invariant 1 ε carve-out — prevents zero-clamping a hot buffer when the
* ISV producer hasn't run yet).

View File

@@ -2541,7 +2541,7 @@ extern "C" __global__ void experience_env_step(
* (2026-04-29): |final_pnl| compounds across trades through Layer 4-9
* credits, creating a multi-trade optimism loop (bonus → Q → reward →
* bonus). Bound it to the direction-branch Q-scale EMA
* (ISV[Q_DIR_ABS_REF_INDEX=21]) so the term's unbounded multiplicand
* (ISV[Q_DIR_ABS_REF_INDEX]) so the term's unbounded multiplicand
* is the stable, gradient-decoupled Q-scale, not the trade-cumulative
* shaped reward. ISV[21] is already produced by q_stats_kernel.cu.
* NOTE: ISV slot index passed as kernel arg (no magic numbers per
@@ -2602,7 +2602,7 @@ extern "C" __global__ void experience_env_step(
* bonus EMA hit 235 by F2 ep2 in the a52d99613 baseline as well —
* pre-existing structural pathology.)
*
* Final fix: bound |reward| by ISV[Q_DIR_ABS_REF_INDEX=21], the
* Final fix: bound |reward| by ISV[Q_DIR_ABS_REF_INDEX], the
* direction-branch Q-scale EMA. The unbounded multiplicand becomes
* the stable, gradient-decoupled Q-scale rather than trade-cumulative
* shaping output. Same surgical pattern as C.4 above — pearl
@@ -4024,7 +4024,7 @@ extern "C" __global__ void hindsight_relabel_kernel(
* 2. Compute E[Q_dir_a] for a=0..b0_size-1 from dir_logits via softmax × support
* 3. Plan 4 Task 2c.3c.6: Q_dir is now adaptively scaled so its per-sample
* RMS over the b0_size action dimension matches the trunk-output RMS
* tracked by ISV[H_S2_RMS_EMA_INDEX=96] (producer kernel
* tracked by ISV[H_S2_RMS_EMA_INDEX] (producer kernel
* `h_s2_rms_ema_update` populates the slot each step). The earlier
* `eq / dz` normalisation was calibrated for the pre-GRN post-ReLU
* h_s2 scale and has been replaced. The scale falls back to
@@ -4111,7 +4111,7 @@ extern "C" __global__ void mag_concat_qdir(
float q_rms = sqrtf(q_sum_sq / (float)b0_size);
/* Pass 3: adaptive scale — write each eq_a so Q_dir's per-sample RMS
* matches the trunk-output RMS tracked by ISV[H_S2_RMS_EMA_INDEX=96].
* matches the trunk-output RMS tracked by ISV[H_S2_RMS_EMA_INDEX].
*
* Fallback: when q_rms ≈ 0 (uniform Q across actions ⇒ no spread to
* scale to a target), there is no well-defined RMS-match. Falling

View File

@@ -247,15 +247,15 @@ static SP4_BW_D_H_S2_P99_CUBIN: &[u8] =
/// `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.
/// device-pointer table + 2-element count table in dedicated mapped-pinned
/// buffers (`q_dir_grad_subbuf_table_buf` / `q_dir_grad_subbuf_counts_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] + Wiener state at
/// offset `(Q_DIR_GRAD_BOUND_INDEX - SP4_SLOT_BASE) * 3` floats. 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` reads
/// `ISV[Q_DIR_GRAD_BOUND_INDEX].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"));
@@ -315,7 +315,7 @@ pub(crate) static AUX_HEADS_LOSS_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!(
/// Plan C Phase 2 follow-up A.2 (2026-04-29): q-drift rate ISV producer.
/// Single-thread single-block cold-path kernel mirroring
/// `moe_lambda_eff_kernel.cu` / `kelly_cap_update_kernel.cu`. Reads
/// ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] plus host-passed
/// ISV[Q_ABS_REF_INDEX] + ISV[Q_DIR_ABS_REF_INDEX] plus host-passed
/// `q_mean_curr` / `q_mean_prev` scalars; writes
/// ISV[Q_DRIFT_RATE_INDEX=129] = `clip(|Δq|/denom, 0, 4)`. Consumer:
/// `tau_update_kernel.cu` multiplies the cosine-scheduled `tau_base` by
@@ -1033,7 +1033,7 @@ pub const MOE_LAMBDA_EFF_INDEX: usize = 128;
/// Per-epoch normalised q-drift rate ∈ [0, 4]. Producer:
/// `q_drift_rate_ema_update` GPU kernel (cold-path cadence, single-block
/// single-thread). Reads host-passed `q_mean_curr`, `q_mean_prev` plus
/// ISV[Q_ABS_REF_INDEX=16] and ISV[Q_DIR_ABS_REF_INDEX=21] to compute
/// ISV[Q_ABS_REF_INDEX] and ISV[Q_DIR_ABS_REF_INDEX] to compute
/// `drift_rate = |q_curr q_prev| / max(|q_prev|, |Q|_mag + |Q|_dir, 1e-6)`
/// clipped to `[0, 4]`. Consumer: `tau_update_kernel.cu` multiplies the
/// cosine-scheduled `tau_base` by `1 / (1 + ISV[129])` so the effective
@@ -3075,6 +3075,25 @@ pub struct GpuDqnTrainer {
/// `const int*`.
pub(crate) oracle_subbuf_counts_buf: MappedI32Buffer,
/// SP4 Layer C #260 follow-up (2026-05-01): dedicated `q_dir_grad`
/// sub-buffer device-pointer table — separate from
/// `oracle_subbuf_table_buf` to eliminate cross-launcher temporal
/// coupling. Mapped-pinned u64 buffer of exactly 2 entries
/// (`d_value_logits.raw_ptr()`, `d_adv_logits.raw_ptr()`); the
/// `q_dir_grad_p99_update` kernel always operates on n_sub=2.
///
/// Cross-boundary data passage only — the host writes ptrs/counts,
/// the kernel reads them. No formula on the host (no
/// `feedback_no_cpu_compute_strict` violation). Populated by
/// `launch_sp4_q_dir_grad_p99` via volatile writes once per launch.
pub(crate) q_dir_grad_subbuf_table_buf: MappedU64Buffer,
/// SP4 Layer C #260 follow-up (2026-05-01): companion to
/// `q_dir_grad_subbuf_table_buf`. Mapped-pinned i32 buffer of exactly
/// 2 entries holding the per-sub-buffer element counts
/// (`d_value_logits.len()`, `d_adv_logits.len()`).
pub(crate) q_dir_grad_subbuf_counts_buf: MappedI32Buffer,
/// SP2: fused NaN-check buffer pointer table. 12 device pointers (u64) for
/// slots 24-35. Host+device-visible mapped-pinned buffer; populated via
/// host-side write (no HtoD copy) per
@@ -3683,7 +3702,7 @@ pub struct GpuDqnTrainer {
// ── Plan 4 Task 2c.3c.5: H_S2 RMS EMA producer ────────────────────
/// Single-block kernel (256 threads, shmem-reduction, no atomicAdd). Reduces
/// RMS = sqrt(sum_sq / (B*SH2)) over `save_h_s2 [B, SH2]` and EMA-updates
/// ISV[H_S2_RMS_EMA_INDEX=96] with α = 0.05. Cold-path-cadence (one launch
/// ISV[H_S2_RMS_EMA_INDEX] with α = 0.05. Cold-path-cadence (one launch
/// per training step, alongside `reward_component_ema`). Producer-only —
/// 2c.3c.6 wires the consumer in `mag_concat_qdir`'s adaptive-scale path.
/// Loaded from `h_s2_rms_ema_kernel.cubin`.
@@ -3768,16 +3787,17 @@ pub struct GpuDqnTrainer {
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-
/// ISV[Q_DIR_GRAD_BOUND_INDEX]. 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
/// count tables in dedicated `q_dir_grad_subbuf_table_buf` /
/// `q_dir_grad_subbuf_counts_buf` (separate from
/// `oracle_subbuf_table_buf` to eliminate cross-launcher temporal
/// coupling), 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,
@@ -4104,13 +4124,17 @@ pub struct GpuDqnTrainer {
q_dir_bin_means_reduce_kernel: CudaFunction,
// ── ISV core buffers (pinned device-mapped, GPU read/write) ──
// isv_signals_pinned is sized for ISV_TOTAL_DIM (60) — the extra slots
// slots [23..31] carry per-branch Q-support centres/half-widths; [43..47)
// are the D.2 per-branch gamma slots; [47..50) Plan 1 C.6 static-config;
// [50..58) are the Plan 2 C.1 Q-quantile EMAs; [58..60) layout fingerprint.
// The network (w_isv_fc1) and history rotation still consume only the first
// ISV_NETWORK_DIM (23) slots; the tail is broadcast-bus scratchpad only.
isv_signals_pinned: *mut f32, // [ISV_TOTAL_DIM = 60]
// isv_signals_pinned is sized for ISV_TOTAL_DIM — see the
// `const ISV_TOTAL_DIM` definition + `ISV_TOTAL_DIM=` line in the
// layout-fingerprint docstring for the authoritative value. Slots
// [23..31] carry per-branch Q-support centres/half-widths; [43..47)
// are the D.2 per-branch gamma slots; [47..50) Plan 1 C.6 static-
// config; [50..58) are the Plan 2 C.1 Q-quantile EMAs; [58..60)
// layout fingerprint; SP4 [131..ISV_TOTAL_DIM) for Layer A/B/C
// ISV-driven bounds. The network (w_isv_fc1) and history rotation
// still consume only the first ISV_NETWORK_DIM slots; the tail is
// broadcast-bus scratchpad only.
isv_signals_pinned: *mut f32, // [ISV_TOTAL_DIM]
isv_signals_dev_ptr: u64,
isv_history_pinned: *mut f32, // [ISV_K * 12 = 48] — history rotates slots [0..11] only
isv_history_dev_ptr: u64,
@@ -5301,7 +5325,7 @@ impl GpuDqnTrainer {
/// `source_ptr` is the [B, SH2] buffer to use as first columns — either
/// save_h_s2 (before VSN) or vsn_masked_buf (after VSN).
///
/// Plan 4 Task 2c.3c.6: kernel now consumes ISV[H_S2_RMS_EMA_INDEX=96]
/// Plan 4 Task 2c.3c.6: kernel now consumes ISV[H_S2_RMS_EMA_INDEX]
/// (populated each step by `launch_h_s2_rms_ema`, the producer wired in
/// 2c.3c.5) so the Q_dir tail of the concat is adaptively rescaled to
/// match the trunk-output RMS regardless of GRN drift across training.
@@ -7676,11 +7700,10 @@ impl GpuDqnTrainer {
// 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.
// Ordering invariant: the slot 35 NaN check (`check_nan_f32(...,
// 35)`) fires BEFORE this clamp, so any NaN entry is captured in
// the regression sentinel before sanitisation. 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()?;
@@ -7759,7 +7782,8 @@ impl GpuDqnTrainer {
// the regression sentinel: `run_nan_checks_post_backward`
// (fused_training.rs:1540) runs AFTER this clamp, so without this
// pre-clamp check the slot 26 flag would observe sanitised zeros
// and never fire. Mirrors the slot 35 check at line 6949 pattern.
// and never fire. Mirrors the same slot-35 NaN-check + clamp
// pattern in `apply_iqn_trunk_gradient`.
self.check_nan_f32(self.ptrs.iqn_trunk_m, trunk_grad_total, 26)?;
self.launch_clamp_finite_f32(
self.ptrs.iqn_trunk_m, trunk_grad_total, max_abs_input,
@@ -8943,7 +8967,7 @@ impl GpuDqnTrainer {
/// `save_h_s2 [B, SH2]` buffer and writes the step observation to
/// `producer_step_scratch_buf[SCRATCH_IDX=40]`. Then launches the GPU
/// `apply_pearls_ad_kernel` on the same stream — applies Pearls A+D
/// to ISV[H_S2_RMS_EMA_INDEX=96] using
/// to ISV[H_S2_RMS_EMA_INDEX] using
/// `wiener_state_buf[40*3..40*3+3]` (slot 40 in the SP4 producer
/// scratch / Wiener-state layout, first retrofit slot after the 40 SP4
/// producers).
@@ -9011,6 +9035,90 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP4 helper: dispatch a single-buffer p99 producer + chained Pearls A+D.
///
/// Used by the SP4 single-buffer histogram-p99 producers (target_q,
/// h_s2, bw_d_h_s2). The kernel is expected to take exactly 4 args in
/// this order: `(const float* buf, int count, float* scratch, int
/// scratch_idx)`. Launch config is 1 block × 256 threads with the
/// histogram contract's `(256/32) * 256 * 4 = 8192` bytes of shared
/// memory. After the kernel writes its step p99 to
/// `producer_step_scratch_buf[scratch_idx]`, this method chains the
/// GPU `apply_pearls_ad_kernel` on the same stream — applies Pearls
/// A+D to ISV[isv_idx] using `wiener_state_buf[(isv_idx -
/// SP4_SLOT_BASE) * 3 .. +3]`. All Pearls A+D state is on-device; no
/// host sync (graph-capture-compatible by construction).
///
/// Multi-sub-buffer launchers (`launch_sp4_q_dir_grad_p99`,
/// `launch_sp4_param_group_oracles_all_groups`) and shape-distinct
/// producers (`launch_sp4_grad_norm_p99` with 1-thread blocks,
/// `launch_sp4_atom_pos_p99_all_branches` with 4-iteration loop +
/// batched n_slots=4 Pearls) keep their bespoke launch logic.
pub(crate) fn launch_sp4_p99_producer_single_buf(
&self,
kernel: &CudaFunction,
buf_ptr: u64,
count: i32,
scratch_idx: usize,
isv_idx: usize,
name_for_err: &'static str,
) -> Result<(), MLError> {
use crate::cuda_pipeline::sp4_isv_slots::SP4_SLOT_BASE;
use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
// Shared memory: 8 warps × 256 bins × sizeof(int) = 8192 bytes.
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 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(kernel)
.arg(&buf_ptr)
.arg(&count)
.arg(&scratch_dev)
.arg(&scratch_idx_arg)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("{name_for_err}: {e}")))?;
}
// GPU Pearls A+D: stream-ordered with the producer kernel, no host
// sync (graph-capture-compatible). Wiener offset for SP4 producers:
// `(isv_idx - SP4_SLOT_BASE) * 3`. Kernel-side `step_obs == 0`
// short-circuit handles cold-start (before the producer's source
// buffer has been populated).
let isv_idx_arg: i32 = isv_idx as i32;
let wiener_offset: i32 = ((isv_idx - SP4_SLOT_BASE) * SP4_WIENER_FLOATS_PER_SLOT) as i32;
debug_assert!(
(wiener_offset as usize + SP4_WIENER_FLOATS_PER_SLOT) <= SP4_WIENER_TOTAL_FLOATS,
"wiener offset {wiener_offset}+{SP4_WIENER_FLOATS_PER_SLOT} exceeds \
SP4_WIENER_TOTAL_FLOATS={SP4_WIENER_TOTAL_FLOATS}"
);
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev,
scratch_idx_arg,
self.isv_signals_dev_ptr,
isv_idx_arg,
self.wiener_state_buf.dev_ptr,
wiener_offset,
1,
)?;
}
Ok(())
}
/// SP4 Layer A Task A5: producer for ISV[TARGET_Q_BOUND_INDEX=131].
///
/// Reads `denoise_target_q_buf` (target-network Q-values [B, 12]), launches
@@ -9032,8 +9140,7 @@ impl GpuDqnTrainer {
/// `10 × Q_ABS_REF.max(1.0)`. Behaviour unchanged; the ISV slot now
/// holds an EMA but is not yet read.
pub fn launch_sp4_target_q_p99(&self) -> Result<(), MLError> {
use crate::cuda_pipeline::sp4_isv_slots::{TARGET_Q_BOUND_INDEX, SP4_SLOT_BASE};
use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
use crate::cuda_pipeline::sp4_isv_slots::TARGET_Q_BOUND_INDEX;
// Producer-step-scratch slot assignment. Stable layout — Tasks A6-A11
// extend with subsequent indices in producer-launch order:
@@ -9053,56 +9160,17 @@ impl 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_BOUND → ISV[171] (Layer C #260)
// 70 = Q_DIR_GRAD_BOUND → ISV[172] (Layer C #260)
const SCRATCH_IDX: usize = 0;
// Shared memory: 8 warps × 256 bins × sizeof(int) = 8192 bytes.
// Matches the contract documented in `sp4_histogram_p99.cuh` and
// mirrored by `sp4_histogram_p99_test_kernel.cu`'s launch in
// `tests/sp4_producer_unit_tests.rs`.
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 count = self.denoise_target_q_buf.len() as i32;
let buf_ptr = self.denoise_target_q_buf.raw_ptr();
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.target_q_p99_update)
.arg(&buf_ptr)
.arg(&count)
.arg(&scratch_dev)
.arg(&scratch_idx_arg)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("target_q_p99_update launch: {e}")))?;
}
// GPU Pearls A+D: same-stream ordering with the producer kernel.
// Wiener offset for SP4 producers: (isv_idx - SP4_SLOT_BASE) * 3.
let isv_idx = TARGET_Q_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(())
self.launch_sp4_p99_producer_single_buf(
&self.target_q_p99_update,
self.denoise_target_q_buf.raw_ptr(),
self.denoise_target_q_buf.len() as i32,
SCRATCH_IDX,
TARGET_Q_BOUND_INDEX,
"target_q_p99_update launch",
)
}
/// SP4 Layer A Task A6: producer for ISV[ATOM_POS_BOUND[0..4]].
@@ -9668,62 +9736,19 @@ impl GpuDqnTrainer {
/// in Layer B once all producers (A5-A9) land. Behaviour unchanged from
/// before this commit.
pub fn launch_sp4_h_s2_p99(&self) -> Result<(), MLError> {
use crate::cuda_pipeline::sp4_isv_slots::{H_S2_BOUND_INDEX, SP4_SLOT_BASE};
use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
use crate::cuda_pipeline::sp4_isv_slots::H_S2_BOUND_INDEX;
// Producer-step-scratch slot 39 is reserved for H_S2_BOUND per the
// stable layout documented on `launch_sp4_target_q_p99`.
const SCRATCH_IDX: usize = 39;
// 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 count = self.save_h_s2.len() as i32;
let buf_ptr = self.save_h_s2.raw_ptr();
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.h_s2_p99_update)
.arg(&buf_ptr)
.arg(&count)
.arg(&scratch_dev)
.arg(&scratch_idx_arg)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("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 169
// → (169-131)*3 = 114. Kernel-side `step_obs == 0` short-circuit
// handles cold-start (before first online forward populates
// `save_h_s2`).
let isv_idx = 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(())
self.launch_sp4_p99_producer_single_buf(
&self.h_s2_p99_update,
self.save_h_s2.raw_ptr(),
self.save_h_s2.len() as i32,
SCRATCH_IDX,
H_S2_BOUND_INDEX,
"h_s2_p99_update launch",
)
}
/// SP4 Layer C #260 (2026-05-01): producer for ISV[BW_D_H_S2_BOUND_INDEX=171].
@@ -9742,75 +9767,32 @@ impl GpuDqnTrainer {
///
/// **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
/// Consumer at `apply_iqn_trunk_gradient` 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`).
/// fresh; the producer launches AFTER the slot-35 NaN check
/// (`check_nan_f32(..., 35)`) in `apply_iqn_trunk_gradient` (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;
use crate::cuda_pipeline::sp4_isv_slots::BW_D_H_S2_BOUND_INDEX;
// 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(())
self.launch_sp4_p99_producer_single_buf(
&self.bw_d_h_s2_p99_update,
self.ptrs.bw_d_h_s2,
(b * sh2) as i32,
SCRATCH_IDX,
BW_D_H_S2_BOUND_INDEX,
"bw_d_h_s2_p99_update launch",
)
}
/// SP4 Layer C #260 (2026-05-01): producer for ISV[Q_DIR_GRAD_BOUND_INDEX=172].
@@ -9830,29 +9812,24 @@ impl GpuDqnTrainer {
///
/// **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
/// Consumer at `launch_cublas_backward_to` 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`).
/// Sub-buffer pointer table is the dedicated 2-entry mapped-pinned
/// `q_dir_grad_subbuf_table_buf` + `q_dir_grad_subbuf_counts_buf`
/// (allocated in the constructor). Separate from
/// `oracle_subbuf_table_buf` to eliminate the previous implicit
/// "must run before param_group_oracle" temporal coupling: each
/// launcher now owns its own descriptor tables.
///
/// 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.
/// NaN checks (`check_nan_f32(..., 24)` + `check_nan_f32(..., 25)`)
/// — 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;
@@ -9864,13 +9841,10 @@ impl GpuDqnTrainer {
// 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;
// q_dir_grad always operates on exactly 2 sub-buffers
// (`d_value_logits` `d_adv_logits`); the descriptor tables
// (`q_dir_grad_subbuf_table_buf` + `q_dir_grad_subbuf_counts_buf`)
// are sized for exactly N_SUB entries — no spare tail to zero.
const N_SUB: usize = 2;
let cfg = LaunchConfig {
@@ -9879,46 +9853,37 @@ impl GpuDqnTrainer {
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.
// Populate the host-visible side of the dedicated mapped-pinned
// tables. Layout is flat (no 4 × K_MAX ptr-array bookkeeping):
// entries [0..N_SUB) hold the sub-buffer device pointers; the
// count table holds the element counts at the same offsets.
//
// 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).
// Safety: `q_dir_grad_subbuf_table_buf.host_ptr` is valid for
// exactly N_SUB u64s; `q_dir_grad_subbuf_counts_buf.host_ptr` is
// valid for exactly N_SUB 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.
let table_host = self.q_dir_grad_subbuf_table_buf.host_ptr;
let counts_host = self.q_dir_grad_subbuf_counts_buf.host_ptr;
std::ptr::write_volatile(table_host.add(0), p_val);
std::ptr::write_volatile(table_host.add(1), p_adv);
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;
// Device pointers into the dedicated tables — flat layout, no
// K_MAX stride. No cross-launcher coupling: this launcher is the
// sole owner of these descriptor buffers.
let sub_buf_ptrs_dev = self.q_dir_grad_subbuf_table_buf.dev_ptr;
let sub_counts_dev = self.q_dir_grad_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;
@@ -9966,7 +9931,7 @@ impl GpuDqnTrainer {
/// 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].
///
/// Reads ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] plus the
/// Reads ISV[Q_ABS_REF_INDEX] + ISV[Q_DIR_ABS_REF_INDEX] plus the
/// host-passed `q_mean_curr` / `q_mean_prev` scalars; writes
/// `clip(|q_curr q_prev| / max(|q_prev|, ISV[16]+ISV[21], 1e-6), 0, 4)`
/// to ISV[129]. Consumer (`tau_update_kernel.cu`) multiplies the
@@ -11749,7 +11714,7 @@ impl GpuDqnTrainer {
// Plan 4 Task 2c.3c.5: load h_s2_rms_ema kernel (cold-path, per-step).
// Single-block 256-thread shmem-reduction kernel; no atomicAdd. Producer
// for ISV[H_S2_RMS_EMA_INDEX=96]; consumer wired in 2c.3c.6.
// for ISV[H_S2_RMS_EMA_INDEX]; consumer wired in 2c.3c.6.
let h_s2_rms_ema_kernel = {
let module = stream.context().load_cubin(H_S2_RMS_EMA_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("h_s2_rms_ema cubin load: {e}")))?;
@@ -11838,7 +11803,7 @@ impl GpuDqnTrainer {
// `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).
// `apply_iqn_trunk_gradient`.
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}")))?;
@@ -11853,7 +11818,7 @@ impl GpuDqnTrainer {
// `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).
// `launch_cublas_backward_to`.
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}")))?;
@@ -13368,6 +13333,24 @@ impl GpuDqnTrainer {
"SP4 oracle_subbuf_counts_buf alloc: {e}"
)))?;
// SP4 Layer C #260 follow-up (2026-05-01): dedicated 2-entry
// sub-buffer descriptor tables for `launch_sp4_q_dir_grad_p99`.
// Allocated separately from `oracle_subbuf_table_buf` to eliminate
// the implicit "must run before param_group_oracle" temporal
// coupling — each launcher owns its own descriptor tables.
// n_sub is fixed at 2 (`d_value_logits` `d_adv_logits`).
const Q_DIR_GRAD_SUBBUF_COUNT: usize = 2;
let q_dir_grad_subbuf_table_buf =
unsafe { MappedU64Buffer::new(Q_DIR_GRAD_SUBBUF_COUNT) }
.map_err(|e| MLError::ModelError(format!(
"SP4 q_dir_grad_subbuf_table_buf alloc: {e}"
)))?;
let q_dir_grad_subbuf_counts_buf =
unsafe { MappedI32Buffer::new(Q_DIR_GRAD_SUBBUF_COUNT) }
.map_err(|e| MLError::ModelError(format!(
"SP4 q_dir_grad_subbuf_counts_buf alloc: {e}"
)))?;
// v8: PopArt running statistics buffers
let popart_mean = stream.alloc_zeros::<f32>(1)
.map_err(|e| MLError::ModelError(format!("alloc popart_mean: {e}")))?;
@@ -14760,6 +14743,8 @@ impl GpuDqnTrainer {
producer_step_scratch_buf,
oracle_subbuf_table_buf,
oracle_subbuf_counts_buf,
q_dir_grad_subbuf_table_buf,
q_dir_grad_subbuf_counts_buf,
nan_check_buf_ptrs,
nan_check_buf_lens,
pruning_epoch: prune_ep,
@@ -17503,11 +17488,14 @@ impl GpuDqnTrainer {
/// SP1 Phase C surgical fix: launch in-place finite clamp on `buf_ptr`.
///
/// Replaces NaN/Inf with 0 and clamps finite |v| to ±max_abs (host-side
/// `f32`). 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 applied at the call site for uninitialised ISV state
/// (Invariant 1 carve-out per `feedback_isv_for_adaptive_bounds`).
/// `f32`). max_abs is sourced from per-call-site SP4 ISV-driven bounds
/// (`ISV[BW_D_H_S2_BOUND_INDEX]` for the IQN trunk path,
/// `ISV[Q_DIR_GRAD_BOUND_INDEX]` for the Bottleneck Linear path), each
/// driven by a per-step p99 producer (`launch_sp4_bw_d_h_s2_p99` /
/// `launch_sp4_q_dir_grad_p99`) chained on the producer's stream. An
/// `EPS_CLAMP_FLOOR` is applied at the call site for uninitialised
/// ISV state (Invariant 1 carve-out per
/// `feedback_isv_for_adaptive_bounds`).
///
/// Pattern matches the existing `isfinite`-guard precedent in
/// `iqn_backward_per_sample` (line 612) and `iqn_adam_kernel` (line 873).
@@ -21175,7 +21163,8 @@ impl GpuDqnTrainer {
// the regression sentinel: `run_nan_checks_post_backward`
// (fused_training.rs:1540) runs AFTER this clamp, so without this
// pre-clamp check the slot 32 flag would observe sanitised zeros
// and never fire. Mirrors the slot 35 check at line 6949 pattern.
// and never fire. Mirrors the same slot-35 NaN-check + clamp
// pattern in `apply_iqn_trunk_gradient`.
if self.config.bottleneck_dim > 0 {
let bn_d_concat_len = self.bn_d_concat_buf.len();
let bn_d_concat_ptr = self.bn_d_concat_buf.raw_ptr();

View File

@@ -3922,7 +3922,7 @@ impl GpuExperienceCollector {
// C.4 / D.4b bonus optimism break (2026-04-29):
// direction-branch Q-scale EMA slot index. The C.4 timing
// bonus and D.4b regime penalty bound their |reward|/|pnl|
// multiplicands by ISV[Q_DIR_ABS_REF_INDEX=21] to break
// multiplicands by ISV[Q_DIR_ABS_REF_INDEX] to break
// the cross-trade compounding optimism loop. ISV[21] is
// produced by q_stats_kernel.cu — no new producer needed.
// Per pearl_one_unbounded_signal_per_reward.md +

View File

@@ -11,21 +11,22 @@
// fresh value.
//
// Migrated from SP1-Phase-C hardcoded formula
// `1e6 × ISV[Q_DIR_ABS_REF_INDEX=21].max(1.0)` per
// `1e6 × ISV[Q_DIR_ABS_REF_INDEX].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).
// 2-element device-pointer table + 2-element count table in dedicated
// mapped-pinned buffers (`q_dir_grad_subbuf_table_buf` /
// `q_dir_grad_subbuf_counts_buf`, allocated in the constructor — separate
// from `oracle_subbuf_table_buf` to eliminate cross-launcher temporal
// coupling). 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] +
// Wiener state at the slot's scratch index plus a wiener offset of
// `(Q_DIR_GRAD_BOUND_INDEX - SP4_SLOT_BASE) * 3` floats. 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

View File

@@ -53,12 +53,12 @@ pub const L1_LAMBDA_TRUNK_INDEX: usize = 170;
/// `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
/// `1e6 × ISV[H_S2_RMS_EMA_INDEX].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
/// `gpu_dqn_trainer.rs::apply_iqn_trunk_gradient` reads
/// `ISV[BW_D_H_S2_BOUND_INDEX].max(EPS_CLAMP_FLOOR)` directly.
pub const BW_D_H_S2_BOUND_INDEX: usize = 171;
@@ -67,15 +67,14 @@ pub const BW_D_H_S2_BOUND_INDEX: usize = 171;
/// accomplices).
///
/// Migrated from SP1-Phase-C hardcoded formula
/// `1e6 × ISV[Q_DIR_ABS_REF_INDEX=21].max(1.0)` per
/// `1e6 × ISV[Q_DIR_ABS_REF_INDEX].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.
/// 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).

View File

@@ -2926,3 +2926,5 @@ Files touched (Layer C #260):
Refs: task #260, baseline commit `1c150d190` (Layer A + B + fix-up
+ L40S smoke pass).
follow-up: symbolic doc-anchors + dedicated q_dir_grad sub-buffer table + p99 helper extraction