diff --git a/crates/ml/src/cuda_pipeline/sp20_stats_compute.rs b/crates/ml/src/cuda_pipeline/sp20_stats_compute.rs index f194f387a..dbad2271a 100644 --- a/crates/ml/src/cuda_pipeline/sp20_stats_compute.rs +++ b/crates/ml/src/cuda_pipeline/sp20_stats_compute.rs @@ -2,18 +2,26 @@ //! SP20 Phase 1.1 (2026-05-09) — `sp20_stats_compute_kernel` launcher. //! //! Component 5 / Kernel 3 of the SP20 design. Produces two scalar f32 -//! outputs from a `[B, 3]` `aux_logits` tensor: +//! outputs from a `[B, K=2]` `aux_logits` tensor: //! //! out[0] = aux_conf_p50 — median over the batch of -//! `max_c softmax(logits[i, *])[c] - 1/3` +//! `max_c softmax(logits[i, *])[c] - 1/2` //! out[1] = aux_conf_std — population std over the batch of the same //! `aux_conf` row signal //! -//! `aux_conf` measures peak class confidence above the K=3 uniform -//! baseline (so 0 ⇔ uniform, 2/3 ⇔ fully concentrated). The downstream +//! `aux_conf` measures peak class confidence above the K=2 uniform +//! baseline (so 0 ⇔ uniform, 1/2 ⇔ fully concentrated). The downstream //! [`sp20_emas_compute_kernel`] (Phase 1.2) Wiener-blends these stats //! into ISV slots [510..520); this kernel itself does NOT write to ISV. //! +//! K=2 fixup (2026-05-09, post Phase 1.3): the kernel originally +//! assumed K=3 `{short, hold, long}`; the production aux head emits +//! K=2 `{down, up}` per `gpu_aux_heads.rs:61` +//! (`AUX_NEXT_BAR_K = 2`, established by SP13 B1.1a). The K=3 layout +//! was an error in the SP19+20 spec — this launcher and its kernel +//! were retargeted to K=2 atomically with the test rewrites and +//! audit-doc fixup so production aux logits are read correctly. +//! //! ── Hard rules covered ───────────────────────────────────────────── //! - `feedback_no_atomicadd` — block tree-reduce + per-warp tile //! histogram (the `sp4_histogram_p99` pattern adapted for p50). @@ -31,9 +39,11 @@ //! the rest of the SP20 reward chain, per //! `feedback_no_partial_refactor`. -/// Row width of `aux_logits` — fixed by the SP14-C aux head's 3-class -/// (down/flat/up) direction output. Kernel reads `[B, AUX_K_CLASSES]`. -pub const AUX_K_CLASSES: usize = 3; +/// Row width of `aux_logits` — fixed by the SP14-C aux head's 2-class +/// (down/up) direction output (`AUX_NEXT_BAR_K = 2` in +/// `gpu_aux_heads.rs:61`, established by SP13 B1.1a). Kernel reads +/// `[B, AUX_K_CLASSES]`. +pub const AUX_K_CLASSES: usize = 2; /// Block size used at every internal kernel pass (max-reduce, sum, /// sum-of-squares, histogram-of-warp-tiles). MUST stay 256: matches the @@ -135,13 +145,14 @@ mod tests { use super::*; #[test] - fn aux_k_classes_is_three() { + fn aux_k_classes_matches_production_aux_head() { // Locks the contract against silent changes — the SP14-C aux - // head emits 3 logits per row, and the kernel hardcodes that - // width via `AUX_K_CLASSES`. If the head ever moves to a - // different K, this constant + the kernel's `SP20_K_CLASSES` - // macro both must move together. - assert_eq!(AUX_K_CLASSES, 3); + // head emits `AUX_NEXT_BAR_K = 2` logits per row (down/up, + // SP13 B1.1a), and the kernel hardcodes that width via + // `AUX_K_CLASSES`. If the head ever moves to a different K, + // this constant + the kernel's `SP20_K_CLASSES` macro both + // must move together. + assert_eq!(AUX_K_CLASSES, 2); } #[test] diff --git a/crates/ml/src/cuda_pipeline/sp20_stats_compute_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_stats_compute_kernel.cu index 1d5e09f07..05e6bac0f 100644 --- a/crates/ml/src/cuda_pipeline/sp20_stats_compute_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_stats_compute_kernel.cu @@ -6,18 +6,25 @@ * and `sp20_controllers_compute` — land in Phase 1.2 and 1.3 and consume * this kernel's outputs as their inputs). * - * Reads a `[B, 3]` row-major `aux_logits` tensor (the SP14-C aux head's - * direction logits — `n_classes = 3` for {down, flat, up}) and produces - * two scalar f32 outputs: + * Reads a `[B, K=2]` row-major `aux_logits` tensor (the SP14-C aux head's + * direction logits — `n_classes = AUX_NEXT_BAR_K = 2` for {down, up}, per + * `gpu_aux_heads.rs:61`, established by SP13 B1.1a) and produces two + * scalar f32 outputs: * * out[0] = aux_conf_p50 — median over the batch of - * aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/3 + * aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/2 * out[1] = aux_conf_std — population std over the batch of aux_conf[i] * - * `aux_conf` is the per-row peak-confidence above the uniform baseline - * (uniform softmax at K=3 gives 1/3 per class, so subtracting 1/3 makes - * the signal zero when the head is maximally uncertain and approaches - * 2/3 when the head is fully concentrated on one class). + * `aux_conf` is the per-row peak-confidence above the K=2 uniform + * baseline (uniform softmax at K=2 gives 1/2 per class, so subtracting + * 1/2 makes the signal zero when the head is maximally uncertain and + * approaches 1/2 when the head is fully concentrated on one class). + * + * K=2 fixup (2026-05-09, post Phase 1.3): the kernel originally assumed + * K=3 `{short, hold, long}` with baseline `1/3`. The production aux head + * emits K=2 per `AUX_NEXT_BAR_K = 2` (SP13 B1.1a). Wiring K=2 production + * aux into a K=3 kernel = OOB reads + corrupt stats. This kernel was + * retargeted to K=2 atomically with the launcher, tests, and audit doc. * * Producer cadence: hot-path (per-training-step). Subsequent kernels * (sp20_emas_compute, sp20_controllers_compute) Wiener-blend these into @@ -38,20 +45,20 @@ * `MappedF32Buffer<2>` written via `__threadfence_system()` for * PCIe-visible coherence; host reads after stream sync. * - `pearl_symmetric_clamp_audit` — the `aux_conf` formula is itself - * bounded by softmax composition (max ∈ [1/3, 1] ⇒ aux_conf ∈ - * [0, 2/3]); no runtime clamp needed on the per-row signal. + * bounded by softmax composition (max ∈ [1/2, 1] ⇒ aux_conf ∈ + * [0, 1/2]); no runtime clamp needed on the per-row signal. * - `feedback_no_cpu_compute_strict` — every operation lives in the * kernel; no CPU fallback for the softmax / max / median / std math. * * Algorithm (single block, 256 threads): * - * Pass A: compute `aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/3` + * Pass A: compute `aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/2` * for i in [0, B), strided across threads, into a flat * per-row scratch `aux_conf_buf [B]` in dynamic shmem. - * Numerically-stable softmax: `m = max(l[i,0], l[i,1], l[i,2]); - * e_c = exp(l[i,c] - m); s = e_0 + e_1 + e_2; + * Numerically-stable softmax: `m = max(l[i,0], l[i,1]); + * e_c = exp(l[i,c] - m); s = e_0 + e_1; * sm_c = e_c / s; max_sm = max_c sm_c`. Then - * aux_conf[i] = max_sm - 1.0/3.0. + * aux_conf[i] = max_sm - 0.5f. * __syncthreads() after the pass. * * Pass B: block tree-reduce sum(aux_conf) and sum(aux_conf^2) for @@ -84,8 +91,8 @@ * With B=128, shmem ≈ 8192 + 512 = 8704 bytes (well under L40S 48 KB). * * Args: - * aux_logits — `[B, 3]` row-major aux head logits. f32 device ptr. - * Row stride is 3 floats; contiguous per row. + * aux_logits — `[B, K=2]` row-major aux head logits. f32 device ptr. + * Row stride is 2 floats; contiguous per row. * out — `MappedF32Buffer<2>` device ptr for [p50, std]. * B — batch size (rows of aux_logits). * ══════════════════════════════════════════════════════════════════════════ */ @@ -94,11 +101,11 @@ #include "sp4_histogram_p99.cuh" #define SP20_STATS_BLOCK 256 -#define SP20_K_CLASSES 3 -#define SP20_UNIFORM_K3 0.333333343f /* 1.0f / 3.0f rounded to f32 */ +#define SP20_K_CLASSES 2 +#define SP20_UNIFORM_K2 0.5f /* 1.0f / 2.0f exactly representable */ extern "C" __global__ void sp20_stats_compute_kernel( - const float* __restrict__ aux_logits, /* [B, 3] row-major */ + const float* __restrict__ aux_logits, /* [B, K=2] row-major */ float* __restrict__ out, /* MappedF32Buffer<2> */ int B) { @@ -141,25 +148,22 @@ extern "C" __global__ void sp20_stats_compute_kernel( float* s_aux_conf = reinterpret_cast( s_warp_tiles + warps * SP4_HIST_BINS); - /* ── Pass A: per-row softmax-max → aux_conf[i] ──────────────────── */ + /* ── Pass A: per-row K=2 softmax-max → aux_conf[i] ──────────────── */ for (int i = tid; i < B; i += bdim) { const int row_off = i * SP20_K_CLASSES; const float l0 = aux_logits[row_off + 0]; const float l1 = aux_logits[row_off + 1]; - const float l2 = aux_logits[row_off + 2]; /* Numerically-stable softmax: subtract row-max before exp. */ - const float m = fmaxf(l0, fmaxf(l1, l2)); + const float m = fmaxf(l0, l1); const float e0 = __expf(l0 - m); const float e1 = __expf(l1 - m); - const float e2 = __expf(l2 - m); - const float denom = e0 + e1 + e2; + const float denom = e0 + e1; const float sm0 = e0 / denom; const float sm1 = e1 / denom; - const float sm2 = e2 / denom; - const float max_sm = fmaxf(sm0, fmaxf(sm1, sm2)); + const float max_sm = fmaxf(sm0, sm1); - s_aux_conf[i] = max_sm - SP20_UNIFORM_K3; + s_aux_conf[i] = max_sm - SP20_UNIFORM_K2; } __syncthreads(); @@ -213,7 +217,7 @@ extern "C" __global__ void sp20_stats_compute_kernel( */ /* Pass C.1: max-reduce of |aux_conf| via the same tile alias. - * aux_conf is non-negative by construction (max_sm ≥ 1/3 ⇒ + * aux_conf is non-negative by construction (max_sm ≥ 1/2 ⇒ * aux_conf ≥ 0), so |x| = x; we still use fabsf to guarantee the * sign discipline that sp4_histogram_p99 demands. */ float local_max = 0.0f; diff --git a/crates/ml/tests/sp20_stats_compute_test.rs b/crates/ml/tests/sp20_stats_compute_test.rs index 42f9f06a8..4bb195bc0 100644 --- a/crates/ml/tests/sp20_stats_compute_test.rs +++ b/crates/ml/tests/sp20_stats_compute_test.rs @@ -2,17 +2,25 @@ //! //! Verifies the fused stats producer emits //! `[aux_conf_p50, aux_conf_std]` consistent with the closed-form CPU -//! oracle for two characteristic distributions of `aux_logits [B, 3]`: +//! oracle for two characteristic distributions of `aux_logits [B, K=2]`: //! -//! 1. **Uniform**: `aux_logits = 0` everywhere → softmax = [1/3,1/3,1/3] -//! → `aux_conf = max - 1/3 = 0` → p50 ≈ 0, std ≈ 0. +//! 1. **Uniform**: `aux_logits = 0` everywhere → softmax = [0.5, 0.5] +//! → `aux_conf = max - 0.5 = 0` → p50 ≈ 0, std ≈ 0. //! 2. **Concentrated**: `aux_logits[*, 0] = 10`, rest = 0 → softmax -//! ≈ [1, 0, 0] (one-hot) → `aux_conf ≈ 2/3` → p50 ≈ 0.667, std ≈ 0. +//! ≈ [1, 0] (one-hot) → `aux_conf ≈ 0.5` → p50 ≈ 0.5, std ≈ 0. //! //! Tests also cover a heterogeneous distribution with non-trivial p50 //! and std so the kernel is exercised on both the histogram path and //! the running-sum/sum-of-squares paths together. //! +//! K=2 fixup (2026-05-09, post Phase 1.3): the kernel originally assumed +//! K=3 `{down, flat, up}` but the production aux head emits K=2 +//! `{down, up}` per `gpu_aux_heads.rs:61` (`AUX_NEXT_BAR_K = 2`, +//! established by SP13 B1.1a). This test file was rewritten to exercise +//! the K=2 contract; the CPU oracle, expected p50/std ranges, and the +//! degenerate-uniform branch all reflect the K=2 baseline `1/2 = 0.5` +//! and the resulting `aux_conf ∈ [0, 0.5]` range. +//! //! Per `feedback_no_atomicadd` the kernel uses block tree-reduce + per- //! warp tile binning. Per `feedback_no_htod_htoh_only_mapped_pinned` //! all CPU↔GPU buffers are `MappedF32Buffer`. @@ -54,26 +62,24 @@ mod gpu { .expect("load sp20_stats_compute_kernel function") } - /// CPU oracle: per-row softmax → max → minus 1/3, then median + std. + /// CPU oracle: per-row K=2 softmax → max → minus 1/2, then median + std. /// Mirrors the kernel's math exactly so the tolerances below pin /// fp32 rounding only, not algorithmic deltas. fn cpu_oracle(aux_logits: &[f32], b: usize) -> (f32, f32) { assert_eq!(aux_logits.len(), b * AUX_K_CLASSES); - let one_third = 1.0_f32 / 3.0_f32; + let one_half = 0.5_f32; let mut conf: Vec = (0..b) .map(|i| { let off = i * AUX_K_CLASSES; let l0 = aux_logits[off]; let l1 = aux_logits[off + 1]; - let l2 = aux_logits[off + 2]; - let m = l0.max(l1).max(l2); + let m = l0.max(l1); let e0 = (l0 - m).exp(); let e1 = (l1 - m).exp(); - let e2 = (l2 - m).exp(); - let s = e0 + e1 + e2; - let max_sm = (e0 / s).max(e1 / s).max(e2 / s); - max_sm - one_third + let s = e0 + e1; + let max_sm = (e0 / s).max(e1 / s); + max_sm - one_half }) .collect(); @@ -81,7 +87,7 @@ mod gpu { // semantics of the kernel (the sorted-array median is what the // histogram converges to as bin_width → 0; for our two // characteristic tests below the values are concentrated at - // 0 or 2/3 which the histogram resolves exactly). + // 0 or 0.5 which the histogram resolves exactly). conf.sort_by(|a, b| a.partial_cmp(b).unwrap_or(std::cmp::Ordering::Equal)); let median_idx = (b + 1) / 2 - 1; // 0-based ⌈B/2⌉ - 1 let median = conf[median_idx]; @@ -125,8 +131,8 @@ mod gpu { /// Test 1 — uniform-logit distribution. /// - /// `aux_logits = 0` everywhere ⇒ softmax = [1/3, 1/3, 1/3] for every - /// row ⇒ `aux_conf = 1/3 - 1/3 = 0` for every row. + /// `aux_logits = 0` everywhere ⇒ softmax = [0.5, 0.5] for every + /// row ⇒ `aux_conf = 0.5 - 0.5 = 0` for every row. /// Expected: p50 ≈ 0, std ≈ 0. The kernel's degenerate-distribution /// branch (step_max == 0) writes both outputs to 0 directly. #[test] @@ -152,17 +158,17 @@ mod gpu { /// Test 2 — concentrated-logit distribution (varied confidence). /// /// Per-row class-0 logit varies linearly from 0 to 3 across the - /// batch, with classes 1 and 2 held at 0. softmax(class 0) varies - /// from 1/3 (logit 0, uniform) to ≈ 0.85 (logit 3), so `aux_conf` - /// varies from 0 to ≈ 0.52 — a 0.5 range spread across well over - /// half of the 256 histogram bins (per + /// batch, with class 1 held at 0. softmax(class 0) varies from 0.5 + /// (logit 0, uniform) to ≈ 0.953 (logit 3), so `aux_conf` varies + /// from 0 to ≈ 0.453 — a 0.45 range spread across well over half of + /// the 256 histogram bins (per /// `pearl_sp4_histogram_warp_tile_undercount`: lockstep-uniform /// inputs cause warp-tile collisions in the non-atomic per-warp /// histogram tiles, so test data MUST vary per row across many /// bins, not just within one bin's saturating-softmax tail). /// /// Logits are kept under the saturating-softmax regime (max < 4) - /// because the softmax tail compresses 0.85 → 1.0 confidences + /// because the softmax tail compresses 0.95 → 1.0 confidences /// into a range smaller than one bin_width (≈ 0.4% of step_max), /// re-creating the warp-tile collision pattern this test /// explicitly avoids. @@ -178,7 +184,7 @@ mod gpu { let mut aux_logits = vec![0.0_f32; B * AUX_K_CLASSES]; for i in 0..B { // Linear ramp 0 → 3 on class 0 — confidence ramps from - // uniform (0) to a non-saturating peak (≈ 0.85). + // uniform (0.5) to a non-saturating peak (≈ 0.953). let frac = (i as f32) / ((B - 1) as f32); aux_logits[i * AUX_K_CLASSES] = 3.0 * frac; } @@ -186,8 +192,8 @@ mod gpu { let (p50_gpu, std_gpu) = run_kernel(B, &aux_logits); let (p50_cpu, std_cpu) = cpu_oracle(&aux_logits, B); - // The GPU histogram bin_width ≈ step_max / 256 ≈ 0.52 / 256 ≈ - // 0.002; agreement to within a few bin_widths is the design + // The GPU histogram bin_width ≈ step_max / 256 ≈ 0.45 / 256 ≈ + // 0.0018; agreement to within a few bin_widths is the design // tolerance for the histogram path. let p50_tol = 0.02_f32; assert!( @@ -196,12 +202,12 @@ mod gpu { (p50_gpu - p50_cpu).abs(), ); // CPU oracle sanity: with logits ramping 0 → 3, the median - // aux_conf is the row at logit ≈ 1.5 → softmax ≈ 0.578 → - // aux_conf ≈ 0.245. We pin it to the (0.15, 0.40) range to - // catch contract drift while tolerating fp32 / softmax math. + // aux_conf is the row at logit ≈ 1.5 → softmax ≈ 0.818 → + // aux_conf ≈ 0.318. We pin it to (0.10, 0.40) to catch + // contract drift while tolerating fp32 / softmax math. assert!( - p50_cpu > 0.15_f32 && p50_cpu < 0.40_f32, - "CPU oracle: concentrated p50 should land in (0.15, 0.40), got {p50_cpu}", + p50_cpu > 0.10_f32 && p50_cpu < 0.40_f32, + "CPU oracle: concentrated p50 should land in (0.10, 0.40), got {p50_cpu}", ); // std non-trivial; CPU and GPU within fp32 noise. let std_tol = 5e-3_f32; @@ -210,80 +216,94 @@ mod gpu { "concentrated: std_gpu={std_gpu} vs cpu_oracle={std_cpu} (diff={:.3e}, tol={std_tol})", (std_gpu - std_cpu).abs(), ); - // CPU oracle std: the aux_conf vector spans ≈ 0.52 with a - // monotonic ramp, so std ≈ 0.52 / sqrt(12) ≈ 0.15 for a - // uniform-distribution-style spread. + // CPU oracle std: the aux_conf vector spans ≈ 0.45 with a + // monotonic-but-saturating ramp, so std lands roughly in the + // 0.05-0.20 range for a sigmoid-shaped spread. assert!( - std_cpu > 0.05_f32 && std_cpu < 0.25_f32, - "CPU oracle: concentrated std should land in (0.05, 0.25), got {std_cpu}", + std_cpu > 0.05_f32 && std_cpu < 0.20_f32, + "CPU oracle: concentrated std should land in (0.05, 0.20), got {std_cpu}", ); } /// Test 3 — heterogeneous distribution exercising both p50 and std. /// - /// Half the rows get jittered hot logits (aux_conf ≈ 0.667) and - /// half get jittered uniform logits (aux_conf ≈ 0). The per-row - /// `aux_conf` is therefore a 50/50 mix of values around 0 and - /// values around 0.667. The median lands at the boundary (since - /// `(B+1)/2 - 1 = 64` indexes the upper end of the lower cluster), - /// and the std is `≈ 1/3` (half-distance between the two - /// concentration clusters). + /// Two clusters with ramped (not lockstep) aux_conf within each + /// cluster: + /// - Lower half (rows 0..64): logit-0 ramps 0.0 → 1.0, logit-1 + /// held at 0 ⇒ aux_conf ramps 0 → sigmoid(1.0)-0.5 ≈ 0.231 + /// - Upper half (rows 64..128): logit-0 ramps 2.5 → 4.0, logit-1 + /// held at 0 ⇒ aux_conf ramps sigmoid(2.5)-0.5 ≈ 0.424 to + /// sigmoid(4.0)-0.5 ≈ 0.491 /// - /// Per-row jitter is required to avoid the + /// The two ramps are well-separated so the median of the sorted + /// aux_conf vector lands at the top of the lower cluster + /// (≈ 0.231). Std is dominated by the inter-cluster gap and lands + /// near 0.15. + /// + /// Per-row ramping (rather than per-row jitter on a constant + /// cluster center) ensures each warp's 32 lanes compute distinct + /// bin indices, avoiding the /// `pearl_sp4_histogram_warp_tile_undercount` lockstep-uniform - /// trap (32 lanes in a warp racing on the same bin without - /// atomicAdd undercounts the histogram). Production aux logits - /// are naturally jittered by the per-sample net forward, so - /// jittered test data IS the realistic-input model. + /// trap where 32 lanes racing on the same non-atomic tile slot + /// undercount the histogram. Production aux logits are naturally + /// distributed across bins, so spread test data IS the realistic- + /// input model. #[test] #[ignore = "requires GPU"] fn heterogeneous_logits_match_cpu_oracle() { const B: usize = 128; let mut aux_logits = vec![0.0_f32; B * AUX_K_CLASSES]; + // Lower cluster: l0 ramps 0 → 1.0; aux_conf 0 → 0.231. for i in 0..(B / 2) { - // Hot half: jitter ±0.5 around 10.0 on class 0. - let jitter = 0.5 * ((i as f32) / ((B / 2) as f32) - 0.5); - aux_logits[i * AUX_K_CLASSES] = 10.0 + jitter; + let frac = (i as f32) / ((B / 2 - 1) as f32); + aux_logits[i * AUX_K_CLASSES + 0] = 1.0 * frac; + // l1 = 0 } - // Uniform half: jitter all three logits ±0.05 around 0 so - // each row's max varies slightly across bins. + // Upper cluster: l0 ramps 2.5 → 4.0; aux_conf 0.424 → 0.491. for i in (B / 2)..B { - let k = (i - B / 2) as f32; - let scale = 0.05; - aux_logits[i * AUX_K_CLASSES + 0] = scale * (k * 0.10).sin(); - aux_logits[i * AUX_K_CLASSES + 1] = scale * (k * 0.13).cos(); - aux_logits[i * AUX_K_CLASSES + 2] = scale * (k * 0.07).sin(); + let k = i - B / 2; + let frac = (k as f32) / ((B / 2 - 1) as f32); + aux_logits[i * AUX_K_CLASSES + 0] = 2.5 + 1.5 * frac; + // l1 = 0 } let (p50_gpu, std_gpu) = run_kernel(B, &aux_logits); let (p50_cpu, std_cpu) = cpu_oracle(&aux_logits, B); - // p50 sanity: the sorted aux_conf vector has 64 small-near-zero - // values then 64 values ≈ 0.667; ⌈B/2⌉-th element (1-based) - // = 64th = the last "near-zero", so the median is small in - // both the CPU oracle and the GPU histogram. With uniform-half - // jitter ≈ 0.05, the per-row aux_conf there is at most - // (max softmax 0.34) - 1/3 ≈ 0.01, so p50 < 0.02. Tolerate - // up to a couple of bin_widths above that. + // p50 sanity: the sorted aux_conf vector has 64 lower-cluster + // values (0 → 0.231) then 64 upper-cluster values (0.424 → + // 0.491). ⌈B/2⌉-th element (1-based) = 64th = the largest + // lower-cluster value ≈ 0.231. Both CPU and GPU should agree + // within a few bin_widths (bin_width = step_max/256 ≈ + // 0.491/256 ≈ 0.002). let p50_tol = 0.02_f32; assert!( (p50_gpu - p50_cpu).abs() < p50_tol, "het: p50_gpu={p50_gpu} vs cpu_oracle={p50_cpu} (diff={:.3e}, tol={p50_tol})", (p50_gpu - p50_cpu).abs(), ); - // std sanity: with the hot cluster at ≈ 0.667 and the uniform - // cluster near 0, var = 0.5 × (0.333)² + 0.5 × (0.333)² ≈ 1/9 - // ⇒ std ≈ 1/3. Jitter perturbs this slightly; allow ±5% of - // 1/3 between the GPU result and the CPU oracle. + // CPU oracle sanity: median should be at the top of the lower + // cluster, ≈ 0.231. Pin to (0.18, 0.27) to catch contract + // drift while tolerating fp32 / sigmoid math. + assert!( + p50_cpu > 0.18_f32 && p50_cpu < 0.27_f32, + "CPU oracle: het p50 should land in (0.18, 0.27), got {p50_cpu}", + ); + // std sanity: with two well-separated ramped clusters (lower + // 0..0.231, upper 0.424..0.491), the variance is dominated + // by the inter-cluster gap. Jitter perturbs this slightly; + // allow ±5e-3 between the GPU result and the CPU oracle. let std_tol = 5e-3_f32; assert!( (std_gpu - std_cpu).abs() < std_tol, "het: std_gpu={std_gpu} vs cpu_oracle={std_cpu} (diff={:.3e}, tol={std_tol})", (std_gpu - std_cpu).abs(), ); + // CPU oracle std lands roughly in the 0.12–0.20 band given + // the bimodal ramped layout. assert!( - (std_cpu - (1.0_f32 / 3.0_f32)).abs() < 0.02, - "CPU oracle: het std should be ≈ 1/3, got {std_cpu}", + std_cpu > 0.12_f32 && std_cpu < 0.20_f32, + "CPU oracle: het std should land in (0.12, 0.20), got {std_cpu}", ); } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 099de1f4e..821d220a3 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -10762,6 +10762,38 @@ bash scripts/audit_sp18_consumers.sh --check # exit - New GPU oracle test `crates/ml/tests/sp20_stats_compute_test.rs` - Cubin manifest entry in `crates/ml/build.rs` +### K=2 fixup (2026-05-09, post Phase 1.3) + +The original Phase 1.1 implementation assumed the SP14-C aux head +emits 3-class logits `{short, hold, long}` with baseline `1/3`. +This was an error in the SP19+20 spec. The production aux head +emits `K = AUX_NEXT_BAR_K = 2` logits `{down, up}` (per +`crates/ml/src/cuda_pipeline/gpu_aux_heads.rs:61`, established by +SP13 B1.1a). Wiring K=2 production aux into a K=3 kernel = OOB +reads and corrupt stats. This fixup retargets the kernel + +launcher + tests to K=2 atomically: + +- `SP20_K_CLASSES`: `3` → `2`; `SP20_UNIFORM_K3 = 0.333...` → + `SP20_UNIFORM_K2 = 0.5f`. +- Pass A reads 2 logits per row (was 3); softmax sums two + exponentials. +- `aux_conf[i] = max_c softmax(...) - 1/2`; range `[0, 1/2]` + (was `[0, 2/3]`). +- `AUX_K_CLASSES` constant in the launcher: `3` → `2`; launcher + unit test renamed `aux_k_classes_is_three` → + `aux_k_classes_matches_production_aux_head`. +- GPU oracle tests rewritten with K=2 CPU oracle, expected p50/std + ranges, and a redesigned heterogeneous-distribution test using + ramped (not lockstep) clusters to avoid the + `pearl_sp4_histogram_warp_tile_undercount` trap when the entire + hot half collapses to bin 255 in K=2's saturated softmax regime. + +Phase 1.2 (`sp20_emas_compute`) and Phase 1.3 +(`sp20_controllers_compute`) consume the scalar `[p50, std]` outputs +and do NOT carry the K dimension; both regression test suites +(`sp20_emas_compute_test`, `sp20_controllers_compute_test`) continue +to pass unmodified. + ### Purpose Component 5 / Kernel 3 of the SP20 fused-producer chain (Phase 1.4 @@ -10769,16 +10801,17 @@ will land Kernels 1+2 — `sp20_emas_compute` + `sp20_controllers_ compute` — and the production launch site atomically with the training-loop wire-up per `feedback_no_partial_refactor`). -Reads `aux_logits [B, 3]` (the SP14-C aux head's 3-class direction -logits) and emits `[aux_conf_p50, aux_conf_std]` into a -`MappedF32Buffer<2>`, where the per-row signal is +Reads `aux_logits [B, K=2]` (the SP14-C aux head's 2-class direction +logits, per `AUX_NEXT_BAR_K = 2`) and emits +`[aux_conf_p50, aux_conf_std]` into a `MappedF32Buffer<2>`, where +the per-row signal is ``` -aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/3 +aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/2 ``` -`aux_conf` measures peak class confidence above the K=3 uniform -baseline (0 ⇔ uniform, 2/3 ⇔ fully concentrated). The downstream +`aux_conf` measures peak class confidence above the K=2 uniform +baseline (0 ⇔ uniform, 1/2 ⇔ fully concentrated). The downstream EMA producer (Phase 1.2) Wiener-blends p50 + std into ISV slots within [510..520) reserved by SP20 (`f5eed1fa7`). @@ -10787,7 +10820,7 @@ within [510..520) reserved by SP20 (`f5eed1fa7`). Single-block, 256-thread kernel; one fused stream of `aux_logits` for both stats per `pearl_fused_per_group_statistics_oracle`. - - **Pass A**: per-row numerically-stable softmax → max → minus 1/3, + - **Pass A**: per-row numerically-stable softmax → max → minus 1/2, write `aux_conf[i]` into per-row scratch in dynamic shmem. - **Pass B1**: block tree-reduce `sum(aux_conf)`. Reuses `s_bins[0..256]` (the histogram tile) as a float-via-int reduce @@ -10855,19 +10888,25 @@ the wire-up will consume. (all `#[ignore = "requires GPU"]`): 1. `uniform_logits_emit_zero_stats` — `aux_logits = 0` ⇒ - softmax = [1/3, 1/3, 1/3] ⇒ aux_conf = 0 ⇒ p50 = std = 0. + softmax = [1/2, 1/2] ⇒ aux_conf = 0 ⇒ p50 = std = 0. Exercises the kernel's degenerate-distribution branch. 2. `varied_confidence_logits_match_cpu_oracle` — class-0 logit - ramps 0 → 3, others 0; CPU oracle matches GPU within - bin_width tolerance for p50 (≈ 0.002) and 5e-3 for std. - 3. `heterogeneous_logits_match_cpu_oracle` — half jittered - hot, half jittered uniform; CPU oracle matches GPU within - 0.02 for p50 and 5e-3 for std. + ramps 0 → 3, class-1 held at 0; CPU oracle matches GPU + within bin_width tolerance for p50 (≈ 0.002) and 5e-3 for + std. + 3. `heterogeneous_logits_match_cpu_oracle` — two ramped clusters + (lower 0 → 1.0, upper 2.5 → 4.0 on class-0 logit; class-1 = 0) + so per-warp lanes compute distinct bins (avoiding the + `pearl_sp4_histogram_warp_tile_undercount` trap that K=2's + saturated softmax would otherwise trigger when the entire + hot half collapses to bin 255). CPU oracle matches GPU + within 0.02 for p50 and 5e-3 for std. 4. `empty_batch_writes_zero_stats` — `B = 0` exercises the pre-warmup degenerate guard; kernel writes `[0, 0]` and returns without entering the histogram path. -Plus 4 unit tests in the launcher module (`AUX_K_CLASSES = 3`, +Plus 4 unit tests in the launcher module +(`AUX_K_CLASSES = AUX_NEXT_BAR_K = 2` per the K=2 fixup, block-size discipline, shmem budget under L40S 48 KiB, zero-batch shmem accounting). diff --git a/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md b/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md index 176bffcef..434248bfd 100644 --- a/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md +++ b/docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md @@ -1,5 +1,16 @@ # SP19+20 WR-First Reward Implementation Plan +> **AMENDED 2026-05-09**: K=3 → K=2 retarget for `sp20_stats_compute`. The +> production aux head emits `K = AUX_NEXT_BAR_K = 2` logits `{down, up}` per +> `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs:61` (established by SP13 B1.1a), +> not K=3 `{short, hold, long}`. The Phase 1.1 task pseudocode below shows +> `aux_logits[B, 3]` and `aux_conf = max(...) - 1/3` (range `[0, 2/3]`); these +> should be read as `aux_logits[B, 2]` and `aux_conf = max(...) - 1/2` (range +> `[0, 1/2]`). The kernel + launcher + tests already landed K=2 atomically +> with this amendment — see the SP20 Phase 1.1 K=2 fixup commit (this +> commit's SHA in the audit log) and `docs/dqn-wire-up-audit.md` +> "K=2 fixup" subsection. + > **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. **Goal:** Implement event-driven WR-first reward with multi-horizon directional ground truth, adaptive asymmetric loss aversion, and aux-gated Q-targets to lift WR ≥ 55% and PF ≥ 2.0. diff --git a/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md b/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md index 7c16033c2..e01e356f8 100644 --- a/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md +++ b/docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md @@ -1,5 +1,15 @@ # SP19+20: WR-First Reward + Multi-Horizon Label Utilization +> **AMENDED 2026-05-09**: K=3 → K=2 retarget for `sp20_stats_compute`. The +> production aux head emits `K = AUX_NEXT_BAR_K = 2` logits `{down, up}` per +> `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs:61` (established by SP13 B1.1a), +> not the K=3 `{short, hold, long}` assumed in this spec. All +> `aux_conf = max(softmax(...)) - 1/3` formulas with range `[0, 2/3]` should +> be read as `aux_conf = max(softmax(...)) - 1/2` with range `[0, 1/2]`. See +> the SP20 Phase 1.1 K=2 fixup commit (this commit's SHA in the audit log) +> and `docs/dqn-wire-up-audit.md` "K=2 fixup" subsection for the kernel / +> launcher / test updates that landed atomically with this amendment. + **Spec date:** 2026-05-09 **Branch:** `sp19-20-wr-first` forks from `feat/sp18-combined` HEAD `d39005c6f` **Combines:** SP19 (multi-horizon labels, already landed) + SP20 (WR-first reward design) @@ -154,7 +164,7 @@ Two parallel emissions: ``` Per bar i: - aux_conf_i = max(softmax(aux_logits_i)) - 1/3 # range [0, 2/3] + aux_conf_i = max(softmax(aux_logits_i)) - 1/2 # range [0, 1/2] (K=2; see AMENDED note) cost_scale = ISV[HOLD_COST_SCALE] # adaptive [0.01, 0.5] per_bar_opp_cost_i = -aux_conf_i × cost_scale @@ -244,7 +254,7 @@ The eligibility trace state must reset at every trade close, not carry across. T ``` At each replay batch step: - aux_conf = max(softmax(aux_logits_at_state)) - 1/3 + aux_conf = max(softmax(aux_logits_at_state)) - 1/2 # K=2; see AMENDED note threshold = ISV[AUX_CONF_THRESHOLD] # adaptive temp = ISV[AUX_GATE_TEMP] # adaptive