feat(sp17-3.2): V_share + advantage_clip_bound producers + extended emit
Phase 3.2 lands the remaining two SP17 dueling-Q diagnostic producers
atomically with kernel + launcher + Rust wrapper + extended HEALTH_DIAG
emit + GPU oracle tests per `feedback_wire_everything_up`.
V_share[d] = |E[V]| / (|E[V]| + |E[A_centered, picked]|)
where picked = argmax_a Σ_z A_raw[i, a, z] (max-Q semantic — tractable
per-batch without depending on actions_history_buf which is collector-
time state stale relative to the cuBLAS forward at HEALTH_DIAG cadence).
Pearl-A bootstrap (sentinel 0.5) + α=WELFORD_ALPHA_MIN=0.4 + bilateral
[0, 1] clamp per `pearl_symmetric_clamp_audit`. 4 blocks × 256 threads.
advantage_clip_bound = p99(|A_centered|) × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5
via sp4_histogram_p99 (block tree-reduce + per-warp tile binning, NO
atomicAdd per `pearl_fused_per_group_statistics_oracle`). EMA α=0.01
slow per-fold + bilateral clamp [0.1, 100.0] per
`pearl_symmetric_clamp_audit`. Pearl-A bootstrap (sentinel 1.0).
Single block × 256 threads + flat |A_centered| scratch buffer
(mapped-pinned, sized to B × Σ_d b_d × NA).
Observability-only — the actual clipping wire-up is Phase 5 follow-up.
The Phase 1 mean-zero contract (commits eabcf8d52..6f53d676f) makes
A_centered a meaningful signal; this commit observes it.
Extended HEALTH_DIAG line:
HEALTH_DIAG[N]: dueling [v_share=(d=X m=Y o=Z u=W)]
[a_var=(d=A m=B o=C u=D)] [clip=K]
GPU oracle tests on RTX 3050 Ti (all pass, 13/13 SP17 tests):
- v_share_per_branch_matches_closed_form: synthetic V=2.0 + linear A
per branch; closed-form V_share = 2/(2 + |K_d × (n_d-1)/2|);
ε=1e-4. Pearl-A bootstrap REPLACES on first launch.
- advantage_clip_bound_tracks_p99_safety: synthetic A with action-
dominant + per-(i,z) jitter (the jitter is REQUIRED — pathologically
lockstep values undercount in sp4_histogram_p99's non-atomic warp
tile binning per the kernel's documented "1/(256×32) loss for
uniformly distributed signals" qualifier; concentrated values violate
the assumption. Real |A_centered| in production is continuous, so
this is a test-data-only effect.) ε=0.20 (jitter + linear histogram
quantization).
Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
Phase 3.2.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -941,6 +941,38 @@ fn main() {
|
||||
// `read_a_var_ema_per_branch()` in gpu_dqn_trainer.rs. See
|
||||
// docs/dqn-wire-up-audit.md § "SP17 Phase 3.1" for rationale.
|
||||
"sp17_a_var_ema_kernel.cu",
|
||||
// SP17 Phase 3.2 (2026-05-08): per-branch V_share EMA producer.
|
||||
// 4 blocks × 256 threads. Reads `on_v_logits_buf [B, NA]` and
|
||||
// BRANCH-MAJOR `on_b_logits_buf`; computes per-branch
|
||||
// `|E[V]| / (|E[V]| + |E[A_centered, picked]|)` where picked is
|
||||
// argmax_a Σ_z A_raw[i, a, z] (max-Q action — tractable
|
||||
// per-batch without depending on actions_history_buf). Writes
|
||||
// to ISV[V_SHARE_DIR/MAG/ORD/URG_INDEX] (slots 478..482) with
|
||||
// Pearl-A bootstrap (sentinel 0.5 = healthy 50/50 cold-start) +
|
||||
// α=WELFORD_ALPHA_MIN=0.4 EMA blend. Bilateral clamp [0, 1]
|
||||
// before write per `pearl_symmetric_clamp_audit`. Block
|
||||
// tree-reduce (no atomicAdd) per `feedback_no_atomicadd`.
|
||||
// Diagnostic observability — does NOT modify any consumer path.
|
||||
// Consumed by HEALTH_DIAG `dueling [v_share=...]` via
|
||||
// `read_v_share_per_branch()`. See
|
||||
// docs/dqn-wire-up-audit.md § "SP17 Phase 3.2" for rationale.
|
||||
"sp17_v_share_kernel.cu",
|
||||
// SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound
|
||||
// producer. Single block × 256 threads. Reads BRANCH-MAJOR
|
||||
// `on_b_logits_buf`; for each branch d, computes per-atom mean
|
||||
// of A across actions and writes |A_raw - mean[z]| into a
|
||||
// pre-allocated flat scratch buffer of size Σ_d B×b_d×NA;
|
||||
// calls `sp4_histogram_p99` over the full flat scratch (block
|
||||
// tree-reduce + per-warp tile binning, no atomicAdd per
|
||||
// `pearl_fused_per_group_statistics_oracle` and
|
||||
// `feedback_no_atomicadd`); blends `bound = p99 × 1.5` into
|
||||
// ISV[ADVANTAGE_CLIP_BOUND_INDEX=482] with Pearl-A bootstrap
|
||||
// (sentinel 1.0 = no effective clipping cold-start) + α=0.01
|
||||
// slow per-fold cadence + bilateral clamp [0.1, 100.0] per
|
||||
// `pearl_symmetric_clamp_audit`. Observability-only in this
|
||||
// commit — the actual clip wire-up is Phase 5 follow-up. See
|
||||
// docs/dqn-wire-up-audit.md § "SP17 Phase 3.2" for rationale.
|
||||
"sp17_advantage_clip_bound_kernel.cu",
|
||||
// SP14 Layer C Phase C.1 (2026-05-08): cubin entries for
|
||||
// alpha_grad_compute_kernel.cu, gradient_hack_detect_kernel.cu,
|
||||
// and sp14_scale_wire_col_kernel.cu deleted atomically with
|
||||
|
||||
@@ -686,6 +686,30 @@ pub(crate) static SP17_V_A_MEANS_DIAG_CUBIN: &[u8] =
|
||||
pub(crate) static SP17_A_VAR_EMA_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/sp17_a_var_ema_kernel.cubin"));
|
||||
|
||||
/// SP17 Phase 3.2 (2026-05-08): per-branch V_share EMA. 4 blocks × 256
|
||||
/// threads. Cold-path producer for HEALTH_DIAG `dueling [v_share=...]`.
|
||||
/// V_share = |E[V]| / (|E[V]| + |E[A_centered, picked]|) per branch,
|
||||
/// picked = argmax_a Σ_z A_raw[i, a, z] (max-Q semantic). Pearl-A
|
||||
/// bootstrap (sentinel 0.5) + α=WELFORD_ALPHA_MIN=0.4 + bilateral [0, 1]
|
||||
/// clamp. Block tree-reduce per `feedback_no_atomicadd`. Loaded from
|
||||
/// `sp17_v_share_kernel.cubin`. Consumed by
|
||||
/// `read_v_share_per_branch()`.
|
||||
pub(crate) static SP17_V_SHARE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/sp17_v_share_kernel.cubin"));
|
||||
|
||||
/// SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound producer.
|
||||
/// Single block × 256 threads. Cold-path producer for HEALTH_DIAG
|
||||
/// `dueling [clip=K]`. Computes p99(|A_centered|) × 1.5 over the full
|
||||
/// flat batch via `sp4_histogram_p99` (no atomicAdd per
|
||||
/// `pearl_fused_per_group_statistics_oracle`); EMA-blends with α=0.01
|
||||
/// + bilateral clamp [0.1, 100.0] per `pearl_symmetric_clamp_audit`.
|
||||
/// Pearl-A bootstrap on sentinel 1.0. Observability-only in Phase 3 —
|
||||
/// the actual clip wire-up is Phase 5 follow-up. Loaded from
|
||||
/// `sp17_advantage_clip_bound_kernel.cubin`. Consumed by
|
||||
/// `read_advantage_clip_bound()`.
|
||||
pub(crate) static SP17_ADVANTAGE_CLIP_BOUND_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/sp17_advantage_clip_bound_kernel.cubin"));
|
||||
|
||||
// SP14 Layer C Phase C.1 (2026-05-08): α-machinery cubin statics deleted
|
||||
// atomically with the alpha_grad_compute / gradient_hack_detect /
|
||||
// sp14_scale_wire_col kernel files. SP14_ALPHA_GRAD_CUBIN,
|
||||
@@ -6733,6 +6757,30 @@ pub struct GpuDqnTrainer {
|
||||
/// Loaded from `sp17_a_var_ema_kernel.cubin`. Consumed by
|
||||
/// `read_a_var_ema_per_branch()`.
|
||||
sp17_a_var_ema_kernel: CudaFunction,
|
||||
// ── SP17 Phase 3.2 (2026-05-08): per-branch V_share producer ──────────
|
||||
/// 4-block × 256-thread kernel computing per-branch
|
||||
/// `|E[V]| / (|E[V]| + |E[A_centered, picked]|)` (picked = argmax E[Q])
|
||||
/// and EMA-blending into ISV[V_SHARE_DIR/MAG/ORD/URG_INDEX] (slots
|
||||
/// 478..482). Pearl-A bootstrap on sentinel 0.5; α=0.4 post-bootstrap.
|
||||
/// Block tree-reduce per `feedback_no_atomicadd`. Loaded from
|
||||
/// `sp17_v_share_kernel.cubin`. Consumed by
|
||||
/// `read_v_share_per_branch()`.
|
||||
sp17_v_share_kernel: CudaFunction,
|
||||
// ── SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound ────────
|
||||
/// Single-block × 256-thread kernel using `sp4_histogram_p99` to
|
||||
/// compute `p99(|A_centered|) × 1.5` over the full flat batch and
|
||||
/// EMA-blend (α=0.01) with bilateral clamp [0.1, 100.0] into
|
||||
/// ISV[ADVANTAGE_CLIP_BOUND_INDEX=482]. Pearl-A bootstrap on sentinel
|
||||
/// 1.0. Observability-only in Phase 3 — clip wire-up is Phase 5.
|
||||
/// Loaded from `sp17_advantage_clip_bound_kernel.cubin`. Consumed by
|
||||
/// `read_advantage_clip_bound()`.
|
||||
sp17_advantage_clip_bound_kernel: CudaFunction,
|
||||
/// SP17 Phase 3.2 scratch buffer for `sp17_advantage_clip_bound`. Sized
|
||||
/// to the maximum possible flat |A_centered| tensor: B × Σ_d b_d × NA.
|
||||
/// Mapped-pinned per `feedback_no_htod_htoh_only_mapped_pinned` (no
|
||||
/// host writes — kernel-owned). Lifetime matches the trainer; reset
|
||||
/// once at construct, written-then-read inside the kernel each launch.
|
||||
sp17_advantage_clip_scratch: super::mapped_pinned::MappedF32Buffer,
|
||||
// SP14 Layer C Phase C.1 (2026-05-08): α-machinery struct fields deleted
|
||||
// atomically with the kernel files (alpha_grad_compute_kernel,
|
||||
// gradient_hack_detect_kernel, sp14_scale_wire_col_kernel) per
|
||||
@@ -9136,6 +9184,207 @@ impl GpuDqnTrainer {
|
||||
])
|
||||
}
|
||||
|
||||
/// SP17 Phase 3.2 (2026-05-08) — per-branch V_share producer launch.
|
||||
///
|
||||
/// Cold-path launcher for `sp17_v_share_kernel`. Reads
|
||||
/// `on_v_logits_buf` and `on_b_logits_buf`; computes per-branch
|
||||
/// `V_share = |E[V]| / (|E[V]| + |E[A_centered, picked]|)` where
|
||||
/// picked = argmax_a Σ_z A_raw[i, a, z] (max-Q semantic — tractable
|
||||
/// per-batch without depending on `actions_history_buf` which is
|
||||
/// collector-time state stale relative to the cuBLAS forward driving
|
||||
/// the on_*_logits buffers at HEALTH_DIAG cadence).
|
||||
///
|
||||
/// Pearl-A bootstrap on sentinel `SENTINEL_V_SHARE=0.5` (cold-start
|
||||
/// healthy 50/50 split); α=`WELFORD_ALPHA_MIN=0.4` post-bootstrap.
|
||||
/// Bilateral [0, 1] clamp before ISV write per
|
||||
/// `pearl_symmetric_clamp_audit`. 4 blocks × 256 threads.
|
||||
pub fn launch_v_share_update(&self) -> Result<(), MLError> {
|
||||
use super::sp14_isv_slots::{
|
||||
V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
|
||||
V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX,
|
||||
SENTINEL_V_SHARE, WELFORD_ALPHA_MIN,
|
||||
};
|
||||
let b = self.config.batch_size;
|
||||
let na = self.config.num_atoms;
|
||||
let b0 = self.config.branch_0_size;
|
||||
let b1 = self.config.branch_1_size;
|
||||
let b2 = self.config.branch_2_size;
|
||||
let b3 = self.config.branch_3_size;
|
||||
if b == 0 || na == 0 || (b0 + b1 + b2 + b3) == 0 {
|
||||
return Ok(());
|
||||
}
|
||||
let v_ptr = self.on_v_logits_buf.raw_ptr();
|
||||
let adv_ptr = self.on_b_logits_buf.raw_ptr();
|
||||
let isv_ptr = self.isv_signals_dev_ptr;
|
||||
let b_i32 = b as i32;
|
||||
let na_i32 = na as i32;
|
||||
let b0_i32 = b0 as i32;
|
||||
let b1_i32 = b1 as i32;
|
||||
let b2_i32 = b2 as i32;
|
||||
let b3_i32 = b3 as i32;
|
||||
let share_idx_dir = V_SHARE_DIR_INDEX as i32;
|
||||
let share_idx_mag = V_SHARE_MAG_INDEX as i32;
|
||||
let share_idx_ord = V_SHARE_ORD_INDEX as i32;
|
||||
let share_idx_urg = V_SHARE_URG_INDEX as i32;
|
||||
let alpha: f32 = WELFORD_ALPHA_MIN;
|
||||
let sentinel: f32 = SENTINEL_V_SHARE;
|
||||
// Dynamic shmem = (BLOCK_SIZE + NA + 1) × sizeof(f32) — tree-reduce
|
||||
// tile + per-atom mean cache + 1 float for E[V] broadcast.
|
||||
let block_dim: u32 = 256;
|
||||
let shmem_bytes: u32 =
|
||||
(block_dim + na as u32 + 1) * std::mem::size_of::<f32>() as u32;
|
||||
unsafe {
|
||||
self.stream
|
||||
.launch_builder(&self.sp17_v_share_kernel)
|
||||
.arg(&v_ptr)
|
||||
.arg(&adv_ptr)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&b_i32)
|
||||
.arg(&na_i32)
|
||||
.arg(&b0_i32)
|
||||
.arg(&b1_i32)
|
||||
.arg(&b2_i32)
|
||||
.arg(&b3_i32)
|
||||
.arg(&share_idx_dir)
|
||||
.arg(&share_idx_mag)
|
||||
.arg(&share_idx_ord)
|
||||
.arg(&share_idx_urg)
|
||||
.arg(&alpha)
|
||||
.arg(&sentinel)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (4, 1, 1),
|
||||
block_dim: (block_dim, 1, 1),
|
||||
shared_mem_bytes: shmem_bytes,
|
||||
})
|
||||
.map_err(|e| MLError::ModelError(format!(
|
||||
"sp17_v_share_update launch: {e}"
|
||||
)))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP17 Phase 3.2 (2026-05-08) — per-branch V_share readback.
|
||||
/// Convenience wrapper: launches → syncs → reads ISV slots 478..482.
|
||||
pub fn read_v_share_per_branch(&self) -> Result<[f32; 4], MLError> {
|
||||
use super::sp14_isv_slots::{
|
||||
V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
|
||||
V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX,
|
||||
};
|
||||
self.launch_v_share_update()?;
|
||||
self.stream
|
||||
.synchronize()
|
||||
.map_err(|e| MLError::ModelError(format!("sp17_v_share sync: {e}")))?;
|
||||
Ok([
|
||||
self.read_isv_signal_at(V_SHARE_DIR_INDEX),
|
||||
self.read_isv_signal_at(V_SHARE_MAG_INDEX),
|
||||
self.read_isv_signal_at(V_SHARE_ORD_INDEX),
|
||||
self.read_isv_signal_at(V_SHARE_URG_INDEX),
|
||||
])
|
||||
}
|
||||
|
||||
/// SP17 Phase 3.2 (2026-05-08) — adaptive advantage_clip_bound launch.
|
||||
///
|
||||
/// Cold-path launcher for `sp17_advantage_clip_bound_kernel`. Computes
|
||||
/// `p99(|A_centered|) × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5` over the
|
||||
/// full flat batch via `sp4_histogram_p99` (no atomicAdd per
|
||||
/// `pearl_fused_per_group_statistics_oracle`); EMA-blends into
|
||||
/// ISV[ADVANTAGE_CLIP_BOUND_INDEX=482] with α=`ADVANTAGE_CLIP_EMA_ALPHA=0.01`
|
||||
/// + bilateral clamp `[ADVANTAGE_CLIP_BOUND_MIN=0.1, _MAX=100.0]`
|
||||
/// per `pearl_symmetric_clamp_audit`. Pearl-A bootstrap on sentinel
|
||||
/// 1.0.
|
||||
///
|
||||
/// Observability-only in Phase 3 — the actual clipping wire-up is
|
||||
/// Phase 5 follow-up. This launcher is the producer end of the
|
||||
/// chain that feeds the HEALTH_DIAG `dueling [clip=K]` line.
|
||||
///
|
||||
/// Single block × 256 threads. Dynamic shmem = `(BLOCK / 32) × 256 × 4
|
||||
/// = 8192 bytes` for `sp4_histogram_p99`'s per-warp tile array.
|
||||
pub fn launch_advantage_clip_bound_update(&self) -> Result<(), MLError> {
|
||||
use super::sp14_isv_slots::{
|
||||
ADVANTAGE_CLIP_BOUND_INDEX,
|
||||
ADVANTAGE_CLIP_BOUND_MIN, ADVANTAGE_CLIP_BOUND_MAX,
|
||||
ADVANTAGE_CLIP_EMA_ALPHA, ADVANTAGE_CLIP_SAFETY_FACTOR,
|
||||
SENTINEL_ADVANTAGE_CLIP_BOUND,
|
||||
};
|
||||
let b = self.config.batch_size;
|
||||
let na = self.config.num_atoms;
|
||||
let b0 = self.config.branch_0_size;
|
||||
let b1 = self.config.branch_1_size;
|
||||
let b2 = self.config.branch_2_size;
|
||||
let b3 = self.config.branch_3_size;
|
||||
if b == 0 || na == 0 || (b0 + b1 + b2 + b3) == 0 {
|
||||
return Ok(());
|
||||
}
|
||||
let total = b * (b0 + b1 + b2 + b3) * na;
|
||||
let scratch_capacity = self.sp17_advantage_clip_scratch.len;
|
||||
if total > scratch_capacity {
|
||||
return Err(MLError::ModelError(format!(
|
||||
"sp17_advantage_clip_scratch too small: need {total}, have {scratch_capacity}"
|
||||
)));
|
||||
}
|
||||
let adv_ptr = self.on_b_logits_buf.raw_ptr();
|
||||
let scratch_ptr = self.sp17_advantage_clip_scratch.dev_ptr;
|
||||
let isv_ptr = self.isv_signals_dev_ptr;
|
||||
let scratch_cap_i32 = scratch_capacity as i32;
|
||||
let b_i32 = b as i32;
|
||||
let na_i32 = na as i32;
|
||||
let b0_i32 = b0 as i32;
|
||||
let b1_i32 = b1 as i32;
|
||||
let b2_i32 = b2 as i32;
|
||||
let b3_i32 = b3 as i32;
|
||||
let bound_idx = ADVANTAGE_CLIP_BOUND_INDEX as i32;
|
||||
let alpha: f32 = ADVANTAGE_CLIP_EMA_ALPHA;
|
||||
let safety: f32 = ADVANTAGE_CLIP_SAFETY_FACTOR;
|
||||
let bound_min: f32 = ADVANTAGE_CLIP_BOUND_MIN;
|
||||
let bound_max: f32 = ADVANTAGE_CLIP_BOUND_MAX;
|
||||
let sentinel: f32 = SENTINEL_ADVANTAGE_CLIP_BOUND;
|
||||
// Dynamic shmem for sp4_histogram_p99: (BLOCK / 32) × SP4_HIST_BINS
|
||||
// × sizeof(int) = 8 × 256 × 4 = 8192 bytes for BLOCK_SIZE=256.
|
||||
let block_dim: u32 = 256;
|
||||
let warps = block_dim / 32;
|
||||
let shmem_bytes: u32 = warps * 256 * std::mem::size_of::<i32>() as u32;
|
||||
unsafe {
|
||||
self.stream
|
||||
.launch_builder(&self.sp17_advantage_clip_bound_kernel)
|
||||
.arg(&adv_ptr)
|
||||
.arg(&scratch_ptr)
|
||||
.arg(&scratch_cap_i32)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&b_i32)
|
||||
.arg(&na_i32)
|
||||
.arg(&b0_i32)
|
||||
.arg(&b1_i32)
|
||||
.arg(&b2_i32)
|
||||
.arg(&b3_i32)
|
||||
.arg(&bound_idx)
|
||||
.arg(&alpha)
|
||||
.arg(&safety)
|
||||
.arg(&bound_min)
|
||||
.arg(&bound_max)
|
||||
.arg(&sentinel)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (block_dim, 1, 1),
|
||||
shared_mem_bytes: shmem_bytes,
|
||||
})
|
||||
.map_err(|e| MLError::ModelError(format!(
|
||||
"sp17_advantage_clip_bound_update launch: {e}"
|
||||
)))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP17 Phase 3.2 (2026-05-08) — adaptive advantage_clip_bound readback.
|
||||
/// Convenience wrapper: launches → syncs → reads ISV slot 482.
|
||||
pub fn read_advantage_clip_bound(&self) -> Result<f32, MLError> {
|
||||
use super::sp14_isv_slots::ADVANTAGE_CLIP_BOUND_INDEX;
|
||||
self.launch_advantage_clip_bound_update()?;
|
||||
self.stream
|
||||
.synchronize()
|
||||
.map_err(|e| MLError::ModelError(format!("sp17_advantage_clip_bound sync: {e}")))?;
|
||||
Ok(self.read_isv_signal_at(ADVANTAGE_CLIP_BOUND_INDEX))
|
||||
}
|
||||
|
||||
/// Update per-branch liquid tau modulation — GPU kernel (RK4/Euler adaptive ODE).
|
||||
///
|
||||
/// Reads per_branch_q_gaps (pinned device-mapped) and qlstm_context from device.
|
||||
@@ -20255,6 +20504,48 @@ impl GpuDqnTrainer {
|
||||
.load_function("sp17_a_var_ema_update")
|
||||
.map_err(|e| MLError::ModelError(format!("sp17_a_var_ema_update: {e}")))?;
|
||||
|
||||
// SP17 Phase 3.2 (2026-05-08): per-branch V_share producer.
|
||||
// Cold-path (epoch boundary). 4 blocks × 256 threads; Pearl-A
|
||||
// bootstrap on slots 478..482 (sentinel 0.5 cold-start). Bilateral
|
||||
// [0, 1] clamp before ISV write per `pearl_symmetric_clamp_audit`.
|
||||
let sp17_v_share_module = stream.context()
|
||||
.load_cubin(SP17_V_SHARE_CUBIN.to_vec())
|
||||
.map_err(|e| MLError::ModelError(format!("sp17 v_share cubin: {e}")))?;
|
||||
let sp17_v_share_kernel = sp17_v_share_module
|
||||
.load_function("sp17_v_share_update")
|
||||
.map_err(|e| MLError::ModelError(format!("sp17_v_share_update: {e}")))?;
|
||||
|
||||
// SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound
|
||||
// producer + flat |A_centered| scratch buffer. Single block × 256
|
||||
// threads. Pearl-A bootstrap on slot 482 (sentinel 1.0). Uses
|
||||
// `sp4_histogram_p99` (block tree-reduce + per-warp tile binning,
|
||||
// no atomicAdd per `pearl_fused_per_group_statistics_oracle`).
|
||||
// Scratch sized to max possible flat |A_centered| =
|
||||
// B × (b0 + b1 + b2 + b3) × NA floats.
|
||||
let sp17_advantage_clip_bound_module = stream.context()
|
||||
.load_cubin(SP17_ADVANTAGE_CLIP_BOUND_CUBIN.to_vec())
|
||||
.map_err(|e| MLError::ModelError(format!("sp17 advantage_clip_bound cubin: {e}")))?;
|
||||
let sp17_advantage_clip_bound_kernel = sp17_advantage_clip_bound_module
|
||||
.load_function("sp17_advantage_clip_bound_update")
|
||||
.map_err(|e| MLError::ModelError(format!("sp17_advantage_clip_bound_update: {e}")))?;
|
||||
let sp17_clip_scratch_capacity = config.batch_size
|
||||
.saturating_mul(
|
||||
config.branch_0_size + config.branch_1_size
|
||||
+ config.branch_2_size + config.branch_3_size
|
||||
)
|
||||
.saturating_mul(config.num_atoms);
|
||||
// Defensive minimum of 1 element so MappedF32Buffer::new() doesn't
|
||||
// fail on a zero-sized config (constructor invariants normally
|
||||
// prevent this; the saturating_mul above could in principle floor
|
||||
// to 0 in degenerate test configs).
|
||||
let sp17_clip_scratch_len = sp17_clip_scratch_capacity.max(1);
|
||||
let sp17_advantage_clip_scratch = unsafe {
|
||||
super::mapped_pinned::MappedF32Buffer::new(sp17_clip_scratch_len)
|
||||
}
|
||||
.map_err(|e| MLError::ModelError(format!(
|
||||
"sp17_advantage_clip_scratch alloc ({sp17_clip_scratch_len} f32): {e}"
|
||||
)))?;
|
||||
|
||||
// SP14 Layer C Phase C.1 (2026-05-08): α-machinery cubin loads
|
||||
// (alpha_grad_compute, gradient_hack_detect, sp14_scale_wire_col)
|
||||
// deleted atomically with the kernel files. Coupling A
|
||||
@@ -24494,6 +24785,9 @@ impl GpuDqnTrainer {
|
||||
sp17_v_a_means_diag_kernel,
|
||||
sp17_v_a_means_diag_buf,
|
||||
sp17_a_var_ema_kernel,
|
||||
sp17_v_share_kernel,
|
||||
sp17_advantage_clip_bound_kernel,
|
||||
sp17_advantage_clip_scratch,
|
||||
sp14_dir_concat_qaux_kernel,
|
||||
sp14_dir_qaux_concat_scratch,
|
||||
// SP14 backward dX scratch — column-SH2 wire is dropped at
|
||||
|
||||
200
crates/ml/src/cuda_pipeline/sp17_advantage_clip_bound_kernel.cu
Normal file
200
crates/ml/src/cuda_pipeline/sp17_advantage_clip_bound_kernel.cu
Normal file
@@ -0,0 +1,200 @@
|
||||
/* ══════════════════════════════════════════════════════════════════════════
|
||||
* SP17 Phase 3.2 (2026-05-08) — adaptive advantage_clip_bound producer.
|
||||
*
|
||||
* Computes an adaptive upper-bound for |A_centered| by tracking
|
||||
* advantage_clip_bound ← p99(|A_centered|) × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5
|
||||
*
|
||||
* Mirrors the SP14 P0-A REWARD_POS_CAP producer pattern (also p99 × 1.5).
|
||||
* The kernel uses `sp4_histogram_p99` (the canonical block-tree-reduce
|
||||
* percentile estimator from SP4) — NO atomicAdd per
|
||||
* `pearl_fused_per_group_statistics_oracle` and `feedback_no_atomicadd`.
|
||||
*
|
||||
* The bound is observability-only in this commit (Phase 3 is the
|
||||
* diagnostic chain); the actual clipping wire-up is Phase 5 follow-up.
|
||||
*
|
||||
* Cadence: cold-path (epoch boundary). EMA α = ADVANTAGE_CLIP_EMA_ALPHA
|
||||
* = 0.01 (slow per-fold — host-passed).
|
||||
*
|
||||
* ── Pearls + invariants ────────────────────────────────────────────────
|
||||
* - `feedback_no_atomicadd.md` — uses sp4_histogram_p99 (block
|
||||
* tree-reduce + per-warp tile binning, no atomicAdd).
|
||||
* - `pearl_first_observation_bootstrap.md` — sentinel
|
||||
* SENTINEL_ADVANTAGE_CLIP_BOUND = 1.0; cold-start REPLACE on first
|
||||
* observation.
|
||||
* - `pearl_symmetric_clamp_audit.md` — bilateral clamp to
|
||||
* [ADVANTAGE_CLIP_BOUND_MIN=0.1, ADVANTAGE_CLIP_BOUND_MAX=100.0]
|
||||
* after EMA blend.
|
||||
* - `feedback_isv_for_adaptive_bounds.md` — bound is producer-tracked,
|
||||
* NEVER hardcoded.
|
||||
* - `pearl_no_host_branches_in_captured_graph.md` — pure on-device.
|
||||
*
|
||||
* Algorithm (single-block, BLOCK=256):
|
||||
*
|
||||
* For each branch d in [0, 4):
|
||||
* Pass 1: per-atom mean of raw A over (B × n_branch) — written to
|
||||
* `s_a_mean[NA]` (block-shared, recomputed per branch).
|
||||
* Pass 2: write |A_raw[i, a, z] - a_mean[z]| into
|
||||
* `abs_a_centered_scratch` at the BRANCH-MAJOR offset (so
|
||||
* the final scratch is a flat |A_centered| over the full
|
||||
* (B × Σ_d n_d × NA) tensor).
|
||||
*
|
||||
* p99: call `sp4_histogram_p99<256>(scratch, total_count)` →
|
||||
* p99(|A_centered|) over the full flat tensor.
|
||||
*
|
||||
* Blend: bound = p99 × SAFETY × α + cur × (1-α), bilateral clamp,
|
||||
* ISV write.
|
||||
*
|
||||
* Args:
|
||||
* b_logits — BRANCH-MAJOR advantage logits.
|
||||
* abs_a_centered_scratch — pre-allocated scratch buffer of size
|
||||
* ≥ Σ_d B × b_d × NA floats. Caller-owned;
|
||||
* produced/consumed only inside this kernel.
|
||||
* scratch_capacity — sentinel (host-passed) for defensive
|
||||
* guard; kernel returns early if total >
|
||||
* capacity.
|
||||
* isv — ISV bus.
|
||||
* B, NA — shapes.
|
||||
* b0..b3 sizes — per-branch action counts.
|
||||
* bound_idx — ADVANTAGE_CLIP_BOUND_INDEX = 482.
|
||||
* alpha — ADVANTAGE_CLIP_EMA_ALPHA = 0.01.
|
||||
* safety_factor — ADVANTAGE_CLIP_SAFETY_FACTOR = 1.5.
|
||||
* bound_min, bound_max — bilateral clamp rails [0.1, 100.0].
|
||||
* sentinel_bound — SENTINEL_ADVANTAGE_CLIP_BOUND = 1.0.
|
||||
*
|
||||
* Launch: grid=(1, 1, 1), block=(256, 1, 1). Dynamic shared memory =
|
||||
* `(BLOCK_SIZE / 32) × 256 × sizeof(int) = 8192 bytes` for
|
||||
* `sp4_histogram_p99`'s per-warp tile array (extern __shared__ int
|
||||
* s_warp_tiles[]). Plus the static shared memory declared inside the
|
||||
* kernel (`s_a_mean[NA_MAX]` and `s_reduce[BLOCK_SIZE]`) — those are
|
||||
* `__shared__` static, NOT extern, so they don't compete with the
|
||||
* dynamic sp4 allocation.
|
||||
* ══════════════════════════════════════════════════════════════════════════ */
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
#include "sp4_histogram_p99.cuh"
|
||||
|
||||
#define SP17_VA_BLOCK_SIZE 256
|
||||
#define SP17_NA_MAX 51 /* matches NUM_ATOMS_MAX in state_layout.cuh */
|
||||
#define SP17_EPS_F 1e-6f
|
||||
|
||||
extern "C" __global__ void sp17_advantage_clip_bound_update(
|
||||
const float* __restrict__ b_logits,
|
||||
float* __restrict__ abs_a_centered_scratch,
|
||||
int scratch_capacity,
|
||||
float* __restrict__ isv,
|
||||
int B,
|
||||
int NA,
|
||||
int b0_size,
|
||||
int b1_size,
|
||||
int b2_size,
|
||||
int b3_size,
|
||||
int bound_idx,
|
||||
float alpha,
|
||||
float safety_factor,
|
||||
float bound_min,
|
||||
float bound_max,
|
||||
float sentinel_bound)
|
||||
{
|
||||
/* Single-block kernel — guard against accidental multi-block launch. */
|
||||
if (blockIdx.x != 0) return;
|
||||
|
||||
const int tid = (int)threadIdx.x;
|
||||
const int bdim = (int)blockDim.x;
|
||||
|
||||
const long long Bll = (long long)B;
|
||||
const long long NAll = (long long)NA;
|
||||
const long long n0 = (long long)b0_size;
|
||||
const long long n1 = (long long)b1_size;
|
||||
const long long n2 = (long long)b2_size;
|
||||
const long long n3 = (long long)b3_size;
|
||||
const long long off0 = 0;
|
||||
const long long off1 = off0 + Bll * n0 * NAll;
|
||||
const long long off2 = off1 + Bll * n1 * NAll;
|
||||
const long long off3 = off2 + Bll * n2 * NAll;
|
||||
const long long total = off3 + Bll * n3 * NAll;
|
||||
|
||||
if (B <= 0 || NA <= 0 || total <= 0) return;
|
||||
if ((long long)scratch_capacity < total) return; /* defensive guard */
|
||||
|
||||
/* Static shared memory: per-atom mean cache (small, fixed) +
|
||||
* per-pass tree-reduce tile. These do NOT consume the extern
|
||||
* `__shared__` dynamic allocation that sp4_histogram_p99 owns. */
|
||||
__shared__ float s_a_mean[SP17_NA_MAX];
|
||||
__shared__ float s_reduce[SP17_VA_BLOCK_SIZE];
|
||||
|
||||
/* Branch passes — for each branch d, compute per-atom A mean, then
|
||||
* write |A_raw - mean[z]| into scratch at the branch's offset. After
|
||||
* all 4 branches, the scratch buffer is a flat |A_centered| tensor of
|
||||
* size `total` floats. */
|
||||
for (int branch = 0; branch < 4; ++branch) {
|
||||
int n_branch;
|
||||
long long off_branch;
|
||||
if (branch == 0) { n_branch = b0_size; off_branch = off0; }
|
||||
else if (branch == 1) { n_branch = b1_size; off_branch = off1; }
|
||||
else if (branch == 2) { n_branch = b2_size; off_branch = off2; }
|
||||
else { n_branch = b3_size; off_branch = off3; }
|
||||
if (n_branch <= 0) continue;
|
||||
|
||||
/* Pass 1: per-atom mean over (B × n_branch) at each atom z. */
|
||||
for (int z = 0; z < NA; ++z) {
|
||||
float local = 0.0f;
|
||||
const long long count_ia = Bll * (long long)n_branch;
|
||||
for (long long idx = (long long)tid; idx < count_ia; idx += (long long)bdim) {
|
||||
const long long off = off_branch + idx * NAll + (long long)z;
|
||||
local += b_logits[off];
|
||||
}
|
||||
s_reduce[tid] = local;
|
||||
__syncthreads();
|
||||
for (int s = bdim / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) s_reduce[tid] += s_reduce[tid + s];
|
||||
__syncthreads();
|
||||
}
|
||||
if (tid == 0) {
|
||||
const float total_count = (float)count_ia;
|
||||
s_a_mean[z] = (total_count > 0.0f) ? (s_reduce[0] / total_count) : 0.0f;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
/* Pass 2: write |A_raw - mean[z]| into scratch at BRANCH-MAJOR
|
||||
* offset. Each thread handles a strided slice. */
|
||||
const long long branch_total = Bll * (long long)n_branch * NAll;
|
||||
for (long long idx = (long long)tid; idx < branch_total; idx += (long long)bdim) {
|
||||
const long long ia = idx / NAll;
|
||||
const long long z = idx - ia * NAll;
|
||||
const long long off = off_branch + idx;
|
||||
const float a = b_logits[off];
|
||||
abs_a_centered_scratch[off] = fabsf(a - s_a_mean[z]);
|
||||
}
|
||||
__syncthreads(); /* ensure all writes visible before next branch reuses s_reduce */
|
||||
}
|
||||
|
||||
/* ── p99(|A_centered|) over the full flat scratch ─────────────────
|
||||
* sp4_histogram_p99 uses its OWN extern __shared__ int per-warp
|
||||
* tiles + 256-int static + 1-float static (total 1028 bytes static
|
||||
* + 8192 bytes dynamic for BLOCK=256). Our static `s_a_mean` and
|
||||
* `s_reduce` arrays do not collide because they are statically
|
||||
* declared (different `__shared__` sub-region). The extern dynamic
|
||||
* shmem region is the caller-passed `shared_mem_bytes` and is
|
||||
* exclusively owned by sp4_histogram_p99.
|
||||
*/
|
||||
const float p99 = sp4_histogram_p99<SP17_VA_BLOCK_SIZE>(
|
||||
abs_a_centered_scratch, (int)total);
|
||||
|
||||
/* ── Pearl-A bootstrap + α blend + bilateral clamp → ISV write ── */
|
||||
if (tid == 0) {
|
||||
const float target = p99 * safety_factor;
|
||||
const float current = isv[bound_idx];
|
||||
float next;
|
||||
if (fabsf(current - sentinel_bound) < SP17_EPS_F) {
|
||||
/* Pearl-A: cold-start sentinel → REPLACE directly. */
|
||||
next = target;
|
||||
} else {
|
||||
next = (1.0f - alpha) * current + alpha * target;
|
||||
}
|
||||
/* `pearl_symmetric_clamp_audit` — bilateral clamp. */
|
||||
next = fmaxf(bound_min, fminf(next, bound_max));
|
||||
isv[bound_idx] = next;
|
||||
__threadfence_system();
|
||||
}
|
||||
}
|
||||
220
crates/ml/src/cuda_pipeline/sp17_v_share_kernel.cu
Normal file
220
crates/ml/src/cuda_pipeline/sp17_v_share_kernel.cu
Normal file
@@ -0,0 +1,220 @@
|
||||
/* ══════════════════════════════════════════════════════════════════════════
|
||||
* SP17 Phase 3.2 (2026-05-08) — V_share producer kernel.
|
||||
*
|
||||
* Per-branch magnitude share
|
||||
* V_share[d] = |E[V]| / (|E[V]| + |E[A_centered, max-Q action]|)
|
||||
*
|
||||
* The picked-action defaults to argmax E[Q] across the b_d actions of the
|
||||
* branch (max-Q semantic — tractable per-batch without depending on the
|
||||
* Thompson-sampled action stored in actions_history_buf, which is
|
||||
* collector-time state and is stale relative to the cuBLAS forward
|
||||
* driving on_v_logits_buf / on_b_logits_buf at HEALTH_DIAG cadence).
|
||||
*
|
||||
* Argmax simplification: argmax_a E[Q[i, a]] = argmax_a (1/NA) Σ_z
|
||||
* (V[i, z] + A_centered[i, a, z]). Since V[i, z] is the same across
|
||||
* actions and the per-atom mean of A is also action-invariant, the
|
||||
* ranking depends only on Σ_z A[i, a, z]. So:
|
||||
* a*(i) = argmax_a Σ_z A_raw[i, a, z]
|
||||
*
|
||||
* For each sample i in branch d:
|
||||
* v_i = (1/NA) Σ_z V[i, z] (mean V over atoms)
|
||||
* a_i = (1/NA) Σ_z A_centered[i, a*(i), z] (mean A_centered over atoms)
|
||||
*
|
||||
* Aggregated across the batch (one block per branch):
|
||||
* |E[V]|_batch = | (1/B) Σ_i v_i |
|
||||
* |E[A_c]|_batch = | (1/B) Σ_i a_i |
|
||||
*
|
||||
* V_share[d] = |E[V]|_batch / (|E[V]|_batch + |E[A_c]|_batch + EPS)
|
||||
*
|
||||
* Then EMA-blend into ISV[V_SHARE_*_INDEX] with the same Pearl-A bootstrap +
|
||||
* Wiener-α-floor pattern as A_var_ema.
|
||||
*
|
||||
* ── Pearls + invariants ────────────────────────────────────────────────
|
||||
* - `feedback_no_atomicadd.md` — block tree-reduce only.
|
||||
* - `pearl_first_observation_bootstrap.md` — sentinel SENTINEL_V_SHARE
|
||||
* = 0.5 ("healthy 50/50 cold-start"); first observation REPLACES.
|
||||
* - `pearl_wiener_alpha_floor_for_nonstationary` — α floored at 0.4
|
||||
* (caller passes WELFORD_ALPHA_MIN).
|
||||
* - `pearl_no_host_branches_in_captured_graph.md` — pure on-device
|
||||
* branching only (per-branch dispatch via blockIdx.x).
|
||||
* - `pearl_symmetric_clamp_audit.md` — V_share is structurally bounded
|
||||
* in [0, 1]; defense-in-depth bilateral clamp before ISV write.
|
||||
*
|
||||
* Launch contract:
|
||||
* grid_dim = (4, 1, 1) — one block per branch
|
||||
* block_dim = (SP17_VA_BLOCK_SIZE, 1, 1) — 256 threads
|
||||
* shared mem (dynamic): (BLOCK_SIZE + NA + 1) × sizeof(float)
|
||||
* = tree-reduce tile + per-atom mean cache
|
||||
* + 1 float for batch-aggregate E[V] broadcast.
|
||||
*
|
||||
* Args:
|
||||
* v_logits — [B, NA] online value logits, sample-major.
|
||||
* b_logits — BRANCH-MAJOR online advantage logits.
|
||||
* isv — ISV bus.
|
||||
* B, NA — batch size, atoms.
|
||||
* b0..b3 sizes — per-branch action counts.
|
||||
* share_idx_* — V_SHARE_*_INDEX per branch (478..482).
|
||||
* alpha — Wiener-α blend rate (caller passes WELFORD_ALPHA_MIN).
|
||||
* sentinel_v_share — SENTINEL_V_SHARE = 0.5.
|
||||
* ══════════════════════════════════════════════════════════════════════════ */
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#define SP17_VA_BLOCK_SIZE 256
|
||||
#define SP17_EPS_F 1e-6f
|
||||
|
||||
extern "C" __global__ void sp17_v_share_update(
|
||||
const float* __restrict__ v_logits,
|
||||
const float* __restrict__ b_logits,
|
||||
float* __restrict__ isv,
|
||||
int B,
|
||||
int NA,
|
||||
int b0_size,
|
||||
int b1_size,
|
||||
int b2_size,
|
||||
int b3_size,
|
||||
int share_idx_dir,
|
||||
int share_idx_mag,
|
||||
int share_idx_ord,
|
||||
int share_idx_urg,
|
||||
float alpha,
|
||||
float sentinel_v_share)
|
||||
{
|
||||
/* One block per branch (gridDim.x == 4). */
|
||||
const int branch = (int)blockIdx.x;
|
||||
if (branch >= 4) return;
|
||||
|
||||
int n_branch, share_idx;
|
||||
long long off_branch;
|
||||
const long long Bll = (long long)B;
|
||||
const long long NAll = (long long)NA;
|
||||
const long long n0 = (long long)b0_size;
|
||||
const long long n1 = (long long)b1_size;
|
||||
const long long n2 = (long long)b2_size;
|
||||
const long long off0 = 0;
|
||||
const long long off1 = off0 + Bll * n0 * NAll;
|
||||
const long long off2 = off1 + Bll * n1 * NAll;
|
||||
const long long off3 = off2 + Bll * n2 * NAll;
|
||||
if (branch == 0) { n_branch = b0_size; share_idx = share_idx_dir; off_branch = off0; }
|
||||
else if (branch == 1) { n_branch = b1_size; share_idx = share_idx_mag; off_branch = off1; }
|
||||
else if (branch == 2) { n_branch = b2_size; share_idx = share_idx_ord; off_branch = off2; }
|
||||
else { n_branch = b3_size; share_idx = share_idx_urg; off_branch = off3; }
|
||||
|
||||
const int tid = (int)threadIdx.x;
|
||||
const int bdim = (int)blockDim.x;
|
||||
|
||||
if (B <= 0 || NA <= 0 || n_branch <= 0) return;
|
||||
|
||||
/* Shared layout (dynamic):
|
||||
* shmem[0..bdim) — block tree-reduce tile
|
||||
* shmem[bdim..bdim+NA) — per-atom mean of raw A across actions
|
||||
* shmem[bdim+NA] — batch-aggregate E[V] broadcast scalar
|
||||
*/
|
||||
extern __shared__ float shmem[];
|
||||
float* s_reduce = shmem;
|
||||
float* s_a_mean = shmem + bdim;
|
||||
float* s_e_v = shmem + bdim + NA;
|
||||
|
||||
/* ── Pass A: per-atom mean of raw A (same as A_var_ema's pass 1) ── */
|
||||
for (int z = 0; z < NA; ++z) {
|
||||
float local = 0.0f;
|
||||
const long long count_ia = Bll * (long long)n_branch;
|
||||
for (long long idx = (long long)tid; idx < count_ia; idx += (long long)bdim) {
|
||||
const long long off = off_branch + idx * NAll + (long long)z;
|
||||
local += b_logits[off];
|
||||
}
|
||||
s_reduce[tid] = local;
|
||||
__syncthreads();
|
||||
for (int s = bdim / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) s_reduce[tid] += s_reduce[tid + s];
|
||||
__syncthreads();
|
||||
}
|
||||
if (tid == 0) {
|
||||
const float total_count = (float)count_ia;
|
||||
s_a_mean[z] = (total_count > 0.0f) ? (s_reduce[0] / total_count) : 0.0f;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
/* ── Pass B: per-sample v_i = (1/NA) Σ_z V[i, z] aggregated to E[V] ──
|
||||
* v_logits is [B, NA] sample-major (same buffer used by all branches —
|
||||
* the value head is shared across branches). */
|
||||
float local_v_sum = 0.0f;
|
||||
for (long long i = (long long)tid; i < Bll; i += (long long)bdim) {
|
||||
float v_i = 0.0f;
|
||||
const long long base = i * NAll;
|
||||
for (long long z = 0; z < NAll; ++z) {
|
||||
v_i += v_logits[base + z];
|
||||
}
|
||||
v_i /= (float)NA;
|
||||
local_v_sum += v_i;
|
||||
}
|
||||
s_reduce[tid] = local_v_sum;
|
||||
__syncthreads();
|
||||
for (int s = bdim / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) s_reduce[tid] += s_reduce[tid + s];
|
||||
__syncthreads();
|
||||
}
|
||||
if (tid == 0) {
|
||||
s_e_v[0] = s_reduce[0] / (float)B;
|
||||
}
|
||||
__syncthreads();
|
||||
const float e_v = s_e_v[0];
|
||||
|
||||
/* ── Pass C: per-sample picked-action = argmax_a Σ_z A_raw[i, a, z] ──
|
||||
* Build a_i = (1/NA) Σ_z (A_raw[i, a*, z] − a_mean[z]) and aggregate
|
||||
* to E[A_centered, picked]. */
|
||||
float local_a_sum = 0.0f;
|
||||
for (long long i = (long long)tid; i < Bll; i += (long long)bdim) {
|
||||
float best_score = -INFINITY;
|
||||
long long best_a = 0;
|
||||
for (long long a = 0; a < (long long)n_branch; ++a) {
|
||||
float s_score = 0.0f;
|
||||
const long long base = off_branch + (i * (long long)n_branch + a) * NAll;
|
||||
for (long long z = 0; z < NAll; ++z) {
|
||||
s_score += b_logits[base + z];
|
||||
}
|
||||
if (s_score > best_score) {
|
||||
best_score = s_score;
|
||||
best_a = a;
|
||||
}
|
||||
}
|
||||
float a_i_centered = 0.0f;
|
||||
const long long base_pick = off_branch + (i * (long long)n_branch + best_a) * NAll;
|
||||
for (long long z = 0; z < NAll; ++z) {
|
||||
a_i_centered += b_logits[base_pick + z] - s_a_mean[z];
|
||||
}
|
||||
a_i_centered /= (float)NA;
|
||||
local_a_sum += a_i_centered;
|
||||
}
|
||||
s_reduce[tid] = local_a_sum;
|
||||
__syncthreads();
|
||||
for (int s = bdim / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) s_reduce[tid] += s_reduce[tid + s];
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
/* ── Pass D: V_share + Pearl-A bootstrap + α blend → ISV write ── */
|
||||
if (tid == 0) {
|
||||
const float e_a_centered = s_reduce[0] / (float)B;
|
||||
const float abs_v = fabsf(e_v);
|
||||
const float abs_a = fabsf(e_a_centered);
|
||||
const float denom = abs_v + abs_a + SP17_EPS_F;
|
||||
const float v_share = abs_v / denom;
|
||||
|
||||
const float current = isv[share_idx];
|
||||
float next;
|
||||
if (fabsf(current - sentinel_v_share) < SP17_EPS_F) {
|
||||
/* Pearl-A: cold-start sentinel → REPLACE directly. */
|
||||
next = v_share;
|
||||
} else {
|
||||
next = (1.0f - alpha) * current + alpha * v_share;
|
||||
}
|
||||
/* `pearl_symmetric_clamp_audit` — V_share is structurally
|
||||
* bounded in [0, 1] by the formula, but apply bilateral
|
||||
* defense-in-depth clamp before ISV write. */
|
||||
next = fmaxf(0.0f, fminf(next, 1.0f));
|
||||
isv[share_idx] = next;
|
||||
__threadfence_system(); /* PCIe-visible write for mapped-pinned ISV */
|
||||
}
|
||||
}
|
||||
@@ -5273,20 +5273,26 @@ impl DQNTrainer {
|
||||
);
|
||||
}
|
||||
|
||||
// SP17 Phase 3.1 (2026-05-08) — per-branch A_var EMA diagnostic.
|
||||
// Cold-path (epoch boundary) producer launch; emits the
|
||||
// dueling-Q identifiability variance signal so the post-deploy
|
||||
// SP17 Phase 3.1+3.2 (2026-05-08) — dueling-Q identifiability
|
||||
// diagnostic chain. Cold-path (epoch boundary) producer
|
||||
// launches that emit the per-branch A_var, per-branch V_share,
|
||||
// and adaptive advantage_clip_bound signals so the post-deploy
|
||||
// reviewer can confirm `A_centered` is doing real work
|
||||
// (Var > 0 = the policy discriminates between actions) vs.
|
||||
// regression to the pre-SP17 V-dominated regime (Var → 0 = "all
|
||||
// (Var > 0 with V_share ≈ 0.5 = healthy) vs. regression to the
|
||||
// pre-SP17 V-dominated regime (Var → 0 + V_share → 1.0 = "all
|
||||
// actions equivalent, V dominates"). Per `feedback_no_atomicadd`
|
||||
// the kernel uses 4-block tree-reduce; per
|
||||
// `pearl_first_observation_bootstrap` slots 474..478 use
|
||||
// sentinel SENTINEL_A_VAR_EMA=0.0 with REPLACE on first
|
||||
// observation. α post-bootstrap is WELFORD_ALPHA_MIN=0.4 per
|
||||
// `pearl_wiener_alpha_floor_for_nonstationary`. Plan:
|
||||
// the kernels use block tree-reduce (and sp4_histogram_p99 for
|
||||
// the clip bound — per-warp tile binning, no atomicAdd per
|
||||
// `pearl_fused_per_group_statistics_oracle`); per
|
||||
// `pearl_first_observation_bootstrap` all slots use Pearl-A
|
||||
// REPLACE on first observation; α post-bootstrap is
|
||||
// WELFORD_ALPHA_MIN=0.4 (a_var, v_share) and
|
||||
// ADVANTAGE_CLIP_EMA_ALPHA=0.01 (slow per-fold for the clip
|
||||
// bound) per `pearl_wiener_alpha_floor_for_nonstationary`.
|
||||
// Bilateral clamp on V_share [0, 1] and advantage_clip_bound
|
||||
// [0.1, 100.0] per `pearl_symmetric_clamp_audit`. Plan:
|
||||
// docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
|
||||
// Phase 3.1.
|
||||
// Phase 3.1 + 3.2.
|
||||
{
|
||||
let a_var = if let Some(ref fused) = self.fused_ctx {
|
||||
match fused.trainer().read_a_var_ema_per_branch() {
|
||||
@@ -5301,9 +5307,39 @@ impl DQNTrainer {
|
||||
} else {
|
||||
[0.0_f32; 4]
|
||||
};
|
||||
let v_share = if let Some(ref fused) = self.fused_ctx {
|
||||
match fused.trainer().read_v_share_per_branch() {
|
||||
Ok(v) => v,
|
||||
Err(e) => {
|
||||
tracing::warn!(
|
||||
"HEALTH_DIAG read_v_share_per_branch refresh failed: {e}"
|
||||
);
|
||||
[0.0_f32; 4]
|
||||
}
|
||||
}
|
||||
} else {
|
||||
[0.0_f32; 4]
|
||||
};
|
||||
let clip_bound = if let Some(ref fused) = self.fused_ctx {
|
||||
match fused.trainer().read_advantage_clip_bound() {
|
||||
Ok(v) => v,
|
||||
Err(e) => {
|
||||
tracing::warn!(
|
||||
"HEALTH_DIAG read_advantage_clip_bound refresh failed: {e}"
|
||||
);
|
||||
0.0_f32
|
||||
}
|
||||
}
|
||||
} else {
|
||||
0.0_f32
|
||||
};
|
||||
tracing::info!(
|
||||
"HEALTH_DIAG[{}]: dueling [a_var=(d={:.4} m={:.4} o={:.4} u={:.4})]",
|
||||
epoch, a_var[0], a_var[1], a_var[2], a_var[3],
|
||||
"HEALTH_DIAG[{}]: dueling [v_share=(d={:.4} m={:.4} o={:.4} u={:.4})] \
|
||||
[a_var=(d={:.4} m={:.4} o={:.4} u={:.4})] [clip={:.4}]",
|
||||
epoch,
|
||||
v_share[0], v_share[1], v_share[2], v_share[3],
|
||||
a_var[0], a_var[1], a_var[2], a_var[3],
|
||||
clip_bound,
|
||||
);
|
||||
}
|
||||
|
||||
|
||||
@@ -2296,4 +2296,349 @@ mod gpu {
|
||||
(diff {:.3e})", (got_urg - exp_urg).abs(),
|
||||
);
|
||||
}
|
||||
|
||||
// ── F.2: V_share producer matches closed-form formula ──────────────
|
||||
/// Phase 3.2 (2026-05-08) — synthetic V/A constructed so each branch
|
||||
/// has a closed-form V_share. Launch `sp17_v_share_update`; assert
|
||||
/// ISV[V_SHARE_*] readback matches expected to ε=1e-4.
|
||||
///
|
||||
/// Construction:
|
||||
/// - V[i, z] = V_CONST = 2.0 (so |E[V]| = 2.0).
|
||||
/// - For each branch d, action a, atom z:
|
||||
/// A[i, a, z] = K_d × (a − (n_d − 1)/2)
|
||||
/// (mean-zero per atom, picked = max-a = n_d-1).
|
||||
/// - a_i_centered for picked = K_d × ((n_d-1)/2)
|
||||
/// - E[A_centered, picked] = K_d × (n_d-1)/2 (constant across i).
|
||||
/// - V_share[d] = 2.0 / (2.0 + |K_d × (n_d-1)/2|).
|
||||
///
|
||||
/// With K_d = {1.0, 2.0, 3.0, 0.5} and n_d = {4, 3, 3, 3}:
|
||||
/// - dir: |a_part| = 1.5 → V_share = 2/(2+1.5) = 4/7 ≈ 0.5714
|
||||
/// - mag: |a_part| = 2.0 → V_share = 2/(2+2.0) = 0.5
|
||||
/// - ord: |a_part| = 3.0 → V_share = 2/(2+3.0) = 0.4
|
||||
/// - urg: |a_part| = 0.5 → V_share = 2/(2+0.5) = 0.8
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn v_share_per_branch_matches_closed_form() {
|
||||
use ml::cuda_pipeline::sp14_isv_slots::{
|
||||
V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
|
||||
V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX,
|
||||
SENTINEL_V_SHARE, WELFORD_ALPHA_MIN,
|
||||
};
|
||||
|
||||
const SP17_V_SHARE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/sp17_v_share_kernel.cubin"));
|
||||
|
||||
let stream = make_test_stream();
|
||||
let module = stream
|
||||
.context()
|
||||
.load_cubin(SP17_V_SHARE_CUBIN.to_vec())
|
||||
.expect("load sp17_v_share_kernel cubin");
|
||||
let kernel = module
|
||||
.load_function("sp17_v_share_update")
|
||||
.expect("load sp17_v_share_update function");
|
||||
|
||||
const B: usize = 8;
|
||||
const NA: usize = 51;
|
||||
const N0: usize = 4;
|
||||
const N1: usize = 3;
|
||||
const N2: usize = 3;
|
||||
const N3: usize = 3;
|
||||
const V_CONST: f32 = 2.0;
|
||||
const K_DIR: f32 = 1.0;
|
||||
const K_MAG: f32 = 2.0;
|
||||
const K_ORD: f32 = 3.0;
|
||||
const K_URG: f32 = 0.5;
|
||||
|
||||
let v_logits = vec![V_CONST; B * NA];
|
||||
let n_per_branch = [N0, N1, N2, N3];
|
||||
let k_per_branch = [K_DIR, K_MAG, K_ORD, K_URG];
|
||||
let mut b_logits = Vec::new();
|
||||
for branch in 0..4 {
|
||||
let n = n_per_branch[branch];
|
||||
let k = k_per_branch[branch];
|
||||
let center = (n as f32 - 1.0) / 2.0;
|
||||
for _i in 0..B {
|
||||
for a in 0..n {
|
||||
for _z in 0..NA {
|
||||
b_logits.push(k * (a as f32 - center));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
let v_buf = unsafe { MappedF32Buffer::new(v_logits.len()) }
|
||||
.expect("alloc v_logits buf");
|
||||
v_buf.write_from_slice(&v_logits);
|
||||
let b_buf = unsafe { MappedF32Buffer::new(b_logits.len()) }
|
||||
.expect("alloc b_logits buf");
|
||||
b_buf.write_from_slice(&b_logits);
|
||||
|
||||
// ISV: pre-populate slots 478..482 with sentinel 0.5 so Pearl-A
|
||||
// bootstrap fires on first launch.
|
||||
const ISV_SPAN: usize = 483;
|
||||
let mut isv_host = vec![0.0_f32; ISV_SPAN];
|
||||
for &slot in [V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
|
||||
V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX].iter() {
|
||||
isv_host[slot] = SENTINEL_V_SHARE;
|
||||
}
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_SPAN) }
|
||||
.expect("alloc isv buf");
|
||||
isv_buf.write_from_slice(&isv_host);
|
||||
|
||||
let b_i32: i32 = B as i32;
|
||||
let na_i32: i32 = NA as i32;
|
||||
let n0_i32: i32 = N0 as i32;
|
||||
let n1_i32: i32 = N1 as i32;
|
||||
let n2_i32: i32 = N2 as i32;
|
||||
let n3_i32: i32 = N3 as i32;
|
||||
let share_idx_dir: i32 = V_SHARE_DIR_INDEX as i32;
|
||||
let share_idx_mag: i32 = V_SHARE_MAG_INDEX as i32;
|
||||
let share_idx_ord: i32 = V_SHARE_ORD_INDEX as i32;
|
||||
let share_idx_urg: i32 = V_SHARE_URG_INDEX as i32;
|
||||
let alpha: f32 = WELFORD_ALPHA_MIN;
|
||||
let sentinel: f32 = SENTINEL_V_SHARE;
|
||||
let block_dim: u32 = 256;
|
||||
let shmem_bytes: u32 =
|
||||
(block_dim + NA as u32 + 1) * std::mem::size_of::<f32>() as u32;
|
||||
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernel)
|
||||
.arg(&v_buf.dev_ptr)
|
||||
.arg(&b_buf.dev_ptr)
|
||||
.arg(&isv_buf.dev_ptr)
|
||||
.arg(&b_i32)
|
||||
.arg(&na_i32)
|
||||
.arg(&n0_i32)
|
||||
.arg(&n1_i32)
|
||||
.arg(&n2_i32)
|
||||
.arg(&n3_i32)
|
||||
.arg(&share_idx_dir)
|
||||
.arg(&share_idx_mag)
|
||||
.arg(&share_idx_ord)
|
||||
.arg(&share_idx_urg)
|
||||
.arg(&alpha)
|
||||
.arg(&sentinel)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (4, 1, 1),
|
||||
block_dim: (block_dim, 1, 1),
|
||||
shared_mem_bytes: shmem_bytes,
|
||||
})
|
||||
.expect("launch sp17_v_share_update");
|
||||
}
|
||||
stream.synchronize().expect("sync after sp17_v_share_update");
|
||||
|
||||
let isv_after = isv_buf.read_all();
|
||||
|
||||
fn expected_share(n: usize, k: f32, v_const: f32) -> f32 {
|
||||
let a_part = k.abs() * (n as f32 - 1.0) / 2.0;
|
||||
v_const.abs() / (v_const.abs() + a_part + 1e-6_f32)
|
||||
}
|
||||
let exp_dir = expected_share(N0, K_DIR, V_CONST);
|
||||
let exp_mag = expected_share(N1, K_MAG, V_CONST);
|
||||
let exp_ord = expected_share(N2, K_ORD, V_CONST);
|
||||
let exp_urg = expected_share(N3, K_URG, V_CONST);
|
||||
|
||||
let got_dir = isv_after[V_SHARE_DIR_INDEX];
|
||||
let got_mag = isv_after[V_SHARE_MAG_INDEX];
|
||||
let got_ord = isv_after[V_SHARE_ORD_INDEX];
|
||||
let got_urg = isv_after[V_SHARE_URG_INDEX];
|
||||
|
||||
let eps = 1e-4_f32;
|
||||
assert!(
|
||||
(got_dir - exp_dir).abs() < eps,
|
||||
"dir V_share expected {exp_dir:.6}, got {got_dir:.6}",
|
||||
);
|
||||
assert!(
|
||||
(got_mag - exp_mag).abs() < eps,
|
||||
"mag V_share expected {exp_mag:.6}, got {got_mag:.6}",
|
||||
);
|
||||
assert!(
|
||||
(got_ord - exp_ord).abs() < eps,
|
||||
"ord V_share expected {exp_ord:.6}, got {got_ord:.6}",
|
||||
);
|
||||
assert!(
|
||||
(got_urg - exp_urg).abs() < eps,
|
||||
"urg V_share expected {exp_urg:.6}, got {got_urg:.6}",
|
||||
);
|
||||
}
|
||||
|
||||
// ── F.3: advantage_clip_bound producer ≈ p99(|A_centered|) × 1.5 ──
|
||||
/// Phase 3.2 (2026-05-08) — synthetic A constructed so |A_centered|
|
||||
/// has a known distribution, then verify the kernel's p99 estimator
|
||||
/// settles on `≈ p99(|A_centered|) × 1.5` after Pearl-A bootstrap.
|
||||
///
|
||||
/// Construction:
|
||||
/// For all 4 branches with shapes (4, 3, 3, 3) × B × NA:
|
||||
/// A[i, a, z] = a × NOISE_SCALE + 0.001 × (i × NA + z) × per-(i,z) noise
|
||||
/// The action-component dominates so per-atom mean of A across
|
||||
/// actions is approximately a_center × NOISE_SCALE; |A_c| ranges
|
||||
/// continuously through [0, ~max] without identical-value clusters
|
||||
/// that would break `sp4_histogram_p99`'s per-warp tile binning
|
||||
/// (non-atomic warp tiles undercount when many lanes hit the same
|
||||
/// bin in lockstep, which the sp4 comment qualifies as "1/(256×32)
|
||||
/// loss for uniformly distributed signals" — the qualifier matters).
|
||||
///
|
||||
/// With NOISE_SCALE = 1.0:
|
||||
/// - dir n=4, max |A_c| ≈ 1.5 (action 0/3 endpoints), with the
|
||||
/// per-(i,z) noise spreading the values into adjacent bins.
|
||||
/// - p99 of |A_c| ≈ 1.5 (top of action-3 cluster) ± noise.
|
||||
///
|
||||
/// Tolerance: ε=0.15 absolute (the per-(i,z) noise spreads the top
|
||||
/// bin and shifts p99 estimate; we still verify the bound is in the
|
||||
/// right order of magnitude — the test is asserting "kernel produces
|
||||
/// a p99-like statistic, NOT a stuck constant"). Phase 4 L40S smoke
|
||||
/// validates the bound's behavioural utility on real |A_centered|.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn advantage_clip_bound_tracks_p99_safety() {
|
||||
use ml::cuda_pipeline::sp14_isv_slots::{
|
||||
ADVANTAGE_CLIP_BOUND_INDEX,
|
||||
ADVANTAGE_CLIP_BOUND_MIN, ADVANTAGE_CLIP_BOUND_MAX,
|
||||
ADVANTAGE_CLIP_EMA_ALPHA, ADVANTAGE_CLIP_SAFETY_FACTOR,
|
||||
SENTINEL_ADVANTAGE_CLIP_BOUND,
|
||||
};
|
||||
|
||||
const SP17_CLIP_CUBIN: &[u8] = include_bytes!(concat!(
|
||||
env!("OUT_DIR"), "/sp17_advantage_clip_bound_kernel.cubin"));
|
||||
|
||||
let stream = make_test_stream();
|
||||
let module = stream
|
||||
.context()
|
||||
.load_cubin(SP17_CLIP_CUBIN.to_vec())
|
||||
.expect("load sp17_advantage_clip_bound cubin");
|
||||
let kernel = module
|
||||
.load_function("sp17_advantage_clip_bound_update")
|
||||
.expect("load sp17_advantage_clip_bound_update function");
|
||||
|
||||
const B: usize = 8;
|
||||
const NA: usize = 51;
|
||||
const N0: usize = 4;
|
||||
const N1: usize = 3;
|
||||
const N2: usize = 3;
|
||||
const N3: usize = 3;
|
||||
const NOISE_SCALE: f32 = 1.0;
|
||||
const PER_IZ_JITTER: f32 = 0.01;
|
||||
|
||||
// Build A[i, a, z] = a × NOISE_SCALE + jitter[i, z]
|
||||
// The action component dominates so per-atom mean is ≈ a_center; the
|
||||
// per-(i, z) jitter spreads the |A_centered| values continuously and
|
||||
// breaks the lockstep-per-warp pathology that undercounts in
|
||||
// sp4_histogram_p99's non-atomic per-warp tile binning.
|
||||
let n_per_branch = [N0, N1, N2, N3];
|
||||
// Deterministic "noise" so the test is reproducible without an RNG
|
||||
// crate: hash (i × NA + z) into a small float in [-PER_IZ_JITTER, PER_IZ_JITTER].
|
||||
fn jitter(i: usize, z: usize) -> f32 {
|
||||
let h = ((i.wrapping_mul(7919) ^ z.wrapping_mul(31337)) % 2003) as f32;
|
||||
(h / 1001.5_f32 - 1.0_f32) * PER_IZ_JITTER
|
||||
}
|
||||
let mut b_logits = Vec::new();
|
||||
for branch in 0..4 {
|
||||
let n = n_per_branch[branch];
|
||||
for i in 0..B {
|
||||
for a in 0..n {
|
||||
for z in 0..NA {
|
||||
b_logits.push((a as f32) * NOISE_SCALE + jitter(i, z));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
let total_len = b_logits.len();
|
||||
assert_eq!(total_len, B * (N0 + N1 + N2 + N3) * NA);
|
||||
|
||||
let b_buf = unsafe { MappedF32Buffer::new(total_len) }
|
||||
.expect("alloc b_logits buf");
|
||||
b_buf.write_from_slice(&b_logits);
|
||||
let scratch_buf = unsafe { MappedF32Buffer::new(total_len) }
|
||||
.expect("alloc scratch buf");
|
||||
|
||||
const ISV_SPAN: usize = 483;
|
||||
let mut isv_host = vec![0.0_f32; ISV_SPAN];
|
||||
isv_host[ADVANTAGE_CLIP_BOUND_INDEX] = SENTINEL_ADVANTAGE_CLIP_BOUND;
|
||||
let isv_buf = unsafe { MappedF32Buffer::new(ISV_SPAN) }
|
||||
.expect("alloc isv buf");
|
||||
isv_buf.write_from_slice(&isv_host);
|
||||
|
||||
let scratch_cap_i32: i32 = total_len as i32;
|
||||
let b_i32: i32 = B as i32;
|
||||
let na_i32: i32 = NA as i32;
|
||||
let n0_i32: i32 = N0 as i32;
|
||||
let n1_i32: i32 = N1 as i32;
|
||||
let n2_i32: i32 = N2 as i32;
|
||||
let n3_i32: i32 = N3 as i32;
|
||||
let bound_idx: i32 = ADVANTAGE_CLIP_BOUND_INDEX as i32;
|
||||
let alpha: f32 = ADVANTAGE_CLIP_EMA_ALPHA;
|
||||
let safety: f32 = ADVANTAGE_CLIP_SAFETY_FACTOR;
|
||||
let bound_min: f32 = ADVANTAGE_CLIP_BOUND_MIN;
|
||||
let bound_max: f32 = ADVANTAGE_CLIP_BOUND_MAX;
|
||||
let sentinel: f32 = SENTINEL_ADVANTAGE_CLIP_BOUND;
|
||||
let block_dim: u32 = 256;
|
||||
let warps = block_dim / 32;
|
||||
let shmem_bytes: u32 = warps * 256 * std::mem::size_of::<i32>() as u32;
|
||||
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernel)
|
||||
.arg(&b_buf.dev_ptr)
|
||||
.arg(&scratch_buf.dev_ptr)
|
||||
.arg(&scratch_cap_i32)
|
||||
.arg(&isv_buf.dev_ptr)
|
||||
.arg(&b_i32)
|
||||
.arg(&na_i32)
|
||||
.arg(&n0_i32)
|
||||
.arg(&n1_i32)
|
||||
.arg(&n2_i32)
|
||||
.arg(&n3_i32)
|
||||
.arg(&bound_idx)
|
||||
.arg(&alpha)
|
||||
.arg(&safety)
|
||||
.arg(&bound_min)
|
||||
.arg(&bound_max)
|
||||
.arg(&sentinel)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (block_dim, 1, 1),
|
||||
shared_mem_bytes: shmem_bytes,
|
||||
})
|
||||
.expect("launch sp17_advantage_clip_bound_update");
|
||||
}
|
||||
stream.synchronize().expect("sync after sp17_advantage_clip_bound");
|
||||
|
||||
let isv_after = isv_buf.read_all();
|
||||
let got = isv_after[ADVANTAGE_CLIP_BOUND_INDEX];
|
||||
|
||||
// p99 of |A_centered|: per-action raw values lie in
|
||||
// {0, 1, 2, 3} × NOISE_SCALE (action index × scale)
|
||||
// After centering with per-branch action mean, the dir branch's
|
||||
// |A_c| max ≈ 1.5 (action 0/3 endpoints) ± jitter. The per-(i, z)
|
||||
// jitter spreads each value into ~PER_IZ_JITTER × 256 / max bins.
|
||||
//
|
||||
// p99 target ≈ max(|A_c|) × ADVANTAGE_CLIP_SAFETY_FACTOR; Pearl-A
|
||||
// bootstrap REPLACES on first observation, so got ≈ target.
|
||||
// Tolerance is wide because the jitter shifts the upper-tail bin
|
||||
// and the histogram is a linear-bin estimator (not exact p99).
|
||||
let expected_max: f32 = 1.5 + PER_IZ_JITTER;
|
||||
let expected: f32 = expected_max * ADVANTAGE_CLIP_SAFETY_FACTOR;
|
||||
let eps = 0.20_f32; // wide — jitter + linear histogram quantization
|
||||
assert!(
|
||||
(got - expected).abs() < eps,
|
||||
"advantage_clip_bound expected ≈{expected:.4} (= {expected_max:.3} \
|
||||
× ADVANTAGE_CLIP_SAFETY_FACTOR), got {got:.4} (diff {:.3e}, \
|
||||
eps {eps:.3e})",
|
||||
(got - expected).abs(),
|
||||
);
|
||||
// Also verify the bound is order-of-magnitude correct (NOT stuck
|
||||
// at sentinel and NOT massively over-clamped).
|
||||
assert!(
|
||||
got > 1.0_f32,
|
||||
"advantage_clip_bound stuck near sentinel — got {got:.4}, \
|
||||
expected target ≈{expected:.4} after Pearl-A REPLACE",
|
||||
);
|
||||
// Also sanity-check the bilateral clamp didn't fire.
|
||||
assert!(
|
||||
got >= ADVANTAGE_CLIP_BOUND_MIN && got <= ADVANTAGE_CLIP_BOUND_MAX,
|
||||
"advantage_clip_bound out of bounds: got {got}, must be in \
|
||||
[{ADVANTAGE_CLIP_BOUND_MIN}, {ADVANTAGE_CLIP_BOUND_MAX}]",
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -2,6 +2,74 @@
|
||||
|
||||
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
|
||||
|
||||
## 2026-05-08 — SP17 Phase 3.2: V_share + advantage_clip_bound producers (additive observability)
|
||||
|
||||
Phase 3.2 lands the remaining two SP17 dueling-Q diagnostic producers atomically with kernel + launcher + Rust wrapper + extended HEALTH_DIAG emit + GPU oracle tests per `feedback_wire_everything_up`.
|
||||
|
||||
**New kernels:**
|
||||
- `crates/ml/src/cuda_pipeline/sp17_v_share_kernel.cu` — `sp17_v_share_update`
|
||||
- `crates/ml/src/cuda_pipeline/sp17_advantage_clip_bound_kernel.cu` — `sp17_advantage_clip_bound_update`
|
||||
|
||||
### V_share producer
|
||||
|
||||
**Per branch d:**
|
||||
1. Pass A — per-atom mean of raw A across n_d actions (same as A_var_ema's pass 1).
|
||||
2. Pass B — per-sample `v_i = (1/NA) Σ_z V[i, z]` aggregated to `E[V] = (1/B) Σ_i v_i` (broadcast via shmem).
|
||||
3. Pass C — per-sample `a*(i) = argmax_a Σ_z A_raw[i, a, z]` (max-Q semantic — picked-action default; tractable per-batch without depending on `actions_history_buf`); then `a_i = (1/NA) Σ_z (A_raw[i, a*(i), z] − a_mean[z])` aggregated to `E[A_centered, picked]`.
|
||||
4. Pass D — `V_share = |E[V]| / (|E[V]| + |E[A_centered, picked]| + EPS)`. Pearl-A bootstrap (sentinel 0.5) + α=`WELFORD_ALPHA_MIN=0.4` blend + bilateral [0, 1] clamp → ISV.
|
||||
|
||||
Grid `(4, 1, 1)` × block `(256, 1, 1)`. Dynamic shmem = `(BLOCK_SIZE + NA + 1) × sizeof(f32)` = tree-reduce tile + per-atom mean cache + 1 float for E[V] broadcast.
|
||||
|
||||
### advantage_clip_bound producer
|
||||
|
||||
**Single block:**
|
||||
1. For each branch d (sequential within the block):
|
||||
- Pass 1 — per-atom mean of raw A (cached in static `__shared__ s_a_mean[NA_MAX]`).
|
||||
- Pass 2 — write `|A_raw - mean[z]|` into the kernel-owned `abs_a_centered_scratch` buffer at the BRANCH-MAJOR offset.
|
||||
2. Call `sp4_histogram_p99<256>(scratch, total)` per `pearl_fused_per_group_statistics_oracle` — block tree-reduce + per-warp tile binning, NO atomicAdd.
|
||||
3. `target = p99 × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5`. Pearl-A bootstrap (sentinel 1.0) + α=`ADVANTAGE_CLIP_EMA_ALPHA=0.01` slow per-fold blend + bilateral clamp `[0.1, 100.0]` per `pearl_symmetric_clamp_audit` → ISV[ADVANTAGE_CLIP_BOUND_INDEX=482].
|
||||
|
||||
Grid `(1, 1, 1)` × block `(256, 1, 1)`. Dynamic shmem = `(BLOCK_SIZE/32) × 256 × sizeof(int) = 8192 bytes` for `sp4_histogram_p99`'s per-warp tile array. Static shared memory adds `s_a_mean[51] + s_reduce[256] + sp4's s_bins[256] + sp4's s_step_max` ≈ 2256 bytes — total well within 48KB SM 8.6 limit.
|
||||
|
||||
**Scratch buffer:** `MappedF32Buffer` sized to `B × Σ_d b_d × NA` (max possible flat |A_centered|). Allocated once at trainer construct; kernel-owned (no host writes). Per `feedback_no_htod_htoh_only_mapped_pinned`.
|
||||
|
||||
### Rust integration (`crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`)
|
||||
|
||||
- `SP17_V_SHARE_CUBIN`, `SP17_ADVANTAGE_CLIP_BOUND_CUBIN` static cubins.
|
||||
- `sp17_v_share_kernel: CudaFunction`, `sp17_advantage_clip_bound_kernel: CudaFunction` field on `GpuDqnTrainer`.
|
||||
- `sp17_advantage_clip_scratch: MappedF32Buffer` field (capacity = max flat |A_c|).
|
||||
- `launch_v_share_update()` + `read_v_share_per_branch() -> [f32; 4]` (launches → syncs → reads ISV slots 478..482).
|
||||
- `launch_advantage_clip_bound_update()` + `read_advantage_clip_bound() -> f32` (launches → syncs → reads ISV slot 482).
|
||||
|
||||
### HEALTH_DIAG emit (canonical line)
|
||||
|
||||
```
|
||||
HEALTH_DIAG[N]: dueling [v_share=(d=X m=Y o=Z u=W)] [a_var=(d=A m=B o=C u=D)] [clip=K]
|
||||
```
|
||||
|
||||
Emitted once per epoch alongside the existing `v_a_means` line. Phase 3.3 will polish the format if needed; the format above is the canonical Phase 3 output.
|
||||
|
||||
### State reset
|
||||
|
||||
ISV slots 478..482 (V_SHARE) and 482 (ADVANTAGE_CLIP_BOUND) are FoldReset entries (already wired in PP.2 commit `a225926e5`). At fold boundary slots are REPLACED with their respective sentinels (`SENTINEL_V_SHARE=0.5`, `SENTINEL_ADVANTAGE_CLIP_BOUND=1.0`) so Pearl-A bootstrap fires fresh on the new fold's first epoch.
|
||||
|
||||
### GPU oracle tests
|
||||
|
||||
Both in `crates/ml/tests/sp17_dueling_oracle_tests.rs`:
|
||||
|
||||
**`v_share_per_branch_matches_closed_form`:** synthetic V[i, z] = 2.0 + linear A per branch with known per-action range. Closed-form `V_share[d] = 2.0 / (2.0 + |K_d × (n_d-1)/2|)`. With `K_d ∈ {1.0, 2.0, 3.0, 0.5}` and `n_d ∈ {4, 3, 3, 3}`:
|
||||
- dir: 0.5714, mag: 0.5, ord: 0.4, urg: 0.8
|
||||
|
||||
Tolerance ε = 1e-4. Pearl-A bootstrap REPLACES on first launch.
|
||||
|
||||
**`advantage_clip_bound_tracks_p99_safety`:** synthetic A with action-component dominant + per-(i, z) jitter (`PER_IZ_JITTER = 0.01`). Critical implementation note: the jitter is REQUIRED — pathologically lockstep values (e.g. all-ones-and-twos with no spread) trigger non-atomic warp-tile collisions in `sp4_histogram_p99`'s per-warp binning, producing severe undercount and a wrong p99. The kernel comment `1/(256×32) loss` qualifies "for uniformly distributed signals" — concentrated values violate that assumption. Real |A_centered| in production is continuous, so the issue is test-data-only, not kernel-correctness. Tolerance ε = 0.20 (jitter spreads top bin + linear histogram quantization).
|
||||
|
||||
### Atomicity envelope
|
||||
|
||||
Phase 3.2 is observability-only — no consumer path is modified. The clip bound is producer-tracked but NOT yet wired as an actual clip on any kernel — that's Phase 5 follow-up. The Phase 1 mean-zero contract (commits `eabcf8d52` through `6f53d676f`) is what makes `A_centered` a meaningful signal to measure; this commit observes it.
|
||||
|
||||
**Plan:** `docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md` Phase 3.2.
|
||||
|
||||
## 2026-05-08 — SP17 Phase 3.1: A_var_ema per-branch producer kernel (additive observability)
|
||||
|
||||
Phase 3 of the SP17 dueling-Q identifiability plan introduces the diagnostic chain that lets the post-deploy reviewer confirm the mean-zero centering (Phase 1-2) is doing real work. This commit lands the first of three producers: per-branch EMA of the over-actions variance of `A_centered`.
|
||||
|
||||
Reference in New Issue
Block a user