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

feat(sp15-p3.1): r_quality + r_discipline split with ISV-driven α + sentinel cold-start

Browse Source 
Per spec §8.2 (3.1) post-amendment-2 fix: ALPHA_SPLIT slot initialized
DIRECTLY to 0.5 in trainer constructor. Formula α = grad_norm_q /
(grad_norm_q + grad_norm_d + ε) takes over only after BOTH grad-norm
EMAs accumulate ≥ N_WARM=100 non-zero observations.

Two kernels in r_quality_discipline_split_kernel.cu (single cubin per
established 1:1-source-to-cubin pattern with multiple kernels):
  - r_quality_discipline_split_kernel: per-step composition + warm count
  - alpha_split_producer_kernel: per-step ALPHA_SPLIT update from grad ratio
    (gated on warm count to prevent premature formula activation)

3 ISV slots (417 ALPHA_SPLIT, 418 GRAD_NORM_QUALITY, 419 GRAD_NORM_DISCIPLINE)
+ sp15_alpha_warm_count [1] mapped-pinned scratch buffer on the trainer
struct. 4 fold-reset registry entries + dispatch arms (one for the
non-ISV warm-count buffer mirrors the sp11_novelty_hash host_slice_mut
pattern).

Per established Phase precedent: kernels + launchers land first; consumer
migration (per-step launches in training_loop.rs reward composition site)
deferred to a follow-up commit per feedback_no_partial_refactor.

Anchor tests: 2.4 cost_sensitivity + 2.6 regime_silences (Phase 2B
contracts) — green via Phase 3.4 regret + 3.2 cost; this commit lands
the split structure they depend on.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-06 15:26:56 +02:00
parent ef373c34d7
commit 2d226e6e76
7 changed files with 493 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
crates/ml/build.rs
View File

@@ -886,6 +886,18 @@ fn main() {
// buffer + bumping s1_input_dim) lands as a follow-up atomic
// commit per the established Phase 1.1-1.4 precedent.
"dd_pct_concat_kernel.cu",
// SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split
// kernel + ISV-driven α producer kernel (single source file, two
// `extern "C" __global__` symbols sharing one cubin per the
// 1:1-source-to-cubin / multi-kernel-per-file convention used by
// `baseline_kernels.cu` and `novelty_simhash_kernel.cu`).
// Composer (`r_quality_discipline_split_kernel`) reads
// ALPHA_SPLIT_INDEX=417 + alpha_warm_count scratch buffer, writes
// r_total + increments warm count. Producer
// (`alpha_split_producer_kernel`) reads grad-norm slots 418/419
// + warm count, writes ALPHA_SPLIT_INDEX clamped to [0.05, 0.95]
// (silent no-op until `*alpha_warm_count >= N_WARM=100`).
"r_quality_discipline_split_kernel.cu",
];
// ALL kernels get common header (BF16 types + wrappers)

