From 45da3eac695c2a44b0225b2c79ddd8ab8e910ee3 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 13:22:01 +0200 Subject: [PATCH] =?UTF-8?q?refactor(sp4):=20#260=20follow-ups=20=E2=80=94?= =?UTF-8?q?=20symbolic=20doc=20anchors,=20dedicated=20q=5Fdir=5Fgrad=20tab?= =?UTF-8?q?le,=20p99=20helper?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- .../src/cuda_pipeline/bw_d_h_s2_p99_kernel.cu | 9 +- .../src/cuda_pipeline/decision_transformer.rs | 2 +- .../src/cuda_pipeline/dqn_utility_kernels.cu | 6 +- .../src/cuda_pipeline/experience_kernels.cu | 8 +- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 499 +++++++++--------- .../cuda_pipeline/gpu_experience_collector.rs | 2 +- .../cuda_pipeline/q_dir_grad_p99_kernel.cu | 19 +- crates/ml/src/cuda_pipeline/sp4_isv_slots.rs | 9 +- docs/dqn-wire-up-audit.md | 2 + 9 files changed, 274 insertions(+), 282 deletions(-) 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 index e09fea902..efe61f6a3 100644 --- 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 @@ -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 diff --git a/crates/ml/src/cuda_pipeline/decision_transformer.rs b/crates/ml/src/cuda_pipeline/decision_transformer.rs index f72589b55..5e2e4d58c 100644 --- a/crates/ml/src/cuda_pipeline/decision_transformer.rs +++ b/crates/ml/src/cuda_pipeline/decision_transformer.rs @@ -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)`, diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu index 5e77032b0..c9cd3016f 100644 --- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu +++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu @@ -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). diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index a893c6dde..56fdfd169 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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 diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index a574ba1bb..a9933b6b1 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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::(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(); diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 28a2e1a6f..651079d13 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -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 + 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 index b30165cd9..3c0b9f0bb 100644 --- a/crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu +++ b/crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu @@ -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 diff --git a/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs index b52b97253..17575628b 100644 --- a/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs @@ -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). diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index e6585dbfe..f1f2001ac 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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