gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fixup(sp20): Phase 1.1 K=3 → K=2 retarget (production aux head is K=2)

Browse Source 
The Phase 1.1 sp20_stats_compute kernel (de922c6a4) 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 — production aux head emits K=2 logits
{down, up} per gpu_aux_heads.rs:61 (AUX_NEXT_BAR_K = 2, established by
SP13 B1.1a). Wiring K=2 production aux into the K=3 kernel = OOB reads +
corrupt stats.

Retargets the kernel + launcher + tests to K=2 atomically per
feedback_no_partial_refactor:

- Kernel: SP20_K_CLASSES 3→2; SP20_UNIFORM_K3 0.333…→SP20_UNIFORM_K2 0.5f;
  Pass A reads 2 logits (was 3); aux_conf range [0, 1/2] (was [0, 2/3]).
- Launcher: AUX_K_CLASSES 3→2; renamed unit test
  aux_k_classes_is_three → aux_k_classes_matches_production_aux_head.
- Tests: rewrote CPU oracle for K=2 input shape and 0.5 baseline; updated
  expected p50/std ranges; redesigned heterogeneous-distribution test to
  use ramped (not lockstep) clusters — K=2's saturated softmax in the hot
  half collapses every row to bin 255, triggering the
  pearl_sp4_histogram_warp_tile_undercount trap; ramped clusters distribute
  bin indices across each warp's 32 lanes so the histogram path matches
  the CPU oracle within bin_width tolerance.

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 suites verified passing unmodified:
- sp20_emas_compute_test: 4 GPU + 1 unit, all pass.
- sp20_controllers_compute_test: 7 GPU, all pass.

Verification (RTX 3050 Ti, sm_86):
- sp20_stats_compute_test: 4 GPU oracle + 4 launcher unit, all pass.
- cargo check -p ml --features cuda: clean (pre-existing warnings only).

Spec + plan amended at top with "AMENDED 2026-05-09" notes; audit doc
Phase 1.1 entry has a "K=2 fixup" subsection documenting the change.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 19:46:17 +02:00
parent 5ded1cb4b9
commit 4e21c38b85
6 changed files with 220 additions and 125 deletions
Download Patch File Download Diff File Expand all files Collapse all files

37
crates/ml/src/cuda_pipeline/sp20_stats_compute.rs
View File

@@ -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]

60
crates/ml/src/cuda_pipeline/sp20_stats_compute_kernel.cu
View File

@@ -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<float*>(
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;

156
crates/ml/tests/sp20_stats_compute_test.rs
View File

@@ -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<f32> = (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}",
);
}

67
docs/dqn-wire-up-audit.md
View File

@@ -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).

11
docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md
View File

@@ -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.

14
docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md
View File

@@ -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