204
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -1187,6 +1187,139 @@ pub fn launch_sp15_baseline_naive_reversion(
Ok(())
}
/// SP15 Phase 3.1 (2026-05-06): cubin shared by the r_quality +
/// r_discipline split composer (`r_quality_discipline_split_kernel`) and
/// the ISV-driven α producer (`alpha_split_producer_kernel`). Per spec
/// §8.2 (3.1) post-amendment-2 fix. Single source file with two
/// `extern "C" __global__` symbols sharing one cubin per the established
/// 1:1-source-to-cubin / multi-kernel-per-file pattern (mirrors the SP15
/// Phase 1.4 baseline cubin and the SP11 SimHash novelty cubin). The two
/// launchers below load this single cubin and resolve different
/// `get_function()` symbols.
///
/// Cold-start protocol:
/// • Constructor writes ALPHA_SPLIT_INDEX (slot 417) DIRECTLY to 0.5
/// (NOT derived from the formula, which would produce 0/0=0 from the
/// zero grad-norm EMAs at boot).
/// • The composer maintains an `alpha_warm_count` scratch buffer
/// ([1] f32 mapped-pinned). Each composer launch increments the
/// count; while `*count < N_WARM=100` the composer overrides α back
/// to 0.5 regardless of producer writes.
/// • The producer is gated on the same warm count and stays a silent
/// no-op until `*count >= N_WARM`; once it activates it writes
/// α = clamp(grad_norm_q / (grad_norm_q + grad_norm_d + ε), 0.05,
/// 0.95) on each step.
///
/// Phase 3.1 lands kernels + launchers + ISV anchor seed + scratch
/// buffer + state-reset registry entries + dispatch arms only;
/// per-step consumer migration in `training_loop.rs` (composing
/// r_total at the reward composition site + the producer launch in the
/// gradient pipeline) is deferred to a follow-up commit per
/// `feedback_no_partial_refactor.md` — kernels + launchers verify in
/// isolation first via the GPU oracle test
/// `r_split_uses_sentinel_alpha_at_cold_start`, mirroring the
/// established Phase 1.1-1.5 atomic pattern.
pub static SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/r_quality_discipline_split_kernel.cubin"));
/// SP15 Phase 3.1 (2026-05-06): launcher for the r_quality + r_discipline
/// split composer. Free function (matches `launch_sp15_dd_state` /
/// `launch_sp15_baseline_*` precedent) so unit/oracle tests can drive the
/// kernel directly without the trainer struct; production callers in
/// `training_loop.rs` invoke the same launcher per step inside the
/// reward-composition site.
///
/// Contract:
/// `r_quality`, `r_discipline`: per-step f32 reward components.
/// `isv`: device ptr to ISV bus (≥443 f32 slots).
/// `r_total_out`: device ptr to f32 [1], caller-allocated.
/// `alpha_warm_count`: device ptr to f32 [1] mapped-pinned
/// scratch — the trainer struct holds the
/// production buffer; tests pass a fresh
/// MappedF32Buffer.
///
/// Single thread, single block — the kernel is a per-step state machine,
/// not a reduction.
pub fn launch_sp15_r_quality_discipline_split(
stream: &Arc<CudaStream>,
r_quality: f32,
r_discipline: f32,
isv: cudarc::driver::sys::CUdeviceptr,
r_total_out: cudarc::driver::sys::CUdeviceptr,
alpha_warm_count: cudarc::driver::sys::CUdeviceptr,
) -> Result<(), MLError> {
let module = stream
.context()
.load_cubin(SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"load sp15_r_quality_discipline_split cubin: {e}"
)))?;
let kernel = module
.load_function("r_quality_discipline_split_kernel")
.map_err(|e| MLError::ModelError(format!(
"load r_quality_discipline_split_kernel function: {e}"
)))?;
unsafe {
stream
.launch_builder(&kernel)
.arg(&r_quality)
.arg(&r_discipline)
.arg(&isv)
.arg(&r_total_out)
.arg(&alpha_warm_count)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"launch r_quality_discipline_split_kernel: {e}"
)))?;
}
Ok(())
}
/// SP15 Phase 3.1 (2026-05-06): launcher for the ISV-driven α producer.
/// Loads the SAME cubin as `launch_sp15_r_quality_discipline_split`
/// (multi-kernel-per-file pattern) and resolves the
/// `alpha_split_producer_kernel` symbol. Reads grad-norm slots
/// 418/419 + alpha_warm_count, writes ALPHA_SPLIT_INDEX (slot 417). No-op
/// until `*alpha_warm_count >= N_WARM=100`.
///
/// Single thread, single block.
pub fn launch_sp15_alpha_split_producer(
stream: &Arc<CudaStream>,
isv: cudarc::driver::sys::CUdeviceptr,
alpha_warm_count: cudarc::driver::sys::CUdeviceptr,
) -> Result<(), MLError> {
let module = stream
.context()
.load_cubin(SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"load sp15_r_quality_discipline_split cubin: {e}"
)))?;
let kernel = module
.load_function("alpha_split_producer_kernel")
.map_err(|e| MLError::ModelError(format!(
"load alpha_split_producer_kernel function: {e}"
)))?;
unsafe {
stream
.launch_builder(&kernel)
.arg(&isv)
.arg(&alpha_warm_count)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"launch alpha_split_producer_kernel: {e}"
)))?;
}
Ok(())
}
/// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty signal — lookup +
/// update kernels sharing one cubin. Lookup reads
/// `1/sqrt(1+count)` ∈ [0, 1] for each (state, action) bucket; update
@@ -3975,6 +4108,24 @@ pub struct GpuDqnTrainer {
sel_t_pinned: *mut i32,
sel_t_dev_ptr: u64,
/// SP15 Phase 3.1 (2026-05-06): warm-up counter for the
/// r_quality + r_discipline split composer (per spec §8.2 (3.1)
/// post-amendment-2 fix). Single-element mapped-pinned buffer
/// initialised to 0.0 in the constructor and reset to 0.0 at fold
/// boundary by the `sp15_alpha_warm_count` registry-arm dispatch
/// (which calls `host_slice_mut().fill(0.0)` mirroring the
/// `sp11_novelty_hash` pattern). Each per-step launch of
/// `r_quality_discipline_split_kernel` increments the count by 1;
/// once `*count >= N_WARM=100` the composer stops overriding α to
/// 0.5 and `alpha_split_producer_kernel` activates from no-op state
/// to write the formula's clamped output to ALPHA_SPLIT_INDEX. The
/// counter is a producer-controlled scalar — the host writes 0.0
/// at construction / fold reset (allowed under
/// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write
/// rule, NOT a host-side compute) and reads it through the mapped
/// page via `read_all()` only in tests.
pub(crate) sp15_alpha_warm_count: super::mapped_pinned::MappedF32Buffer, // [1]
// ── Per-branch Q-gap EMA (trajectory backtracking state) ──
/// [4] per-branch Q-gap EMA — device buffer (read-modify-write for backtracking restore).
per_branch_q_gap_ema_buf: CudaSlice<f32>,
@@ -18612,6 +18763,31 @@ impl GpuDqnTrainer {
dev_ptr
};
// SP15 Phase 3.1 (2026-05-06): warm-up counter for the
// r_quality + r_discipline split composer (per spec §8.2 (3.1)
// post-amendment-2 fix). Single-element mapped-pinned buffer
// initialised to 0.0 by `MappedF32Buffer::new`'s zero-init
// contract — the host writes 0.0 at construction and again at
// fold boundary via the registry-arm dispatch (allowed under
// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write
// rule, NOT a host-side compute). The composer kernel
// (`r_quality_discipline_split_kernel`) increments by 1 each
// step; once `*count >= N_WARM=100` the composer stops
// overriding α and the producer kernel
// (`alpha_split_producer_kernel`) activates from no-op state.
// Production wire-up of the per-step launches in
// `training_loop.rs` is deferred to a follow-up commit per
// `feedback_no_partial_refactor.md`.
let sp15_alpha_warm_count = unsafe { super::mapped_pinned::MappedF32Buffer::new(1) }
.map_err(|e| MLError::ModelError(format!(
"sp15_alpha_warm_count alloc (1 f32): {e}"
)))?;
// `MappedF32Buffer::new` already zero-fills the host side; the
// explicit write here makes the cold-start contract self-evident
// at the call site (mirrors `sel_clip_buf.write_from_slice` for
// its 1.0 fixed-clip default).
sp15_alpha_warm_count.write_from_slice(&[0.0_f32]);
// ── Per-branch Q-gap EMA (trajectory backtracking state) ────────
let per_branch_q_gap_ema_buf = stream.alloc_zeros::<f32>(4) // [4]
.map_err(|e| MLError::ModelError(format!("alloc per_branch_q_gap_ema_buf: {e}")))?;
@@ -20176,6 +20352,33 @@ impl GpuDqnTrainer {
use crate::cuda_pipeline::sp15_isv_slots::OFI_IMPACT_LAMBDA_INDEX;
*sig_ptr.add(OFI_IMPACT_LAMBDA_INDEX) = 2.0e-4_f32;
// SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline
// split slots (per spec §8.2 (3.1) post-amendment-2 fix).
// ALPHA_SPLIT (slot 417) is initialized DIRECTLY to 0.5
// — NOT derived from the formula
// `α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)` at
// boot, which would produce 0/0=0 from the zero
// grad-norm EMAs. The composer kernel
// (`r_quality_discipline_split_kernel`) enforces this
// 0.5 sentinel until `alpha_warm_count >= N_WARM=100`,
// at which point the producer kernel
// (`alpha_split_producer_kernel`) takes over and writes
// the formula's clamped output here on each step. Per
// `feedback_isv_for_adaptive_bounds.md` — the structural
// 0.5 anchor is the cold-start absorber, the runtime
// value is signal-driven. The matching grad-norm slots
// (418, 419) are stateful EMA outputs, FoldReset
// sentinel 0 so Pearl A bootstraps from the new fold's
// first observation per
// `pearl_first_observation_bootstrap.md`.
use crate::cuda_pipeline::sp15_isv_slots::{
ALPHA_SPLIT_INDEX, GRAD_NORM_DISCIPLINE_INDEX,
GRAD_NORM_QUALITY_INDEX,
};
*sig_ptr.add(ALPHA_SPLIT_INDEX) = 0.5_f32;
*sig_ptr.add(GRAD_NORM_QUALITY_INDEX) = 0.0_f32;
*sig_ptr.add(GRAD_NORM_DISCIPLINE_INDEX) = 0.0_f32;
// Layout fingerprint (ISV[58..60)). Compile-time structural hash of
// the slot layout; checkpoint load fails-fast on mismatch.
// Stored as a u64 split across two f32 lanes using raw bit-cast so
@@ -20916,6 +21119,7 @@ impl GpuDqnTrainer {
sel_clip_buf,
sel_t_pinned,
sel_t_dev_ptr,
sp15_alpha_warm_count,
per_branch_q_gap_ema_buf,
last_per_branch_q_gaps: [0.0; 4],
q_mag_means_cached: [0.0_f32; 3],

117
crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu Normal file
View File

@@ -0,0 +1,117 @@
// crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu
//
// SP15 Phase 3.1 — r_total = α × r_quality + (1 − α) × r_discipline
// (per spec §8.2 (3.1) post-amendment-2 fix).
//
// ALPHA_SPLIT slot (ISV[417]) is initialized DIRECTLY to 0.5 in the
// trainer constructor (NOT derived from formula at cold-start, which
// would produce 0/0=0 from zero grad-norm EMAs). The formula
// α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)
// only takes over once BOTH grad-norm EMAs have accumulated ≥ N_WARM
// non-zero observations. Until then, the slot remains at 0.5 and the
// composer reads the constructor-written value through the
// r_quality_discipline_split_kernel below.
//
// Two `extern "C" __global__` symbols share this single cubin (the
// established 1:1-source-to-cubin / multi-kernel-per-file pattern used
// by `baseline_kernels.cu` and `novelty_simhash_kernel.cu`):
//
// r_quality_discipline_split_kernel — per-step composition + warm count
// • Reads ALPHA_SPLIT_INDEX (slot 417)
// • Reads/writes alpha_warm_count scratch buffer ([1])
// • While `*alpha_warm_count < N_WARM`, OVERRIDES alpha to 0.5 (the
// cold-start sentinel) regardless of what the producer kernel may
// have written — this protects the composer from premature formula
// activation when the grad-norm EMAs are still warming up.
// • Writes r_total to r_total_out
//
// alpha_split_producer_kernel — per-step ALPHA_SPLIT update from grad
// ratio
// • Returns immediately if `*alpha_warm_count < N_WARM` (the
// composer alone owns the warm-count increment; the producer is a
// silent no-op until the composer has fully warmed up)
// • Reads GRAD_NORM_QUALITY_INDEX (slot 418), GRAD_NORM_DISCIPLINE_INDEX
// (slot 419)
// • Writes ALPHA_SPLIT_INDEX (slot 417) clamped to [0.05, 0.95] so
// neither signal can totally dominate the composer
//
// Hard rules:
// • feedback_no_atomicadd — single-thread/single-block; no reductions.
// • feedback_isv_for_adaptive_bounds — N_WARM is the only structural
// constant; α itself is signal-driven.
// • feedback_first_observation_bootstrap — alpha_warm_count starts at 0;
// the warm-up window is the EMA cold-start absorber.
#include <cuda_runtime.h>
// Mirror the named slot constants from
// `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs`. Keep these literals in
// lockstep with the Rust constants — the Rust unit test
// `all_sp15_slots_fit_within_isv_total_dim` enforces ISV_TOTAL_DIM bounds
// and `sp15_slot_layout_locked` pins each named index, so any drift here
// shows up as an immediate test failure on the Rust side.
#define ALPHA_SPLIT_INDEX 417
#define GRAD_NORM_QUALITY_INDEX 418
#define GRAD_NORM_DISCIPLINE_INDEX 419
// Welford warm-up count. After both grad-norm EMAs have accumulated this
// many non-zero observations the producer kernel takes over. Until then
// the composer enforces the cold-start sentinel α=0.5.
#define N_WARM 100.0f
extern "C" __global__ void r_quality_discipline_split_kernel(
float r_quality,
float r_discipline,
const float* __restrict__ isv,
float* __restrict__ r_total_out,
float* __restrict__ alpha_warm_count
) {
if (threadIdx.x != 0 || blockIdx.x != 0) return;
// Load the constructor-written sentinel (or, post-warmup, the
// producer-written value).
float alpha = isv[ALPHA_SPLIT_INDEX];
// Cold-start absorber. The composer is the SOLE owner of the warm
// count — the producer reads it but never increments. This guarantees
// exactly one increment per call to this kernel even though both
// kernels alias the same buffer.
float warm = *alpha_warm_count;
if (warm < N_WARM) {
alpha = 0.5f;
*alpha_warm_count = warm + 1.0f;
}
*r_total_out = alpha * r_quality + (1.0f - alpha) * r_discipline;
// Make the slot writes visible to the host on the mapped-pinned
// buffers (same pattern used by the SP15 cost-net / dd-state kernels).
__threadfence_system();
}
extern "C" __global__ void alpha_split_producer_kernel(
float* __restrict__ isv,
const float* __restrict__ alpha_warm_count
) {
if (threadIdx.x != 0 || blockIdx.x != 0) return;
// Wait out the cold-start window. While the composer is overriding
// alpha to 0.5 (warm < N_WARM), writing the formula's output to the
// slot would be wasted work and could leak a transient value if the
// composer somehow ran on a stale warm count — leave the slot at the
// constructor-written 0.5 sentinel during warmup.
if (*alpha_warm_count < N_WARM) return;
float gn_q = isv[GRAD_NORM_QUALITY_INDEX];
float gn_d = isv[GRAD_NORM_DISCIPLINE_INDEX];
const float EPS = 1e-6f;
float alpha = gn_q / (gn_q + gn_d + EPS);
// Clamp to [0.05, 0.95] so neither head can totally dominate the
// composer — bilateral clamp per `pearl_symmetric_clamp_audit.md`.
alpha = fmaxf(0.05f, fminf(alpha, 0.95f));
isv[ALPHA_SPLIT_INDEX] = alpha;
__threadfence_system();
}

47
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -1088,6 +1088,53 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "ISV[DD_PCT_INDEX=406] — SP15 Phase 1.3 (2026-05-06) `clip(current_dd / max(dd_budget, 1e-4), 0, 1)` ∈ [0, 1]. Stateful kernel output; FoldReset sentinel 0 (cold-start at fold boundary; the kernel overwrites with the true value on the new fold's first launch). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.3.",
},
// ── SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split ───
// Three ISV slots [417..420) + one mapped-pinned scratch buffer
// (`sp15_alpha_warm_count`):
// - ALPHA_SPLIT_INDEX=417: Invariant-1 sentinel anchor (NOT a
// stateful EMA). Constructor-seeded to 0.5 per spec §8.2 (3.1)
// post-amendment-2 fix; FoldReset rewrites the same 0.5 anchor
// so the composer's cold-start absorber re-engages on each new
// fold (the warm-count goes back to 0; the producer kernel
// becomes a no-op until `*warm_count >= N_WARM=100` again).
// Per `feedback_isv_for_adaptive_bounds.md` — the structural
// 0.5 anchor is the cold-start absorber, the runtime value is
// signal-driven once warmup completes.
// - GRAD_NORM_QUALITY_INDEX=418: stateful EMA of the r_quality
// gradient L2 norm. FoldReset sentinel 0 so Pearl A
// bootstraps from the new fold's first observation per
// `pearl_first_observation_bootstrap.md`.
// - GRAD_NORM_DISCIPLINE_INDEX=419: stateful EMA of the
// r_discipline gradient L2 norm. FoldReset sentinel 0 same
// rationale as quality.
// - sp15_alpha_warm_count (NOT an ISV slot — single-element
// mapped-pinned scratch buffer on the trainer struct):
// FoldReset sentinel 0.0 via `host_slice_mut().fill(0.0)`
// so the new fold's first composer launch re-enters the
// warm-up window from step 0 and the producer-kernel gate
// re-engages. Mirrors the `sp11_novelty_hash` pattern of
// resetting a non-ISV mapped-pinned buffer at fold boundary.
// Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).
RegistryEntry {
name: "sp15_alpha_split",
category: ResetCategory::FoldReset,
description: "ISV[ALPHA_SPLIT_INDEX=417] — SP15 Phase 3.1 (2026-05-06) blend coefficient for `r_total = α × r_quality + (1 − α) × r_discipline`. Constant Invariant-1 anchor (0.5) per `feedback_isv_for_adaptive_bounds.md` — NOT a stateful EMA. Constructor-seeded to 0.5; FoldReset rewrites the same anchor so the composer's cold-start absorber re-engages on each new fold (per spec §8.2 (3.1) post-amendment-2 fix: deriving α from `grad_norm_q / (grad_norm_q + grad_norm_d + ε)` at boot would produce 0/0=0 from zero grad-norm EMAs). The companion `alpha_split_producer_kernel` is gated on `sp15_alpha_warm_count` and stays a silent no-op until `*warm_count >= N_WARM=100`, at which point it activates and writes the formula's clamped output here on each step. Per-fold ISV refit may overwrite this in later phases. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
},
RegistryEntry {
name: "sp15_grad_norm_quality",
category: ResetCategory::FoldReset,
description: "ISV[GRAD_NORM_QUALITY_INDEX=418] — SP15 Phase 3.1 (2026-05-06) stateful EMA of the r_quality gradient L2 norm. Read by `alpha_split_producer_kernel` (post-warmup) to compute `α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)`. FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md` (also drives the warm-count gate: `alpha_split_producer_kernel` will re-enter no-op state at fold boundary because the companion `sp15_alpha_warm_count` resets in the same arm cluster). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
},
RegistryEntry {
name: "sp15_grad_norm_discipline",
category: ResetCategory::FoldReset,
description: "ISV[GRAD_NORM_DISCIPLINE_INDEX=419] — SP15 Phase 3.1 (2026-05-06) stateful EMA of the r_discipline gradient L2 norm. Read by `alpha_split_producer_kernel` (post-warmup) as the denominator partner of `sp15_grad_norm_quality`. FoldReset sentinel 0 same rationale as quality. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
},
RegistryEntry {
name: "sp15_alpha_warm_count",
category: ResetCategory::FoldReset,
description: "GpuDqnTrainer.sp15_alpha_warm_count [1] f32 mapped-pinned scratch buffer — SP15 Phase 3.1 (2026-05-06) warm-up counter for the r_quality + r_discipline split composer (per spec §8.2 (3.1) post-amendment-2 fix). Initialised to 0.0 in the constructor; incremented by 1 each step inside `r_quality_discipline_split_kernel`; once `*count >= N_WARM=100` the composer stops overriding α to 0.5 and `alpha_split_producer_kernel` activates from no-op state. FoldReset sentinel 0.0 via `host_slice_mut().fill(0.0)` (mirrors the `sp11_novelty_hash` pattern of resetting a non-ISV mapped-pinned buffer at fold boundary) — the new fold's first composer launch must re-enter the warm-up window from step 0 so the cold-start absorber holds α at 0.5 long enough for the new fold's grad-norm EMAs to accumulate ≥100 non-zero observations before the producer kernel takes over. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
},
// SP5 Task A1: Wiener-state companion reset. The wiener_state_buf
// covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) +
// SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3

49
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -8141,6 +8141,55 @@ impl DQNTrainer {
fused.trainer().write_isv_signal_at(DD_PCT_INDEX, 0.0);
}
}
// SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split
// dispatch arms (per spec §8.2 (3.1) post-amendment-2 fix).
// ALPHA_SPLIT_INDEX=417 is an Invariant-1 anchor (NOT a stateful
// EMA) — rewrite the constructor's 0.5 sentinel at fold boundary
// so the composer's cold-start absorber re-engages on each new
// fold; deriving from `grad_norm_q / (grad_norm_q + grad_norm_d
// + ε)` at boot would produce 0/0=0 from zero grad-norm EMAs.
// Per `feedback_isv_for_adaptive_bounds.md`. Per-fold ISV refit
// may overwrite this in later phases.
"sp15_alpha_split" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp15_isv_slots::ALPHA_SPLIT_INDEX;
fused.trainer().write_isv_signal_at(ALPHA_SPLIT_INDEX, 0.5);
}
}
// GRAD_NORM_QUALITY_INDEX=418 / GRAD_NORM_DISCIPLINE_INDEX=419
// are stateful EMAs — FoldReset sentinel 0 so Pearl A
// bootstraps from the new fold's first observation per
// `pearl_first_observation_bootstrap.md`. The companion
// `sp15_alpha_warm_count` reset re-arms the warm-up gate so
// the producer kernel re-enters no-op state until both EMAs
// accumulate ≥100 non-zero observations on the new fold.
"sp15_grad_norm_quality" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp15_isv_slots::GRAD_NORM_QUALITY_INDEX;
fused.trainer().write_isv_signal_at(GRAD_NORM_QUALITY_INDEX, 0.0);
}
}
"sp15_grad_norm_discipline" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp15_isv_slots::GRAD_NORM_DISCIPLINE_INDEX;
fused.trainer().write_isv_signal_at(GRAD_NORM_DISCIPLINE_INDEX, 0.0);
}
}
// sp15_alpha_warm_count is a non-ISV mapped-pinned scratch
// buffer on the trainer struct. Reset to 0.0 via
// `host_slice_mut().fill(0.0)` (allowed under
// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write
// rule, NOT a host-side compute) — mirrors the
// `sp11_novelty_hash` pattern of resetting a non-ISV
// mapped-pinned buffer at fold boundary. The new fold's first
// composer launch must re-enter the warm-up window from step 0
// so the cold-start absorber holds α at 0.5 long enough for
// the new fold's grad-norm EMAs to accumulate.
"sp15_alpha_warm_count" => {
if let Some(ref mut fused) = self.fused_ctx {
fused.trainer_mut().sp15_alpha_warm_count.host_slice_mut().fill(0.0);
}
}
// SP11 Task A2 (2026-05-04): novelty visit-count hash table
// reset arm — closes the A0 deferral per the registry entry's
// docstring. 1M f32 slots, zero-filled via mapped-pinned host

64
crates/ml/tests/sp15_phase1_oracle_tests.rs
View File

@@ -17,10 +17,11 @@ mod gpu {
};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedU32Buffer};
use ml::cuda_pipeline::sp15_isv_slots::{
BASELINE_BUYHOLD_SHARPE_INDEX, BASELINE_HOLD_ONLY_SHARPE_INDEX,
ALPHA_SPLIT_INDEX, BASELINE_BUYHOLD_SHARPE_INDEX, BASELINE_HOLD_ONLY_SHARPE_INDEX,
BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX, BASELINE_NAIVE_REVERSION_SHARPE_INDEX,
COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, DD_PCT_INDEX,
DD_RECOVERY_BARS_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
DD_RECOVERY_BARS_INDEX, GRAD_NORM_DISCIPLINE_INDEX, GRAD_NORM_QUALITY_INDEX,
OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
};
/// Upper bound for any ISV index touched by SP15 oracle tests. Matches
@@ -605,6 +606,65 @@ mod gpu {
);
}
}
/// Test 3.1.a — `r_total = 0.5 × r_quality + 0.5 × r_discipline` at
/// cold start. Per spec §8.2 (3.1) post-amendment-2 fix: the
/// ALPHA_SPLIT slot (ISV[417]) is constructor-written to 0.5 and the
/// composer's warm-count override holds it there until both
/// grad-norm EMAs accumulate ≥ N_WARM=100 observations. This test
/// validates the cold-start behavior in isolation: warm count starts
/// at 0, so the composer must override α to 0.5 regardless of any
/// other ISV state, AND increment the warm count to 1.
#[test]
#[ignore = "requires GPU"]
fn r_split_uses_sentinel_alpha_at_cold_start() {
let stream = make_test_stream();
// Safety: CUDA context active via `make_test_stream` above.
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
.expect("alloc MappedF32Buffer for isv bus");
let mut isv_init = vec![0.0f32; ISV_LEN];
// Mirror the trainer constructor's ALPHA_SPLIT seed of 0.5; the
// grad-norm slots stay at the zero sentinel (cold start, no
// observations yet).
isv_init[ALPHA_SPLIT_INDEX] = 0.5;
isv_init[GRAD_NORM_QUALITY_INDEX] = 0.0;
isv_init[GRAD_NORM_DISCIPLINE_INDEX] = 0.0;
isv_buf.write_from_slice(&isv_init);
let alpha_warm_count_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc MappedF32Buffer for alpha_warm_count");
alpha_warm_count_buf.write_from_slice(&[0.0f32]);
let r_total_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc MappedF32Buffer for r_total");
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_r_quality_discipline_split(
&stream,
/* r_quality = */ 1.0,
/* r_discipline = */ -2.0,
isv_buf.dev_ptr,
r_total_buf.dev_ptr,
alpha_warm_count_buf.dev_ptr,
)
.expect("launch r_quality_discipline_split_kernel");
stream
.synchronize()
.expect("synchronize after r_quality_discipline_split_kernel launch");
// r_total = 0.5 × 1.0 + 0.5 × (-2.0) = -0.5
let r_total = r_total_buf.read_all()[0];
assert!(
(r_total - (-0.5)).abs() < 1e-5,
"r_total = {r_total}, expected -0.5"
);
// Warm count: 0 → 1 after the single composer launch.
let warm = alpha_warm_count_buf.read_all()[0];
assert!(
(warm - 1.0).abs() < 1e-5,
"warm count = {warm}, expected 1"
);
}
}
/// SP15 Phase 1.7 — verify the trainer's `set_test_data_from_slices` API

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

File diff suppressed because one or more lines are too long