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

fix(sp16-t3): Wiener-optimal adaptive α per pearl — hold_cost_scale + min_hold_temperature

Browse Source 
Per train-multi-seed-hjzss validation: SP16 T1+T2 chain was structurally
landed but BEHAVIORALLY INERT in 5-epoch smoke (bit-identical to pfh9n
baseline through epoch 3). Root cause: hardcoded `alpha = 0.05f` in both
producer kernels violates feedback_isv_for_adaptive_bounds AND prevents
convergence in short runs (~60 epochs needed from cold start).

Fix per pearl_wiener_optimal_adaptive_alpha:
  α = diff_var / (diff_var + sample_var + ε)

Where sample_var = running variance of target signal (Welford accumulator)
and diff_var = running variance of consecutive one-step differences.

Cold-start: target jumps 1.0 → 6.4 → 7.0 → high diff_var → α ≈ 0.6+
            → near-bootstrap responsiveness in epochs 1-3
Steady-state: signal stabilizes → diff_var drops → α decays naturally
              → smoothing emerges without hardcoded constant

Adds 12 new ISV slots (6 per producer):
- HCS_TARGET_MEAN/M2, HCS_DIFF_MEAN/M2, HCS_PREV_TARGET, HCS_SAMPLE_COUNT
- MHT_TARGET_MEAN/M2, MHT_DIFF_MEAN/M2, MHT_PREV_TARGET, MHT_SAMPLE_COUNT

ISV_TOTAL_DIM 462 → 474.

Both kernels migrated atomically. Pearl-A bootstrap preserved (sentinel
on prev_blended triggers REPLACE; cold-start α=1.0 when N<3 samples).
Defensive bounds [WELFORD_ALPHA_MIN=0.01, WELFORD_ALPHA_MAX=0.95] on the
Wiener-derived α to guard against denormal/underflow corner cases.

HEALTH_DIAG[N] emit extended with `alpha=...` and `sample_count=...` for
direct trajectory observation in validation smoke.

Behavioral tests verify:
- α high during signal jumps (>0.3 at epoch 3 post-cold-start)
- α low in steady state (mean tail α<0.4 under converging signal)
- Pearl-A bootstrap fires on first observation (Welford state advances
  regardless of REPLACE branch)
- α stays within [WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX] over 50 epochs
  (post-cold-start; cold-start α=1.0 by design)
- No 0.05f hardcoded literal remains in blend math (regression-locked
  via host-only string scan)

5 GPU + host tests pass: sp16_phase3_alpha_high_during_signal_jump,
alpha_low_in_steady_state, pearl_a_bootstrap_first_obs,
alpha_naturally_bounded, no_hardcoded_alpha. sp14 + sp15 oracle suites
unchanged (34 GPU tests + 4 host tests).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-08 18:56:58 +02:00
parent 6dce4ac28b
commit 641aa0dfde
11 changed files with 1469 additions and 256 deletions
Download Patch File Download Diff File Expand all files Collapse all files

51
crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs
View File

@@ -1048,19 +1048,26 @@ impl MinHoldTemperatureUpdateOps {
} }
/// Launch the adaptive MIN_HOLD_TEMPERATURE producer (SP16 Phase 1 /// Launch the adaptive MIN_HOLD_TEMPERATURE producer (SP16 Phase 1
/// revised — hold-rate-overrun driven). /// revised — hold-rate-overrun driven; SP16 T3 — Wiener-optimal α).
/// ///
/// Args: /// Args:
/// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned).
/// - `temp_idx`: MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460). /// - `temp_idx`: MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460).
/// - `hold_rate_observed_idx`: HOLD_RATE_OBSERVED_EMA_INDEX (382). /// - `hold_rate_observed_idx`: HOLD_RATE_OBSERVED_EMA_INDEX (382).
/// - `hold_rate_target_idx`: HOLD_RATE_TARGET_INDEX (381). /// - `hold_rate_target_idx`: HOLD_RATE_TARGET_INDEX (381).
/// - Welford accumulator slot indices (T3): MHT_TARGET_MEAN_INDEX
/// (468), MHT_TARGET_M2_INDEX (469), MHT_DIFF_MEAN_INDEX (470),
/// MHT_DIFF_M2_INDEX (471), MHT_PREV_TARGET_INDEX (472),
/// MHT_SAMPLE_COUNT_INDEX (473).
/// - `sentinel_temp`: SENTINEL_MIN_HOLD_TEMPERATURE (50.0). /// - `sentinel_temp`: SENTINEL_MIN_HOLD_TEMPERATURE (50.0).
/// - `sentinel_observed_hold_rate`: SENTINEL_HOLD_RATE_OBSERVED /// - `sentinel_observed_hold_rate`: SENTINEL_HOLD_RATE_OBSERVED
/// (0.0; matches the SP13 fold-reset registry sentinel for /// (0.0; matches the SP13 fold-reset registry sentinel for
/// `sp13_hold_rate_observed_ema`). /// `sp13_hold_rate_observed_ema`).
/// - `temp_min`/`temp_max`: MIN_HOLD_TEMPERATURE_MIN/MAX (5.0, 50.0). /// - `temp_min`/`temp_max`: MIN_HOLD_TEMPERATURE_MIN/MAX (5.0, 50.0).
/// - `alpha`: MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05). ///
/// SP16 T3: the `alpha` argument is gone — α is computed inside the
/// kernel from the Welford accumulators (Wiener-optimal: α = diff_var
/// / (diff_var + sample_var + ε)).
#[allow(clippy::too_many_arguments)] #[allow(clippy::too_many_arguments)]
pub(crate) fn launch( pub(crate) fn launch(
&self, &self,
@@ -1069,11 +1076,16 @@ impl MinHoldTemperatureUpdateOps {
temp_idx: i32, temp_idx: i32,
hold_rate_observed_idx: i32, hold_rate_observed_idx: i32,
hold_rate_target_idx: i32, hold_rate_target_idx: i32,
target_mean_idx: i32,
target_m2_idx: i32,
diff_mean_idx: i32,
diff_m2_idx: i32,
prev_target_idx: i32,
sample_count_idx: i32,
sentinel_temp: f32, sentinel_temp: f32,
sentinel_observed_hold_rate: f32, sentinel_observed_hold_rate: f32,
temp_min: f32, temp_min: f32,
temp_max: f32, temp_max: f32,
alpha: f32,
) -> Result<(), MLError> { ) -> Result<(), MLError> {
unsafe { unsafe {
stream stream
@@ -1082,11 +1094,16 @@ impl MinHoldTemperatureUpdateOps {
.arg(&temp_idx) .arg(&temp_idx)
.arg(&hold_rate_observed_idx) .arg(&hold_rate_observed_idx)
.arg(&hold_rate_target_idx) .arg(&hold_rate_target_idx)
.arg(&target_mean_idx)
.arg(&target_m2_idx)
.arg(&diff_mean_idx)
.arg(&diff_m2_idx)
.arg(&prev_target_idx)
.arg(&sample_count_idx)
.arg(&sentinel_temp) .arg(&sentinel_temp)
.arg(&sentinel_observed_hold_rate) .arg(&sentinel_observed_hold_rate)
.arg(&temp_min) .arg(&temp_min)
.arg(&temp_max) .arg(&temp_max)
.arg(&alpha)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (1, 1, 1), block_dim: (1, 1, 1),
@@ -1237,19 +1254,27 @@ impl HoldCostScaleUpdateOps {
Ok(Self { update_kernel }) Ok(Self { update_kernel })
} }
/// Launch the adaptive Hold cost scale producer. /// Launch the adaptive Hold cost scale producer (SP16 Phase 2 +
/// T3 Wiener-optimal α).
/// ///
/// Args: /// Args:
/// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned).
/// - `scale_idx`: HOLD_COST_SCALE_INDEX (461). /// - `scale_idx`: HOLD_COST_SCALE_INDEX (461).
/// - `hold_rate_observed_idx`: HOLD_RATE_OBSERVED_EMA_INDEX (382). /// - `hold_rate_observed_idx`: HOLD_RATE_OBSERVED_EMA_INDEX (382).
/// - `hold_rate_target_idx`: HOLD_RATE_TARGET_INDEX (381). /// - `hold_rate_target_idx`: HOLD_RATE_TARGET_INDEX (381).
/// - Welford accumulator slot indices (T3): HCS_TARGET_MEAN_INDEX
/// (462), HCS_TARGET_M2_INDEX (463), HCS_DIFF_MEAN_INDEX (464),
/// HCS_DIFF_M2_INDEX (465), HCS_PREV_TARGET_INDEX (466),
/// HCS_SAMPLE_COUNT_INDEX (467).
/// - `sentinel_scale`: SENTINEL_HOLD_COST_SCALE (0.0). /// - `sentinel_scale`: SENTINEL_HOLD_COST_SCALE (0.0).
/// - `sentinel_observed_hold_rate`: SENTINEL_HOLD_RATE_OBSERVED /// - `sentinel_observed_hold_rate`: SENTINEL_HOLD_RATE_OBSERVED
/// (0.0; matches the SP13 fold-reset registry sentinel — shared /// (0.0; matches the SP13 fold-reset registry sentinel — shared
/// by design with SP16-P1). /// by design with SP16-P1).
/// - `scale_min`/`scale_max`: HOLD_COST_SCALE_MIN/MAX (1.0, 25.0). /// - `scale_min`/`scale_max`: HOLD_COST_SCALE_MIN/MAX (1.0, 25.0).
/// - `alpha`: HOLD_COST_SCALE_EMA_ALPHA (0.05). ///
/// SP16 T3: the `alpha` argument is gone — α is computed inside the
/// kernel from the Welford accumulators (Wiener-optimal: α = diff_var
/// / (diff_var + sample_var + ε)).
#[allow(clippy::too_many_arguments)] #[allow(clippy::too_many_arguments)]
pub(crate) fn launch( pub(crate) fn launch(
&self, &self,
@@ -1258,11 +1283,16 @@ impl HoldCostScaleUpdateOps {
scale_idx: i32, scale_idx: i32,
hold_rate_observed_idx: i32, hold_rate_observed_idx: i32,
hold_rate_target_idx: i32, hold_rate_target_idx: i32,
target_mean_idx: i32,
target_m2_idx: i32,
diff_mean_idx: i32,
diff_m2_idx: i32,
prev_target_idx: i32,
sample_count_idx: i32,
sentinel_scale: f32, sentinel_scale: f32,
sentinel_observed_hold_rate: f32, sentinel_observed_hold_rate: f32,
scale_min: f32, scale_min: f32,
scale_max: f32, scale_max: f32,
alpha: f32,
) -> Result<(), MLError> { ) -> Result<(), MLError> {
unsafe { unsafe {
stream stream
@@ -1271,11 +1301,16 @@ impl HoldCostScaleUpdateOps {
.arg(&scale_idx) .arg(&scale_idx)
.arg(&hold_rate_observed_idx) .arg(&hold_rate_observed_idx)
.arg(&hold_rate_target_idx) .arg(&hold_rate_target_idx)
.arg(&target_mean_idx)
.arg(&target_m2_idx)
.arg(&diff_mean_idx)
.arg(&diff_m2_idx)
.arg(&prev_target_idx)
.arg(&sample_count_idx)
.arg(&sentinel_scale) .arg(&sentinel_scale)
.arg(&sentinel_observed_hold_rate) .arg(&sentinel_observed_hold_rate)
.arg(&scale_min) .arg(&scale_min)
.arg(&scale_max) .arg(&scale_max)
.arg(&alpha)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (1, 1, 1), block_dim: (1, 1, 1),

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

File diff suppressed because one or more lines are too long

223
crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu
View File

@@ -1,117 +1,96 @@
/* ══════════════════════════════════════════════════════════════════════════ /* ══════════════════════════════════════════════════════════════════════════
* SP16 Phase 2 (2026-05-08) — adaptive Hold cost scale producer. * SP16 Phase 2 (2026-05-08) + T3 (Wiener-optimal α, 2026-05-08) — adaptive
* Hold cost scale producer.
* *
* Mirrors the SP16 Phase 1 (revised) MIN_HOLD_TEMPERATURE producer driver * Mirrors the SP16 Phase 1 (revised) MIN_HOLD_TEMPERATURE producer
* pattern (`min_hold_temperature_update_kernel.cu`) — same input signals * driver pattern (`min_hold_temperature_update_kernel.cu`) — same
* (HOLD_RATE_OBSERVED_EMA + HOLD_RATE_TARGET), same overrun-norm shape, * input signals (HOLD_RATE_OBSERVED_EMA + HOLD_RATE_TARGET), same
* same cadence (per-epoch boundary), same Pearl-A bootstrap + Welford * overrun-norm shape, same cadence (per-epoch boundary), same Pearl-A
* EMA, but maps the overrun into the Hold-cost-scale dimension instead * bootstrap, same Wiener-optimal α (T3) — but maps the overrun into
* of the min-hold-temperature dimension. * the Hold-cost-scale dimension instead of the min-hold-temperature
* dimension.
* *
* Why this kernel exists (post-mortem of train-multi-seed-pfh9n): * Mapping (unchanged from T2):
*
* HEALTH_DIAG[2]: hold_pricing observed_rate=0.2501 target=0.2000
* cost=0.006251
*
* The realised Hold rate climbed 0.25 → 0.52 across training while the
* cost penalty stayed at ~0.006 — ~100× smaller than per-bar reward
* magnitudes (popart=0.97, cf=0.65 in `reward_split`). Hold became
* effectively free, allowing the structural Q-bias toward zero-variance
* Hold to dominate. This kernel addresses the cost-magnitude axis: when
* the model is over-holding (overrun > 0), the per-bar Hold cost is
* scaled up multiplicatively at the 3 consumer sites in
* `experience_kernels.cu` (segment_complete branch line ~3089, per-bar
* positioned-Hold branch line ~3553, per-bar flat-Hold branch line ~3617).
*
* Mapping:
* *
* overrun = max(0, observed − target) * overrun = max(0, observed − target)
* overrun_normalized = clamp(overrun / max(target, 0.01), 0, 1) * overrun_normalized = clamp(overrun / max(target, 0.01), 0, 1)
* target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_normalized * target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_normalized
* *
* - observed = 2× target → overrun_norm = 1.0 → scale = 25 * SP16 T3 — Wiener-optimal adaptive α (2026-05-08):
* (effective cost ~0.006 × 25 = 0.15/bar — competitive with per-bar
* reward magnitudes ~0.01-0.1 but well below SP12 cap [-10, +5])
* - observed = target → overrun_norm = 0 → scale = 1
* (no extra penalty — bit-identical to pre-Phase-2 cost magnitude)
* - observed < target → overrun_norm = 0 → scale = 1
* (under-holding — no penalty escalation needed)
* *
* Why bounds [1, 25] are dimensional safety, not tuning: * The pre-T3 implementation hardcoded `const float alpha = 0.05f` for
* - SCALE_MIN=1: the floor is "no extra penalty" — lower bound enforced * the EMA blend. train-multi-seed-hjzss validation showed this α was
* by the formula (overrun is clamped at 0 from below). Below 1 would * far too low to converge in 5-epoch smoke. Per
* mean "discount the cost below baseline" which is the opposite of * `pearl_wiener_optimal_adaptive_alpha`:
* what the controller is supposed to do under overrun. *
* - SCALE_MAX=25: at this ceiling, the effective Hold cost is ~0.15/bar. * α = diff_var / (diff_var + sample_var + ε)
* Above 25× the static base, single Hold bars start dominating per- *
* bar reward magnitudes by an order of magnitude — overcorrection * Where `sample_var` = running variance of target_scale and
* territory. The audit-doc 25× ceiling matches the ratio found in * `diff_var` = running variance of consecutive one-step differences.
* the pfh9n post-mortem (cost ~0.006, reward ~0.97 → ratio ~150 → * Welford accumulators stored in ISV slots [462..468):
* tightening to 6:1 via 25× scale is the conservative point). *
* - HCS_TARGET_MEAN_INDEX (462) — running mean of target_scale
* - HCS_TARGET_M2_INDEX (463) — Welford ΣΔ² (target_scale)
* - HCS_DIFF_MEAN_INDEX (464) — running mean of consecutive diffs
* - HCS_DIFF_M2_INDEX (465) — Welford ΣΔ² (diff)
* - HCS_PREV_TARGET_INDEX (466) — previous-epoch target_scale
* - HCS_SAMPLE_COUNT_INDEX (467) — Welford counter
*
* Cold-start path: when `sample_count < 3`, α = 1.0 → REPLACE. After
* the third observation, sample_var and diff_var are well-defined and
* α emerges from the data.
*
* Defensive bounds [WELFORD_ALPHA_MIN=0.01, WELFORD_ALPHA_MAX=0.95]
* applied to the Wiener-derived α before the blend.
* *
* Algorithm (single-block single-thread — cold path, per-epoch boundary): * Algorithm (single-block single-thread — cold path, per-epoch boundary):
* * Phases 1-7 mirror the MHT producer; only the slot indices and
* Phase 1 (single thread): read observed_hold_rate (slot 382), * bound constants differ. See `min_hold_temperature_update_kernel.cu`
* target_hold_rate (slot 381), and current scale (slot 461) from ISV. * for the canonical algorithm comment.
*
* Phase 2 (sentinel guard): if observed_hold_rate is at sentinel 0.0
* (no per-step Hold observations have landed yet, e.g. epoch 0 of
* any fold before the first hold_rate_observer launch chains in),
* keep ISV slot 461 unchanged — the consumer falls back to scale=1.0
* when slot 461 is at sentinel 0.0 (bit-identical pre-Phase-2 cost
* magnitude). This is the bootstrap-protection equivalent of the
* SP16-P1 dir_acc sentinel guard.
*
* Phase 3 (compute target scale):
*
* overrun = fmaxf(0, observed − target)
* overrun_norm = fminf(1, overrun / fmaxf(target, 0.01))
* target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_norm
* bilateral clamp to [SCALE_MIN, SCALE_MAX]
*
* Phase 4 (Pearl-A bootstrap + Welford-style EMA):
*
* if current is at sentinel 0.0 (within EPS) → REPLACE with target_scale.
* else → blended = (1 − α) × current + α × target_scale
* re-clamp post-blend.
* *
* Pearls + invariants: * Pearls + invariants:
* *
* - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. * - `feedback_no_atomicadd.md` — single-thread kernel; no reductions.
* *
* - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` * - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction`
* pre-loaded at construction; on-device guards for sentinel * pre-loaded at construction; on-device guards.
* detection.
* *
* - `pearl_first_observation_bootstrap.md` — first valid hold-rate * - `pearl_first_observation_bootstrap.md` — first valid hold-rate
* observation replaces sentinel 0.0 directly; no blend. * observation replaces sentinel 0.0 directly; no blend.
* *
* - `pearl_wiener_optimal_adaptive_alpha.md` (NEW T3) — α derived
* from Welford variance accumulators, not hardcoded.
*
* - `feedback_no_stubs.md` — full body, no return-zero placeholders. * - `feedback_no_stubs.md` — full body, no return-zero placeholders.
* *
* - `feedback_isv_for_adaptive_bounds.md` — bounds [1, 25] are * - `feedback_isv_for_adaptive_bounds.md` — bounds [1, 25] are
* dimensional safety floors (not tuning constants). * dimensional safety floors; α is no longer hardcoded.
* *
* - `pearl_symmetric_clamp_audit.md` — bilateral * - `pearl_symmetric_clamp_audit.md` — bilateral
* `fmaxf(lo, fminf(x, hi))` clamp on target_scale before writing * `fmaxf(lo, fminf(x, hi))` clamp on target_scale before writing
* and post-blend. * and post-blend.
* *
* - `feedback_no_partial_refactor.md` — kernel + launcher + 3 * - `feedback_no_partial_refactor.md` — kernel + launcher + 3
* consumer sites + tests + audit doc updated atomically in the * consumer sites + tests + audit doc updated atomically; both
* same commit. * T1 and T2 producers migrate together to Wiener-optimal α.
* *
* Args: * Args:
* isv — `[ISV_TOTAL_DIM]` device f32, modified * isv — `[ISV_TOTAL_DIM]` device f32, modified
* in place at index `scale_idx`. * in place at index `scale_idx` and 6
* Welford slots [462..468).
* scale_idx — HOLD_COST_SCALE_INDEX (461). * scale_idx — HOLD_COST_SCALE_INDEX (461).
* hold_rate_observed_idx — HOLD_RATE_OBSERVED_EMA_INDEX (382). * hold_rate_observed_idx — HOLD_RATE_OBSERVED_EMA_INDEX (382).
* hold_rate_target_idx — HOLD_RATE_TARGET_INDEX (381). * hold_rate_target_idx — HOLD_RATE_TARGET_INDEX (381).
* target_mean_idx — HCS_TARGET_MEAN_INDEX (462).
* target_M2_idx — HCS_TARGET_M2_INDEX (463).
* diff_mean_idx — HCS_DIFF_MEAN_INDEX (464).
* diff_M2_idx — HCS_DIFF_M2_INDEX (465).
* prev_target_idx — HCS_PREV_TARGET_INDEX (466).
* sample_count_idx — HCS_SAMPLE_COUNT_INDEX (467).
* sentinel_scale — SENTINEL_HOLD_COST_SCALE (0.0). * sentinel_scale — SENTINEL_HOLD_COST_SCALE (0.0).
* sentinel_observed_hold_rate — SENTINEL_HOLD_RATE_OBSERVED (0.0; matches * sentinel_observed_hold_rate — SENTINEL_HOLD_RATE_OBSERVED (0.0).
* the SP13 hold-rate observer cold-start —
* shared sentinel by design with SP16-P1).
* scale_min — HOLD_COST_SCALE_MIN (1.0). * scale_min — HOLD_COST_SCALE_MIN (1.0).
* scale_max — HOLD_COST_SCALE_MAX (25.0). * scale_max — HOLD_COST_SCALE_MAX (25.0).
* alpha — HOLD_COST_SCALE_EMA_ALPHA (0.05).
* *
* Launch: grid=(1, 1, 1), block=(1, 1, 1). * Launch: grid=(1, 1, 1), block=(1, 1, 1).
* No shared memory — single-thread cold-path kernel. * No shared memory — single-thread cold-path kernel.
@@ -120,6 +99,10 @@
#include <cuda_runtime.h> #include <cuda_runtime.h>
#define EPS_F 1e-6f #define EPS_F 1e-6f
#define WELFORD_ALPHA_MIN_F 0.01f
#define WELFORD_ALPHA_MAX_F 0.95f
#define WELFORD_MIN_SAMPLE_COUNT_F 3.0f
#define WELFORD_COLD_START_ALPHA_F 1.0f
extern "C" __global__ extern "C" __global__
void hold_cost_scale_update( void hold_cost_scale_update(
@@ -127,11 +110,16 @@ void hold_cost_scale_update(
int scale_idx, int scale_idx,
int hold_rate_observed_idx, int hold_rate_observed_idx,
int hold_rate_target_idx, int hold_rate_target_idx,
int target_mean_idx,
int target_M2_idx,
int diff_mean_idx,
int diff_M2_idx,
int prev_target_idx,
int sample_count_idx,
float sentinel_scale, float sentinel_scale,
float sentinel_observed_hold_rate, float sentinel_observed_hold_rate,
float scale_min, float scale_min,
float scale_max, float scale_max)
float alpha)
{ {
/* Single-thread kernel — only thread 0 of block 0 does anything. */ /* Single-thread kernel — only thread 0 of block 0 does anything. */
if (threadIdx.x != 0 || blockIdx.x != 0) return; if (threadIdx.x != 0 || blockIdx.x != 0) return;
@@ -140,40 +128,64 @@ void hold_cost_scale_update(
const float target_hold_rate = isv[hold_rate_target_idx]; const float target_hold_rate = isv[hold_rate_target_idx];
/* Cold-start sentinel guard: observed hold-rate EMA is at sentinel /* Cold-start sentinel guard: observed hold-rate EMA is at sentinel
* 0.0 (matches the SP13 hold_rate_observer fold-reset sentinel — see * 0.0 — keep the scale slot and every Welford accumulator
* `state_reset_registry.rs::sp13_hold_rate_observed_ema`). Keep the * unchanged. The consumer falls back to scale=1.0 when slot is at
* scale slot unchanged; the consumer falls back to scale=1.0 (bit- * sentinel 0.0 (bit-identical pre-Phase-2 cost magnitude). */
* identical pre-Phase-2 cost magnitude). This mirrors the SP16-P1
* MIN_HOLD_TEMPERATURE_FALLBACK=50.0 sentinel-detect path. */
if (fabsf(observed_hold_rate - sentinel_observed_hold_rate) < EPS_F) { if (fabsf(observed_hold_rate - sentinel_observed_hold_rate) < EPS_F) {
return; return;
} }
/* Hold-rate overrun mapping. The denominator floor of 0.01 keeps /* Hold-rate overrun mapping. */
* the formula numerically stable when the target is configured to
* a tiny value (and protects against accidental sentinel-0 pollution
* on the target slot). target=0.20 (default) → effective denominator
* is the target itself; the floor only matters in pathological
* misconfiguration. Mirrors the SP16-P1 MIN_HOLD_TEMPERATURE producer
* for consistency (same input chain). */
const float overrun = fmaxf(0.0f, observed_hold_rate - target_hold_rate); const float overrun = fmaxf(0.0f, observed_hold_rate - target_hold_rate);
const float target_floor = fmaxf(target_hold_rate, 0.01f); const float target_floor = fmaxf(target_hold_rate, 0.01f);
const float overrun_norm = fminf(1.0f, overrun / target_floor); const float overrun_norm = fminf(1.0f, overrun / target_floor);
float target_scale = scale_min + (scale_max - scale_min) * overrun_norm; float target_scale = scale_min + (scale_max - scale_min) * overrun_norm;
/* Bilateral clamp per `pearl_symmetric_clamp_audit.md`. Defends
* against arithmetic round-off pushing target outside the
* dimensional-safety window. */
target_scale = fmaxf(scale_min, fminf(target_scale, scale_max)); target_scale = fmaxf(scale_min, fminf(target_scale, scale_max));
/* Read Welford accumulator state. */
float sample_count = isv[sample_count_idx];
float target_mean = isv[target_mean_idx];
float target_M2 = isv[target_M2_idx];
float diff_mean = isv[diff_mean_idx];
float diff_M2 = isv[diff_M2_idx];
const float prev_target = isv[prev_target_idx];
/* Update Welford on target_scale. */
sample_count += 1.0f;
const float t_delta = target_scale - target_mean;
target_mean += t_delta / sample_count;
const float t_delta2 = target_scale - target_mean;
target_M2 += t_delta * t_delta2;
/* Update Welford on consecutive diff (need ≥ 2 samples). */
const float diff_count = sample_count - 1.0f;
if (sample_count >= 2.0f) {
const float diff = target_scale - prev_target;
const float d_delta = diff - diff_mean;
diff_mean += d_delta / diff_count;
const float d_delta2 = diff - diff_mean;
diff_M2 += d_delta * d_delta2;
}
/* Compute Wiener-optimal α. */
float alpha;
if (sample_count < WELFORD_MIN_SAMPLE_COUNT_F) {
/* Cold-start REPLACE — keeps producer responsive in epochs
* 1-2 while Welford state stabilises by epoch 3. */
alpha = WELFORD_COLD_START_ALPHA_F;
} else {
const float sample_var = target_M2 / (sample_count - 1.0f);
const float diff_var = (diff_count >= 2.0f)
? diff_M2 / (diff_count - 1.0f)
: diff_M2;
alpha = diff_var / (diff_var + sample_var + EPS_F);
alpha = fmaxf(WELFORD_ALPHA_MIN_F, fminf(alpha, WELFORD_ALPHA_MAX_F));
}
/* Pearl-A first-observation bootstrap. The sentinel is /* Pearl-A first-observation bootstrap. The sentinel is
* SENTINEL_HOLD_COST_SCALE=0.0; any value within EPS of 0 is treated * SENTINEL_HOLD_COST_SCALE=0.0; any value within EPS of 0 is
* as sentinel and REPLACED directly with target_scale. The consumer's * treated as sentinel and REPLACED directly. */
* cold-start fallback (scale=1.0 when isv[461] ≤ 0 or out-of-bounds)
* therefore kicks in for exactly one launch — the producer's first
* valid observation lands a real scale and the consumer reads it
* thereafter. Mirrors the SP14-P1-Producer Pearl-A pattern. */
const float current = isv[scale_idx]; const float current = isv[scale_idx];
float blended; float blended;
if (current <= EPS_F) { if (current <= EPS_F) {
@@ -181,9 +193,16 @@ void hold_cost_scale_update(
} else { } else {
blended = (1.0f - alpha) * current + alpha * target_scale; blended = (1.0f - alpha) * current + alpha * target_scale;
} }
/* Re-clamp post-blend (defensive — handles malformed prior state /* Re-clamp post-blend. */
* outside [scale_min, scale_max]). */
blended = fmaxf(scale_min, fminf(blended, scale_max)); blended = fmaxf(scale_min, fminf(blended, scale_max));
isv[scale_idx] = blended; /* Write back: blended scale + all 6 Welford accumulators + new
* prev_target (= target_scale for the next launch's diff). */
isv[scale_idx] = blended;
isv[sample_count_idx] = sample_count;
isv[target_mean_idx] = target_mean;
isv[target_M2_idx] = target_M2;
isv[diff_mean_idx] = diff_mean;
isv[diff_M2_idx] = diff_M2;
isv[prev_target_idx] = target_scale;
} }

235
crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu
View File

@@ -1,29 +1,18 @@
/* ══════════════════════════════════════════════════════════════════════════ /* ══════════════════════════════════════════════════════════════════════════
* SP16 Phase 1 (revised, 2026-05-08) — adaptive MIN_HOLD_TEMPERATURE producer. * SP16 Phase 1 (revised, 2026-05-08) + T3 (Wiener-optimal α, 2026-05-08) —
* adaptive MIN_HOLD_TEMPERATURE producer.
* *
* SIGNAL SWAP: this kernel was originally driven by * Driver: `realized_hold_rate vs target_hold_rate` overrun. Both values
* `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary * already live on the GPU as ISV slots populated by the existing SP13
* directional accuracy). Post-mortem of train-multi-seed-pfh9n * hold-pricing chain:
* established that slot 373 stays at sentinel 0.5 in Fold 1 (either
* the per-step aux producer doesn't push it off sentinel quickly
* enough or the binary classifier converges within ε of 0.5 in noisy
* environments), causing the kernel's early-return guard at
* `fabsf(short_ema − 0.5) < EPS` to fire on every launch — slot 460
* stayed at sentinel 50 for the entire fold.
*
* The driver is now `realized_hold_rate vs target_hold_rate` overrun.
* Both values already live on the GPU as ISV slots populated by the
* existing SP13 hold-pricing chain:
* *
* - `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` — per-fold EMA of the * - `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` — per-fold EMA of the
* per-step Hold-pick rate `count(action_dir == DIR_HOLD) / B`, * per-step Hold-pick rate, written by `apply_fixed_alpha_ema_kernel`
* written by `apply_fixed_alpha_ema_kernel` chained off * chained off `hold_rate_observer_kernel`.
* `hold_rate_observer_kernel` in the experience collector.
* *
* - `ISV[HOLD_RATE_TARGET_INDEX=381]` — static target (default 0.20), * - `ISV[HOLD_RATE_TARGET_INDEX=381]` — static target (default 0.20).
* constructor-init.
* *
* New mapping: * Mapping (unchanged from T1):
* *
* overrun = max(0, observed − target) * overrun = max(0, observed − target)
* overrun_normalized = clamp(overrun / max(target, 0.01), 0, 1) * overrun_normalized = clamp(overrun / max(target, 0.01), 0, 1)
@@ -41,50 +30,81 @@
* - HIGH T → soft_factor → 0 → no exit penalty (permissive) * - HIGH T → soft_factor → 0 → no exit penalty (permissive)
* - LOW T → soft_factor → 1 → max exit penalty (strict) * - LOW T → soft_factor → 1 → max exit penalty (strict)
* *
* So when the model is over-holding, temperature climbs HIGH (giving an * SP16 T3 — Wiener-optimal adaptive α (2026-05-08):
* exit ramp), and when at/below target it falls LOW (locking in
* commitment). This is a closed-loop control on the actual Hold-rate
* symptom rather than the (currently masked) dir_acc proxy.
* *
* Why this is better than the deleted dir_acc-driven mapping: * The pre-T3 implementation hardcoded `const float alpha = 0.05f` for
* the EMA blend. train-multi-seed-hjzss validation showed this α was
* far too low to converge in 5-epoch smoke (~60 epochs needed from
* cold start). Per `feedback_isv_for_adaptive_bounds.md` the constant
* α was itself an adaptive bound being hardcoded, and per
* `pearl_wiener_optimal_adaptive_alpha`:
* *
* 1. Direct closed-loop control on the actual problem. Hold% is the * α = diff_var / (diff_var + sample_var + ε)
* symptom; hold-rate-overrun directly drives the corrective signal.
* *
* 2. No chained-input-sentinel masking. Hold rates are measured per- * Where `sample_var` = running variance of target_temp and
* epoch from realized actions, never at a "no-data" sentinel after * `diff_var` = running variance of consecutive one-step differences.
* epoch 1. * Welford accumulators stored in ISV slots [468..474):
* *
* 3. Survives fold reset cleanly. Hold-rate measurement starts fresh * - MHT_TARGET_MEAN_INDEX (468) — running mean of target_temp
* in Fold 1 with real data immediately on the first per-step * - MHT_TARGET_M2_INDEX (469) — Welford ΣΔ² (target_temp)
* observation; no need for the dir-acc EMA to climb out of 0.5. * - MHT_DIFF_MEAN_INDEX (470) — running mean of consecutive diffs
* - MHT_DIFF_M2_INDEX (471) — Welford ΣΔ² (diff)
* - MHT_PREV_TARGET_INDEX (472) — previous-epoch target_temp
* - MHT_SAMPLE_COUNT_INDEX (473) — Welford counter
*
* Cold-start path: when `sample_count < 3` (insufficient samples for
* variance estimation), α = 1.0 → REPLACE. This keeps the producer
* responsive in the first 3 epochs of each fold while the Wiener
* formula stabilises. After the third observation, sample_var and
* diff_var are well-defined and α emerges from the data:
*
* - Cold start (target jumps 1.0 → 6.4 → 7.0): diff_var dominates →
* α ≈ 0.6+ → near-bootstrap responsiveness.
* - Steady state: signal stabilises → diff_var drops → α → 0.05 or
* lower → smoothing emerges naturally without a hardcoded constant.
*
* Defensive bounds [WELFORD_ALPHA_MIN=0.01, WELFORD_ALPHA_MAX=0.95]
* applied to the Wiener-derived α before the blend. The Wiener
* formula naturally yields α ∈ [0, 1] but numerical edge cases (both
* variances identically zero within ε) could push α to a denormal
* value or floor at 0 (no learning).
* *
* Algorithm (single-block single-thread — cold path, per-epoch boundary): * Algorithm (single-block single-thread — cold path, per-epoch boundary):
* *
* Phase 1 (single thread): read observed_hold_rate (slot 382), * Phase 1 (read inputs): observed_hold_rate (slot 382),
* target_hold_rate (slot 381), and current temp (slot 460) from ISV. * target_hold_rate (slot 381), current temp (slot 460), and all 6
* Welford accumulators ([468..474)).
* *
* Phase 2 (sentinel guard): if observed_hold_rate is at sentinel 0.0 * Phase 2 (sentinel guard): if observed_hold_rate is at sentinel 0.0
* (no per-step Hold observations have landed yet, e.g. epoch 0 of * (no per-step Hold observations have landed yet), keep ALL slots
* any fold before the first hold_rate_observer launch chains in), * unchanged — the consumer falls back to the kernel-passed scalar.
* keep ISV slot unchanged — sentinel persists; consumer falls back
* to the kernel-passed scalar (seeded at MIN_HOLD_TEMPERATURE_FALLBACK
* =50.0). This is the bootstrap-protection equivalent of the deleted
* dir_acc sentinel guard.
* *
* Phase 3 (compute target temp): * Phase 3 (compute target temp): unchanged from T1.
* *
* overrun = fmaxf(0, observed − target) * Phase 4 (Welford accumulator update):
* overrun_norm = fminf(1, overrun / fmaxf(target, 0.01)) * sample_count += 1
* target_temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × overrun_norm * update target mean/M2 via Welford on target_temp
* bilateral clamp to [TEMP_MIN, TEMP_MAX] * if sample_count ≥ 2: update diff mean/M2 via Welford on
* (target_temp − prev_target).
* *
* Phase 4 (Pearl-A bootstrap + Welford-style EMA): * Phase 5 (Wiener-optimal α):
* if sample_count < 3:
* α = 1.0 (cold-start REPLACE)
* else:
* sample_var = target_M2 / (sample_count − 1)
* diff_var = (diff_count ≥ 2) ? diff_M2 / (diff_count − 1)
* : diff_M2
* α = clamp(diff_var / (diff_var + sample_var + ε),
* WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX)
* *
* Phase 6 (blend):
* if current is at sentinel (within EPS of 50.0) → REPLACE with target. * if current is at sentinel (within EPS of 50.0) → REPLACE with target.
* else → blended = (1 − α) × current + α × target_temp * else → blended = (1 − α) × current + α × target_temp
* re-clamp post-blend. * re-clamp post-blend.
* *
* Phase 7 (write back): blended + all 6 Welford accumulators + new
* prev_target (= target_temp).
*
* Pearls + invariants: * Pearls + invariants:
* *
* - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. * - `feedback_no_atomicadd.md` — single-thread kernel; no reductions.
@@ -96,27 +116,36 @@
* - `pearl_first_observation_bootstrap.md` — first valid hold-rate * - `pearl_first_observation_bootstrap.md` — first valid hold-rate
* observation replaces sentinel directly; no blend. * observation replaces sentinel directly; no blend.
* *
* - `feedback_no_stubs.md` — full body, no return-zero placeholders. * - `pearl_wiener_optimal_adaptive_alpha.md` (NEW T3) — α derived
* from Welford variance accumulators, not hardcoded.
* *
* - `feedback_isv_for_adaptive_bounds.md` — bounds [5, 50] are * - `feedback_isv_for_adaptive_bounds.md` — bounds [5, 50] are
* dimensional safety floors (matches the original schedule range), * dimensional safety floors; α is no longer hardcoded.
* NOT tuning. *
* - `feedback_no_stubs.md` — full body, no return-zero placeholders.
* *
* - `pearl_symmetric_clamp_audit.md` — bilateral * - `pearl_symmetric_clamp_audit.md` — bilateral
* `fmaxf(lo, fminf(x, hi))` clamp on target_temp before writing. * `fmaxf(lo, fminf(x, hi))` clamp on target_temp before writing.
* *
* Args: * Args:
* isv — `[ISV_TOTAL_DIM]` device f32, modified * isv — `[ISV_TOTAL_DIM]` device f32, modified
* in place at index `temp_idx`. * in place at index `temp_idx` and Welford
* slots [target_mean_idx, target_M2_idx,
* diff_mean_idx, diff_M2_idx,
* prev_target_idx, sample_count_idx].
* temp_idx — MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460). * temp_idx — MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460).
* hold_rate_observed_idx — HOLD_RATE_OBSERVED_EMA_INDEX (382). * hold_rate_observed_idx — HOLD_RATE_OBSERVED_EMA_INDEX (382).
* hold_rate_target_idx — HOLD_RATE_TARGET_INDEX (381). * hold_rate_target_idx — HOLD_RATE_TARGET_INDEX (381).
* target_mean_idx — MHT_TARGET_MEAN_INDEX (468).
* target_M2_idx — MHT_TARGET_M2_INDEX (469).
* diff_mean_idx — MHT_DIFF_MEAN_INDEX (470).
* diff_M2_idx — MHT_DIFF_M2_INDEX (471).
* prev_target_idx — MHT_PREV_TARGET_INDEX (472).
* sample_count_idx — MHT_SAMPLE_COUNT_INDEX (473).
* sentinel_temp — SENTINEL_MIN_HOLD_TEMPERATURE (50.0). * sentinel_temp — SENTINEL_MIN_HOLD_TEMPERATURE (50.0).
* sentinel_observed_hold_rate — SENTINEL_HOLD_RATE_OBSERVED (0.0; matches * sentinel_observed_hold_rate — SENTINEL_HOLD_RATE_OBSERVED (0.0).
* the SP13 hold-rate observer cold-start).
* temp_min — MIN_HOLD_TEMPERATURE_MIN (5.0). * temp_min — MIN_HOLD_TEMPERATURE_MIN (5.0).
* temp_max — MIN_HOLD_TEMPERATURE_MAX (50.0). * temp_max — MIN_HOLD_TEMPERATURE_MAX (50.0).
* alpha — MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05).
* *
* Launch: grid=(1, 1, 1), block=(1, 1, 1). * Launch: grid=(1, 1, 1), block=(1, 1, 1).
* No shared memory — single-thread cold-path kernel. * No shared memory — single-thread cold-path kernel.
@@ -125,6 +154,10 @@
#include <cuda_runtime.h> #include <cuda_runtime.h>
#define EPS_F 1e-6f #define EPS_F 1e-6f
#define WELFORD_ALPHA_MIN_F 0.01f
#define WELFORD_ALPHA_MAX_F 0.95f
#define WELFORD_MIN_SAMPLE_COUNT_F 3.0f
#define WELFORD_COLD_START_ALPHA_F 1.0f
extern "C" __global__ extern "C" __global__
void min_hold_temperature_update( void min_hold_temperature_update(
@@ -132,11 +165,16 @@ void min_hold_temperature_update(
int temp_idx, int temp_idx,
int hold_rate_observed_idx, int hold_rate_observed_idx,
int hold_rate_target_idx, int hold_rate_target_idx,
int target_mean_idx,
int target_M2_idx,
int diff_mean_idx,
int diff_M2_idx,
int prev_target_idx,
int sample_count_idx,
float sentinel_temp, float sentinel_temp,
float sentinel_observed_hold_rate, float sentinel_observed_hold_rate,
float temp_min, float temp_min,
float temp_max, float temp_max)
float alpha)
{ {
/* Single-thread kernel — only thread 0 of block 0 does anything. */ /* Single-thread kernel — only thread 0 of block 0 does anything. */
if (threadIdx.x != 0 || blockIdx.x != 0) return; if (threadIdx.x != 0 || blockIdx.x != 0) return;
@@ -145,38 +183,70 @@ void min_hold_temperature_update(
const float target_hold_rate = isv[hold_rate_target_idx]; const float target_hold_rate = isv[hold_rate_target_idx];
/* Cold-start sentinel guard: observed hold-rate EMA is at sentinel /* Cold-start sentinel guard: observed hold-rate EMA is at sentinel
* 0.0 (matches the SP13 hold_rate_observer fold-reset sentinel — see * 0.0 (matches the SP13 hold_rate_observer fold-reset sentinel). Keep
* `state_reset_registry.rs::sp13_hold_rate_observed_ema`). Keep the * the temp slot AND every Welford accumulator unchanged so the
* temp slot unchanged; the consumer will fall back to the kernel- * Welford state machine starts cleanly on the first real
* passed scalar (MIN_HOLD_TEMPERATURE_FALLBACK=50.0) via the * observation of the new fold. */
* sentinel-detect path. Bit-identical to pre-SP16-P1 epoch-0
* behaviour under the deleted dir_acc-driven kernel (which also
* preserved sentinel temp on cold-start). */
if (fabsf(observed_hold_rate - sentinel_observed_hold_rate) < EPS_F) { if (fabsf(observed_hold_rate - sentinel_observed_hold_rate) < EPS_F) {
return; return;
} }
/* Hold-rate overrun mapping. The denominator floor of 0.01 keeps /* Hold-rate overrun mapping. */
* the formula numerically stable when the target is configured to
* a tiny value (and protects against accidental sentinel-0 pollution
* on the target slot). target=0.20 (default) → effective denominator
* is the target itself; the floor only matters in pathological
* misconfiguration. */
const float overrun = fmaxf(0.0f, observed_hold_rate - target_hold_rate); const float overrun = fmaxf(0.0f, observed_hold_rate - target_hold_rate);
const float target_floor = fmaxf(target_hold_rate, 0.01f); const float target_floor = fmaxf(target_hold_rate, 0.01f);
const float overrun_norm = fminf(1.0f, overrun / target_floor); const float overrun_norm = fminf(1.0f, overrun / target_floor);
float target_temp = temp_min + (temp_max - temp_min) * overrun_norm; float target_temp = temp_min + (temp_max - temp_min) * overrun_norm;
/* Bilateral clamp per `pearl_symmetric_clamp_audit.md`. Defends
* against arithmetic round-off pushing target outside the
* dimensional-safety window. */
target_temp = fmaxf(temp_min, fminf(target_temp, temp_max)); target_temp = fmaxf(temp_min, fminf(target_temp, temp_max));
/* Read Welford accumulator state. */
float sample_count = isv[sample_count_idx];
float target_mean = isv[target_mean_idx];
float target_M2 = isv[target_M2_idx];
float diff_mean = isv[diff_mean_idx];
float diff_M2 = isv[diff_M2_idx];
const float prev_target = isv[prev_target_idx];
/* Update Welford on target_temp. */
sample_count += 1.0f;
const float t_delta = target_temp - target_mean;
target_mean += t_delta / sample_count;
const float t_delta2 = target_temp - target_mean;
target_M2 += t_delta * t_delta2;
/* Update Welford on consecutive diff (need ≥ 2 samples). */
const float diff_count = sample_count - 1.0f;
if (sample_count >= 2.0f) {
const float diff = target_temp - prev_target;
const float d_delta = diff - diff_mean;
diff_mean += d_delta / diff_count;
const float d_delta2 = diff - diff_mean;
diff_M2 += d_delta * d_delta2;
}
/* Compute Wiener-optimal α. */
float alpha;
if (sample_count < WELFORD_MIN_SAMPLE_COUNT_F) {
/* Insufficient samples for variance estimation → cold-start
* REPLACE. Keeps producer responsive in epochs 1-2 while the
* Welford state stabilises by epoch 3. */
alpha = WELFORD_COLD_START_ALPHA_F;
} else {
const float sample_var = target_M2 / (sample_count - 1.0f);
const float diff_var = (diff_count >= 2.0f)
? diff_M2 / (diff_count - 1.0f)
: diff_M2;
alpha = diff_var / (diff_var + sample_var + EPS_F);
/* Defensive bounds — Wiener formula naturally yields α ∈ [0, 1]
* but numerical edge cases (both variances identically zero
* within ε) could push α to a denormal or floor at 0. */
alpha = fmaxf(WELFORD_ALPHA_MIN_F, fminf(alpha, WELFORD_ALPHA_MAX_F));
}
/* Pearl-A first-observation bootstrap. The sentinel is /* Pearl-A first-observation bootstrap. The sentinel is
* SENTINEL_MIN_HOLD_TEMPERATURE=50.0 (matches the deleted * SENTINEL_MIN_HOLD_TEMPERATURE=50.0; any value within EPS is
* MIN_HOLD_TEMPERATURE_START); any value within EPS is treated * treated as sentinel and REPLACED directly. Otherwise blend with
* as sentinel and REPLACED directly. Otherwise EMA blend. */ * Wiener-optimal α. */
const float current = isv[temp_idx]; const float current = isv[temp_idx];
float blended; float blended;
if (fabsf(current - sentinel_temp) < EPS_F) { if (fabsf(current - sentinel_temp) < EPS_F) {
@@ -184,9 +254,16 @@ void min_hold_temperature_update(
} else { } else {
blended = (1.0f - alpha) * current + alpha * target_temp; blended = (1.0f - alpha) * current + alpha * target_temp;
} }
/* Re-clamp post-blend (defensive — handles malformed prior state /* Re-clamp post-blend. */
* outside [temp_min, temp_max]). */
blended = fmaxf(temp_min, fminf(blended, temp_max)); blended = fmaxf(temp_min, fminf(blended, temp_max));
isv[temp_idx] = blended; /* Write back: blended temp + all 6 Welford accumulators + new
* prev_target (= target_temp for the next launch's diff). */
isv[temp_idx] = blended;
isv[sample_count_idx] = sample_count;
isv[target_mean_idx] = target_mean;
isv[target_M2_idx] = target_M2;
isv[diff_mean_idx] = diff_mean;
isv[diff_M2_idx] = diff_M2;
isv[prev_target_idx] = target_temp;
} }

177
crates/ml/src/cuda_pipeline/sp14_isv_slots.rs
View File

@@ -426,12 +426,17 @@ pub const SENTINEL_MIN_HOLD_TEMPERATURE: f32 = 50.0;
pub const MIN_HOLD_TEMPERATURE_MIN: f32 = 5.0; pub const MIN_HOLD_TEMPERATURE_MIN: f32 = 5.0;
pub const MIN_HOLD_TEMPERATURE_MAX: f32 = 50.0; pub const MIN_HOLD_TEMPERATURE_MAX: f32 = 50.0;
// EMA blend rate. Moderate — temperature should track the realised // SP16 T3 (2026-05-08): the hardcoded `MIN_HOLD_TEMPERATURE_EMA_ALPHA = 0.05`
// hold-rate EMA (which itself updates per-step) but not perfectly mirror // constant was deleted. Per `feedback_isv_for_adaptive_bounds`, α was
// it (avoid jitter from per-batch hold-rate noise). α=0.05 matches the // itself an adaptive bound being hardcoded; per
// P0-A REWARD_POS_CAP cadence (mid-cadence — slower than per-step // `pearl_wiener_optimal_adaptive_alpha`, α is now derived from Welford
// signals, faster than per-fold beliefs). // accumulators in slots [468..474). The producer kernel
pub const MIN_HOLD_TEMPERATURE_EMA_ALPHA: f32 = 0.05; // `min_hold_temperature_update_kernel.cu` computes
// `α = diff_var / (diff_var + sample_var + ε)` per launch.
// Cold-start path (sample_count < 3) uses α = 1.0 (REPLACE) for the
// first 3 epochs to keep the producer responsive while Welford state
// stabilises. After that the Wiener formula takes over with bounds
// [WELFORD_ALPHA_MIN=0.01, WELFORD_ALPHA_MAX=0.95] for numerical safety.
// SP16 Phase 1 (revised, 2026-05-08) — sentinel for the observed // SP16 Phase 1 (revised, 2026-05-08) — sentinel for the observed
// hold-rate EMA (`ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]`). Matches the // hold-rate EMA (`ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]`). Matches the
@@ -508,17 +513,106 @@ pub const SENTINEL_HOLD_COST_SCALE: f32 = 0.0;
pub const HOLD_COST_SCALE_MIN: f32 = 1.0; pub const HOLD_COST_SCALE_MIN: f32 = 1.0;
pub const HOLD_COST_SCALE_MAX: f32 = 25.0; pub const HOLD_COST_SCALE_MAX: f32 = 25.0;
// EMA blend rate. Mid-cadence (matches SP16-P1 MIN_HOLD_TEMPERATURE // SP16 T3 (2026-05-08): the hardcoded `HOLD_COST_SCALE_EMA_ALPHA = 0.05`
// α=0.05): the scale should track realised Hold-rate overrun on the // constant was deleted. Per `feedback_isv_for_adaptive_bounds` and
// per-epoch cadence (faster than per-fold beliefs, slower than // `pearl_wiener_optimal_adaptive_alpha`, α is now derived from Welford
// per-step signals — same input chain as MIN_HOLD_TEMPERATURE). // accumulators in slots [462..468). See the SP16 T3 comment block at
// Pearl-A bootstrap replaces sentinel directly on first valid // the top of the Welford slots section for the full rationale.
// observation; α=0.05 thereafter.
pub const HOLD_COST_SCALE_EMA_ALPHA: f32 = 0.05;
pub const SP16_P2_HOLD_COST_SCALE_SLOT_BASE: usize = 461; pub const SP16_P2_HOLD_COST_SCALE_SLOT_BASE: usize = 461;
pub const SP16_P2_HOLD_COST_SCALE_SLOT_END: usize = 462; pub const SP16_P2_HOLD_COST_SCALE_SLOT_END: usize = 462;
// ── SP16 T3 (2026-05-08): Wiener-optimal adaptive α — Welford accumulators ─
// Both T1 (MIN_HOLD_TEMPERATURE producer at slot 460) and T2 (HOLD_COST_SCALE
// producer at slot 461) previously hardcoded `alpha = 0.05f` in their EMA
// blend. train-multi-seed-hjzss validation showed the SP16 chain was
// structurally landed but BEHAVIORALLY INERT in 5-epoch smoke — bit-identical
// to the pfh9n baseline through epoch 3 (hold_cost_scale 1.0→1.0→1.27 vs
// target 7.0; min_hold_temp 0.0→5.0→5.50 vs target 16.0). Root cause: with
// α=0.05 the EMA needs ~60 epochs from cold start to converge.
//
// Per `feedback_isv_for_adaptive_bounds.md` the constant α is itself an
// adaptive bound that should not be hardcoded. Per
// `pearl_wiener_optimal_adaptive_alpha`:
//
// α = diff_var / (diff_var + sample_var + ε)
//
// Where `sample_var` = running variance of the target signal and
// `diff_var` = running variance of consecutive one-step differences.
// Cold-start (signal jumps 1.0 → 6.4 → 7.0): diff_var dominates → α ≈ 0.6+
// → near-bootstrap responsiveness in epochs 1-3. Steady-state: signal
// stabilises, diff_var drops, α decays naturally → smoothing emerges
// without a hardcoded constant.
//
// Each producer maintains 6 Welford accumulators (target mean/M2, diff
// mean/M2, prev_target, sample_count). 12 slots total. Pearl-A
// first-observation bootstrap on the prev_blended sentinel (slot 460 sentinel
// 50.0; slot 461 sentinel 0.0). Cold-start fallback when sample_count < 3
// (insufficient samples for variance estimation): α = 1.0 (REPLACE) so the
// producer remains responsive in early epochs.
//
// Slot layout chosen to keep T1 and T2 accumulators contiguous in memory
// (one 12-slot block at 462..474), simplifying the FoldReset registry +
// dispatch arm wiring. Each block of 6 mirrors the Welford state machine
// exactly; the only difference between T1 and T2 is which producer kernel
// reads/writes them (selected by passing the corresponding base index).
//
// HCS = HOLD_COST_SCALE (T2 producer); MHT = MIN_HOLD_TEMPERATURE (T1).
//
// Sentinels: every Welford slot is FoldReset-bootstrap (sentinel 0.0). At
// fold start, sample_count = 0 → cold-start path fires (α = 1.0; first
// valid observation REPLACES). After first observation, sample_count ≥ 1
// and Welford accumulators take over.
pub const HCS_TARGET_MEAN_INDEX: usize = 462; // Welford running mean of target_scale
pub const HCS_TARGET_M2_INDEX: usize = 463; // Welford sum of squared deviations (target)
pub const HCS_DIFF_MEAN_INDEX: usize = 464; // Welford running mean of consecutive diffs
pub const HCS_DIFF_M2_INDEX: usize = 465; // Welford sum of squared deviations (diff)
pub const HCS_PREV_TARGET_INDEX: usize = 466; // previous-epoch target_scale (for diff)
pub const HCS_SAMPLE_COUNT_INDEX: usize = 467; // Welford counter (number of observations)
pub const MHT_TARGET_MEAN_INDEX: usize = 468; // Welford running mean of target_temp
pub const MHT_TARGET_M2_INDEX: usize = 469; // Welford sum of squared deviations (target)
pub const MHT_DIFF_MEAN_INDEX: usize = 470; // Welford running mean of consecutive diffs
pub const MHT_DIFF_M2_INDEX: usize = 471; // Welford sum of squared deviations (diff)
pub const MHT_PREV_TARGET_INDEX: usize = 472; // previous-epoch target_temp (for diff)
pub const MHT_SAMPLE_COUNT_INDEX: usize = 473; // Welford counter (number of observations)
// Sentinel: all 12 Welford slots reset to 0.0 at fold boundary. Combined
// with the cold-start fallback (sample_count < 3 → α = 1.0 REPLACE), this
// means the first observation of each new fold replaces blended directly,
// then the second observation seeds diff statistics, then α adapts from
// the third observation onward — the Wiener formula stabilises by sample 3.
pub const SENTINEL_WELFORD_ZERO: f32 = 0.0;
// Defensive α bounds. The Wiener formula naturally yields α ∈ [0, 1] but
// numerical edge cases (target_M2 = diff_M2 = 0 within ε) could push α
// to a denormal value or floor at 0 (no learning). Clamp to keep the
// blend responsive to genuine signal even under degenerate input
// distributions. Bounds chosen so:
// - WELFORD_ALPHA_MIN=0.01: ensures every observation carries some
// weight, avoiding a permanently-frozen EMA when signals are
// deterministically constant.
// - WELFORD_ALPHA_MAX=0.95: avoids pathological α=1.0 (which is
// equivalent to "no smoothing at all" — only sensible during the
// cold-start path, not steady state).
pub const WELFORD_ALPHA_MIN: f32 = 0.01;
pub const WELFORD_ALPHA_MAX: f32 = 0.95;
// Sample-count threshold for switching from cold-start REPLACE to Wiener
// blend. With < 3 samples we cannot estimate diff variance reliably (need
// at least 2 diffs, which means 3 observations). Below threshold: α=1.0
// (bootstrap). At/above threshold: α from Wiener formula.
pub const WELFORD_MIN_SAMPLE_COUNT: f32 = 3.0;
// Numerical floor — keeps α defined when both variances are identically
// zero (e.g., observed = target exactly, every step). Mirrors the EPS
// floor in `apply_pearls_ad_kernel` and the SP4 wiener buffer producers.
pub const WELFORD_EPS: f32 = 1.0e-6;
pub const SP16_T3_WIENER_SLOT_BASE: usize = 462;
pub const SP16_T3_WIENER_SLOT_END: usize = 474;
#[cfg(test)] #[cfg(test)]
mod tests { mod tests {
use super::*; use super::*;
@@ -754,10 +848,11 @@ mod tests {
// MIN_HOLD_TEMPERATURE_START=50). // MIN_HOLD_TEMPERATURE_START=50).
assert_eq!(MIN_HOLD_TEMPERATURE_MIN, 5.0); assert_eq!(MIN_HOLD_TEMPERATURE_MIN, 5.0);
assert_eq!(MIN_HOLD_TEMPERATURE_MAX, 50.0); assert_eq!(MIN_HOLD_TEMPERATURE_MAX, 50.0);
// EMA α matches mid-cadence (slower than per-step hold-rate EMA, // SP16 T3 (2026-05-08): the hardcoded `MIN_HOLD_TEMPERATURE_EMA_ALPHA
// faster than per-fold beliefs; same as P0-A reward-cap's 0.05 // = 0.05` constant was deleted. α is now Wiener-optimal, derived
// mid-cadence). // from Welford accumulators in slots [468..474) per
assert_eq!(MIN_HOLD_TEMPERATURE_EMA_ALPHA, 0.05); // `pearl_wiener_optimal_adaptive_alpha`. The slot-layout lock
// moved to `sp16_t3_wiener_welford_slot_layout_locked`.
// SP16 Phase 1 (revised) — observed hold-rate EMA sentinel // SP16 Phase 1 (revised) — observed hold-rate EMA sentinel
// matches the SP13 fold-reset registry value (state_reset_ // matches the SP13 fold-reset registry value (state_reset_
// registry.rs `sp13_hold_rate_observed_ema`) so the kernel's // registry.rs `sp13_hold_rate_observed_ema`) so the kernel's
@@ -797,9 +892,10 @@ mod tests {
// with per-bar reward magnitudes ~0.01-0.1). // with per-bar reward magnitudes ~0.01-0.1).
assert_eq!(HOLD_COST_SCALE_MIN, 1.0); assert_eq!(HOLD_COST_SCALE_MIN, 1.0);
assert_eq!(HOLD_COST_SCALE_MAX, 25.0); assert_eq!(HOLD_COST_SCALE_MAX, 25.0);
// EMA α matches SP16-P1 MIN_HOLD_TEMPERATURE mid-cadence (same // SP16 T3 (2026-05-08): the hardcoded `HOLD_COST_SCALE_EMA_ALPHA
// input chain — hold-rate overrun — same target cadence). // = 0.05` constant was deleted. α is now Wiener-optimal, derived
assert_eq!(HOLD_COST_SCALE_EMA_ALPHA, 0.05); // from Welford accumulators in slots [462..468) per
// `pearl_wiener_optimal_adaptive_alpha`.
} }
#[test] #[test]
@@ -812,4 +908,45 @@ mod tests {
SP16_P2_HOLD_COST_SCALE_SLOT_END, ISV_TOTAL_DIM, SP16_P2_HOLD_COST_SCALE_SLOT_END, ISV_TOTAL_DIM,
); );
} }
/// Lock SP16 T3 (2026-05-08) Wiener-optimal Welford accumulator slot
/// layout. 12 contiguous slots [462..474) — 6 per producer (HCS for
/// HOLD_COST_SCALE/T2, MHT for MIN_HOLD_TEMPERATURE/T1). Replaces the
/// hardcoded `α = 0.05f` literal in both producer kernels with a
/// signal-driven Wiener-optimal α per `pearl_wiener_optimal_adaptive_alpha`.
#[test]
fn sp16_t3_wiener_welford_slot_layout_locked() {
assert_eq!(SP16_T3_WIENER_SLOT_BASE, 462);
assert_eq!(SP16_T3_WIENER_SLOT_END, 474);
// HCS block (HOLD_COST_SCALE T2 producer accumulators).
assert_eq!(HCS_TARGET_MEAN_INDEX, 462);
assert_eq!(HCS_TARGET_M2_INDEX, 463);
assert_eq!(HCS_DIFF_MEAN_INDEX, 464);
assert_eq!(HCS_DIFF_M2_INDEX, 465);
assert_eq!(HCS_PREV_TARGET_INDEX, 466);
assert_eq!(HCS_SAMPLE_COUNT_INDEX, 467);
// MHT block (MIN_HOLD_TEMPERATURE T1 producer accumulators).
assert_eq!(MHT_TARGET_MEAN_INDEX, 468);
assert_eq!(MHT_TARGET_M2_INDEX, 469);
assert_eq!(MHT_DIFF_MEAN_INDEX, 470);
assert_eq!(MHT_DIFF_M2_INDEX, 471);
assert_eq!(MHT_PREV_TARGET_INDEX, 472);
assert_eq!(MHT_SAMPLE_COUNT_INDEX, 473);
// Sentinel + bounds — Pearl-A bootstrap on every fold boundary.
assert_eq!(SENTINEL_WELFORD_ZERO, 0.0);
assert_eq!(WELFORD_ALPHA_MIN, 0.01);
assert_eq!(WELFORD_ALPHA_MAX, 0.95);
assert_eq!(WELFORD_MIN_SAMPLE_COUNT, 3.0);
}
#[test]
fn all_sp16_t3_slots_fit_within_isv_total_dim() {
use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM;
assert!(
SP16_T3_WIENER_SLOT_END <= ISV_TOTAL_DIM,
"SP16_T3_WIENER_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for SP16 T3 Wiener-optimal Welford accumulator slots; \
bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).",
SP16_T3_WIENER_SLOT_END, ISV_TOTAL_DIM,
);
}
} }

36
crates/ml/src/cuda_pipeline/state_layout.cuh
View File

@@ -211,7 +211,36 @@
// Cold-start fallback: when slot is at sentinel 0.0 or outside [1, 25], // Cold-start fallback: when slot is at sentinel 0.0 or outside [1, 25],
// consumer uses scale=1.0 (bit-identical pre-Phase-2 cost magnitude). // consumer uses scale=1.0 (bit-identical pre-Phase-2 cost magnitude).
// ──────────────────────────────────────────────────────────────────────────── // ────────────────────────────────────────────────────────────────────────────
#define ISV_HOLD_COST_SCALE_IDX 461 // SP16 Phase 2 adaptive Hold cost scale (sentinel 0.0; bounds [1, 25]; α=0.05) #define ISV_HOLD_COST_SCALE_IDX 461 // SP16 Phase 2 adaptive Hold cost scale (sentinel 0.0; bounds [1, 25]; α=Wiener-optimal per SP16 T3)
// ────────────────────────────────────────────────────────────────────────────
// === SP16 T3 SLOTS (2026-05-08, Wiener-optimal adaptive α) ===
// Mirror of crates/ml/src/cuda_pipeline/sp14_isv_slots.rs slots [462..474).
// Replaces hardcoded `alpha = 0.05f` in both producer kernels (T1 +
// T2) with a signal-driven Wiener-optimal α: α = diff_var / (diff_var
// + sample_var + ε). Each producer maintains 6 Welford accumulators
// (target mean/M2, diff mean/M2, prev_target, sample_count). Per
// `pearl_wiener_optimal_adaptive_alpha`. FoldReset sentinel 0.0 across
// all 12 slots; cold-start fallback (sample_count < 3 → α = 1.0
// REPLACE) keeps the producer responsive in the first 3 epochs of
// each fold while the Wiener formula stabilises.
// ────────────────────────────────────────────────────────────────────────────
#define ISV_HCS_TARGET_MEAN_IDX 462 // Welford mean of target_scale (T2 producer)
#define ISV_HCS_TARGET_M2_IDX 463 // Welford ΣΔ² (target_scale)
#define ISV_HCS_DIFF_MEAN_IDX 464 // Welford mean of consecutive diff (target_scale)
#define ISV_HCS_DIFF_M2_IDX 465 // Welford ΣΔ² (diff)
#define ISV_HCS_PREV_TARGET_IDX 466 // previous-epoch target_scale
#define ISV_HCS_SAMPLE_COUNT_IDX 467 // Welford counter (T2 producer)
#define ISV_MHT_TARGET_MEAN_IDX 468 // Welford mean of target_temp (T1 producer)
#define ISV_MHT_TARGET_M2_IDX 469 // Welford ΣΔ² (target_temp)
#define ISV_MHT_DIFF_MEAN_IDX 470 // Welford mean of consecutive diff (target_temp)
#define ISV_MHT_DIFF_M2_IDX 471 // Welford ΣΔ² (diff)
#define ISV_MHT_PREV_TARGET_IDX 472 // previous-epoch target_temp
#define ISV_MHT_SAMPLE_COUNT_IDX 473 // Welford counter (T1 producer)
#define WELFORD_ALPHA_MIN_F 0.01f
#define WELFORD_ALPHA_MAX_F 0.95f
#define WELFORD_MIN_SAMPLE_COUNT_F 3.0f
#define WELFORD_EPS_F 1.0e-6f
// NOTE: this macro is kept at 383 for the historic SP13-era boundary marker // NOTE: this macro is kept at 383 for the historic SP13-era boundary marker
// and is referenced only in commentary inside .cu/.cuh files. The actual // and is referenced only in commentary inside .cu/.cuh files. The actual
@@ -419,7 +448,10 @@
#define MIN_HOLD_TEMPERATURE_MIN_BOUND 5.0f #define MIN_HOLD_TEMPERATURE_MIN_BOUND 5.0f
#define MIN_HOLD_TEMPERATURE_MAX_BOUND 50.0f #define MIN_HOLD_TEMPERATURE_MAX_BOUND 50.0f
#define MIN_HOLD_TEMPERATURE_DEFAULT 50.0f #define MIN_HOLD_TEMPERATURE_DEFAULT 50.0f
#define MIN_HOLD_TEMPERATURE_EMA_ALPHA 0.05f // SP16 T3 (2026-05-08): the hardcoded `MIN_HOLD_TEMPERATURE_EMA_ALPHA = 0.05f`
// constant was deleted. α is now Wiener-optimal, derived from the Welford
// accumulators in slots [468..474) per `pearl_wiener_optimal_adaptive_alpha`.
// See the SP16 T3 SLOTS section above for the slot indices.
#define DIR_ACC_RANDOM_BASELINE 0.5f #define DIR_ACC_RANDOM_BASELINE 0.5f
// ── Compile-time checks ── // ── Compile-time checks ──

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

@@ -1181,7 +1181,94 @@ impl StateResetRegistry {
RegistryEntry { RegistryEntry {
name: "sp14_audit_4b_min_hold_temperature_adaptive", name: "sp14_audit_4b_min_hold_temperature_adaptive",
category: ResetCategory::FoldReset, category: ResetCategory::FoldReset,
description: "ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] — SP16 Phase 1 (revised, 2026-05-08) adaptive MIN_HOLD_TEMPERATURE (replaces the DELETED epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` formerly at state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch). Producer kernel `min_hold_temperature_update_kernel` reads ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] and ISV[HOLD_RATE_TARGET_INDEX=381] (sp13 hold-pricing chain) and maps the overrun into a [5, 50] temperature via `target_temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp(max(0, observed - target) / max(target, 0.01), 0, 1)`, slow-EMA-blends into slot 460. SIGNAL SWAPPED 2026-05-08 from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] — post-mortem of train-multi-seed-pfh9n showed slot 373 stayed at sentinel 0.5 in Fold 1 (per-step aux producer didn't push it off sentinel quickly enough OR binary classifier converged within ε of 0.5 in noisy environments), so the kernel's early-return-on-sentinel guard pinned slot 460 at 50 for the entire fold. Hold rates are measured per-epoch from realised actions and survive fold reset cleanly. Cold-start sentinel guard preserved: when observed_hold_rate is at SENTINEL_HOLD_RATE_OBSERVED=0.0 (matches sp13_hold_rate_observed_ema fold-reset sentinel), the kernel keeps slot 460 unchanged — consumer falls back to MIN_HOLD_TEMPERATURE_FALLBACK=50.0. FoldReset sentinel SENTINEL_MIN_HOLD_TEMPERATURE=50.0 — Pearl-A first-observation bootstrap (matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start; epoch-0 had T=50 under the old schedule too). α=0.05 mid-cadence EMA. Bounds [5, 50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md` (matches the original schedule range). Per `compute_min_hold_penalty` at trade_physics.cuh:575 (`soft_factor = deficit / (deficit + T)`), HIGH T = forgiving (saturates → 0, no exit penalty), LOW T = sharp (saturates → 1, max exit penalty). Over-holding (overrun > 0) pushes temp HIGH (exit ramp); at-or-below target pushes temp LOW (strict commitment). Closed-loop control on the actual Hold-rate symptom rather than the (sentinel-masked) dir_acc proxy. Consumed by experience_kernels.cu via the kernel-passed `min_hold_temperature` scalar — the training loop seeds that scalar per-epoch via `DQNTrainer::read_min_hold_temperature_from_isv` (cold-start fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when slot is at sentinel).", description: "ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] — SP16 Phase 1 (revised, 2026-05-08) adaptive MIN_HOLD_TEMPERATURE (replaces the DELETED epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` formerly at state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch). Producer kernel `min_hold_temperature_update_kernel` reads ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] and ISV[HOLD_RATE_TARGET_INDEX=381] (sp13 hold-pricing chain) and maps the overrun into a [5, 50] temperature via `target_temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp(max(0, observed - target) / max(target, 0.01), 0, 1)`, slow-EMA-blends into slot 460. SIGNAL SWAPPED 2026-05-08 from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] — post-mortem of train-multi-seed-pfh9n showed slot 373 stayed at sentinel 0.5 in Fold 1 (per-step aux producer didn't push it off sentinel quickly enough OR binary classifier converged within ε of 0.5 in noisy environments), so the kernel's early-return-on-sentinel guard pinned slot 460 at 50 for the entire fold. Hold rates are measured per-epoch from realised actions and survive fold reset cleanly. Cold-start sentinel guard preserved: when observed_hold_rate is at SENTINEL_HOLD_RATE_OBSERVED=0.0 (matches sp13_hold_rate_observed_ema fold-reset sentinel), the kernel keeps slot 460 unchanged — consumer falls back to MIN_HOLD_TEMPERATURE_FALLBACK=50.0. FoldReset sentinel SENTINEL_MIN_HOLD_TEMPERATURE=50.0 — Pearl-A first-observation bootstrap (matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start; epoch-0 had T=50 under the old schedule too). SP16 T3 (2026-05-08): the hardcoded α=0.05 was replaced with Wiener-optimal α from Welford accumulators in slots [468..474). Bounds [5, 50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md` (matches the original schedule range). Per `compute_min_hold_penalty` at trade_physics.cuh:575 (`soft_factor = deficit / (deficit + T)`), HIGH T = forgiving (saturates → 0, no exit penalty), LOW T = sharp (saturates → 1, max exit penalty). Over-holding (overrun > 0) pushes temp HIGH (exit ramp); at-or-below target pushes temp LOW (strict commitment). Closed-loop control on the actual Hold-rate symptom rather than the (sentinel-masked) dir_acc proxy. Consumed by experience_kernels.cu via the kernel-passed `min_hold_temperature` scalar — the training loop seeds that scalar per-epoch via `DQNTrainer::read_min_hold_temperature_from_isv` (cold-start fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when slot is at sentinel).",
},
// ── SP16 T3 (2026-05-08): Wiener-optimal Welford accumulators ────
// 12 ISV slots [462..474) — 6 per producer (HCS for
// HOLD_COST_SCALE/T2 at slot 461, MHT for MIN_HOLD_TEMPERATURE/T1
// at slot 460). Replace the hardcoded `α = 0.05f` literal in both
// producer kernels with a signal-driven Wiener-optimal α per
// `pearl_wiener_optimal_adaptive_alpha`:
//
// α = diff_var / (diff_var + sample_var + ε)
//
// Each block of 6 slots holds:
// - target_mean: Welford running mean of the producer's target
// (target_scale for HCS, target_temp for MHT)
// - target_M2: Welford ΣΔ² (sample variance accumulator)
// - diff_mean: Welford running mean of consecutive one-step diffs
// - diff_M2: Welford ΣΔ² for the diff series
// - prev_target: previous-epoch target value (for next diff calc)
// - sample_count: Welford counter (number of observations)
//
// FoldReset sentinel 0.0 across all 12 slots — combined with the
// cold-start fallback (sample_count < 3 → α = 1.0 REPLACE), this
// means the first observation of each new fold replaces blended
// directly, then the second observation seeds diff statistics, then
// α adapts from the third observation onward. Per
// `feedback_isv_for_adaptive_bounds.md` α was an adaptive bound
// being hardcoded; this lifts the constant into the producer's
// own statistics, eliminating the 60-epoch convergence lag observed
// in train-multi-seed-hjzss validation.
RegistryEntry {
name: "sp16_t3_hcs_target_mean",
category: ResetCategory::FoldReset,
description: "ISV[HCS_TARGET_MEAN_INDEX=462] — SP16 T3 Welford running mean of target_scale for the HOLD_COST_SCALE producer kernel. Sentinel 0.0; first valid observation initialises via Welford increment in `hold_cost_scale_update_kernel.cu`. Combined with the cold-start fallback (sample_count<3 → α=1.0), the first 3 epochs of each fold force REPLACE; from epoch 3 onward Wiener-optimal α emerges from Welford state.",
},
RegistryEntry {
name: "sp16_t3_hcs_target_m2",
category: ResetCategory::FoldReset,
description: "ISV[HCS_TARGET_M2_INDEX=463] — SP16 T3 Welford ΣΔ² (sample-variance accumulator) for target_scale in the HOLD_COST_SCALE producer kernel. Sentinel 0.0; tracks variance of `target_scale` series across each fold. Used to compute `sample_var = M2 / (n - 1)` once `n ≥ 3`.",
},
RegistryEntry {
name: "sp16_t3_hcs_diff_mean",
category: ResetCategory::FoldReset,
description: "ISV[HCS_DIFF_MEAN_INDEX=464] — SP16 T3 Welford running mean of consecutive `target_scale` one-step diffs in the HOLD_COST_SCALE producer kernel. Sentinel 0.0; first diff observation lands when sample_count reaches 2.",
},
RegistryEntry {
name: "sp16_t3_hcs_diff_m2",
category: ResetCategory::FoldReset,
description: "ISV[HCS_DIFF_M2_INDEX=465] — SP16 T3 Welford ΣΔ² for the consecutive-diff series in the HOLD_COST_SCALE producer kernel. Sentinel 0.0; provides `diff_var` half of the Wiener formula α = diff_var / (diff_var + sample_var + ε).",
},
RegistryEntry {
name: "sp16_t3_hcs_prev_target",
category: ResetCategory::FoldReset,
description: "ISV[HCS_PREV_TARGET_INDEX=466] — SP16 T3 previous-epoch `target_scale` for the HOLD_COST_SCALE producer kernel. Sentinel 0.0; written every launch from the freshly-computed `target_scale` so the next launch's diff calculation has a basis.",
},
RegistryEntry {
name: "sp16_t3_hcs_sample_count",
category: ResetCategory::FoldReset,
description: "ISV[HCS_SAMPLE_COUNT_INDEX=467] — SP16 T3 Welford observation counter for the HOLD_COST_SCALE producer kernel. Sentinel 0.0; incremented every launch (excluding the cold-start sentinel-guard early-return). The cold-start fallback (sample_count < WELFORD_MIN_SAMPLE_COUNT=3 → α = 1.0) keeps the producer responsive in epochs 1-2 of each fold while Welford state stabilises.",
},
RegistryEntry {
name: "sp16_t3_mht_target_mean",
category: ResetCategory::FoldReset,
description: "ISV[MHT_TARGET_MEAN_INDEX=468] — SP16 T3 Welford running mean of target_temp for the MIN_HOLD_TEMPERATURE producer kernel. Sentinel 0.0; same algorithm as the HCS counterpart at slot 462 but tracks `target_temp ∈ [5, 50]` instead of `target_scale ∈ [1, 25]`.",
},
RegistryEntry {
name: "sp16_t3_mht_target_m2",
category: ResetCategory::FoldReset,
description: "ISV[MHT_TARGET_M2_INDEX=469] — SP16 T3 Welford ΣΔ² for target_temp in the MIN_HOLD_TEMPERATURE producer kernel. Sentinel 0.0; mirrors slot 463 modulo the producer it serves.",
},
RegistryEntry {
name: "sp16_t3_mht_diff_mean",
category: ResetCategory::FoldReset,
description: "ISV[MHT_DIFF_MEAN_INDEX=470] — SP16 T3 Welford running mean of consecutive `target_temp` diffs in the MIN_HOLD_TEMPERATURE producer kernel. Sentinel 0.0; mirrors slot 464.",
},
RegistryEntry {
name: "sp16_t3_mht_diff_m2",
category: ResetCategory::FoldReset,
description: "ISV[MHT_DIFF_M2_INDEX=471] — SP16 T3 Welford ΣΔ² for the consecutive-diff series in the MIN_HOLD_TEMPERATURE producer kernel. Sentinel 0.0; mirrors slot 465.",
},
RegistryEntry {
name: "sp16_t3_mht_prev_target",
category: ResetCategory::FoldReset,
description: "ISV[MHT_PREV_TARGET_INDEX=472] — SP16 T3 previous-epoch `target_temp` for the MIN_HOLD_TEMPERATURE producer kernel. Sentinel 0.0; mirrors slot 466.",
},
RegistryEntry {
name: "sp16_t3_mht_sample_count",
category: ResetCategory::FoldReset,
description: "ISV[MHT_SAMPLE_COUNT_INDEX=473] — SP16 T3 Welford observation counter for the MIN_HOLD_TEMPERATURE producer kernel. Sentinel 0.0; mirrors slot 467. Cold-start fallback (sample_count < 3 → α = 1.0) ensures the first 3 epochs of each fold force REPLACE so the Welford-derived α takes over from epoch 3 onward without a 60-epoch convergence lag.",
}, },
// ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ──────────── // ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ────────────
// Two ISV slots [407, 408]: // Two ISV slots [407, 408]:

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

@@ -578,7 +578,12 @@ impl DQNTrainer {
use crate::cuda_pipeline::sp13_isv_slots::{ use crate::cuda_pipeline::sp13_isv_slots::{
HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX,
}; };
use crate::cuda_pipeline::sp14_isv_slots::MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX; use crate::cuda_pipeline::sp14_isv_slots::{
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX,
MHT_TARGET_M2_INDEX, MHT_DIFF_M2_INDEX, MHT_SAMPLE_COUNT_INDEX,
WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX, WELFORD_MIN_SAMPLE_COUNT,
WELFORD_EPS,
};
let trainer = fused.trainer(); let trainer = fused.trainer();
let obs = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX); let obs = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX);
let tgt = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX); let tgt = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX);
@@ -587,9 +592,27 @@ impl DQNTrainer {
let denom = tgt.max(0.01); let denom = tgt.max(0.01);
let overrun_norm = (overrun / denom).clamp(0.0, 1.0); let overrun_norm = (overrun / denom).clamp(0.0, 1.0);
let target_temp = 5.0_f32 + (50.0_f32 - 5.0_f32) * overrun_norm; let target_temp = 5.0_f32 + (50.0_f32 - 5.0_f32) * overrun_norm;
// SP16 T3 (2026-05-08): recover the Wiener-optimal α
// from the on-device Welford state for direct
// observation. Mirrors the kernel's algorithm so
// smoke-test trajectories can verify α adapts (e.g.
// α ≈ 0.6+ during epoch 1-3, decays toward
// WELFORD_ALPHA_MIN in steady state).
let n = trainer.read_isv_signal_at(MHT_SAMPLE_COUNT_INDEX);
let target_m2 = trainer.read_isv_signal_at(MHT_TARGET_M2_INDEX);
let diff_m2 = trainer.read_isv_signal_at(MHT_DIFF_M2_INDEX);
let alpha = if n < WELFORD_MIN_SAMPLE_COUNT {
1.0_f32
} else {
let sample_var = target_m2 / (n - 1.0).max(1.0);
let diff_count = (n - 1.0).max(0.0);
let diff_var = if diff_count >= 2.0 { diff_m2 / (diff_count - 1.0) } else { diff_m2 };
(diff_var / (diff_var + sample_var + WELFORD_EPS))
.clamp(WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX)
};
tracing::info!( tracing::info!(
"HEALTH_DIAG[{}]: min_hold_temp_diag obs_hold_rate={:.4} target_hold_rate={:.4} overrun_norm={:.4} target_temp={:.4} blended_temp={:.4}", "HEALTH_DIAG[{}]: min_hold_temp_diag obs_hold_rate={:.4} target_hold_rate={:.4} overrun_norm={:.4} target_temp={:.4} blended_temp={:.4} alpha={:.4} sample_count={}",
epoch, obs, tgt, overrun_norm, target_temp, temp, epoch, obs, tgt, overrun_norm, target_temp, temp, alpha, n as u32,
); );
} }
// SP16 Phase 2 (2026-05-08): hold_cost_scale_diag — // SP16 Phase 2 (2026-05-08): hold_cost_scale_diag —
@@ -608,7 +631,12 @@ impl DQNTrainer {
use crate::cuda_pipeline::sp13_isv_slots::{ use crate::cuda_pipeline::sp13_isv_slots::{
HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX,
}; };
use crate::cuda_pipeline::sp14_isv_slots::HOLD_COST_SCALE_INDEX; use crate::cuda_pipeline::sp14_isv_slots::{
HOLD_COST_SCALE_INDEX,
HCS_TARGET_M2_INDEX, HCS_DIFF_M2_INDEX, HCS_SAMPLE_COUNT_INDEX,
WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX, WELFORD_MIN_SAMPLE_COUNT,
WELFORD_EPS,
};
let trainer = fused.trainer(); let trainer = fused.trainer();
let obs = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX); let obs = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX);
let tgt = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX); let tgt = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX);
@@ -617,9 +645,24 @@ impl DQNTrainer {
let denom = tgt.max(0.01); let denom = tgt.max(0.01);
let overrun_norm = (overrun / denom).clamp(0.0, 1.0); let overrun_norm = (overrun / denom).clamp(0.0, 1.0);
let target_scale = 1.0_f32 + (25.0_f32 - 1.0_f32) * overrun_norm; let target_scale = 1.0_f32 + (25.0_f32 - 1.0_f32) * overrun_norm;
// SP16 T3 (2026-05-08): recover the Wiener-optimal α
// from the on-device Welford state for direct
// observation (mirrors the kernel's formula).
let n = trainer.read_isv_signal_at(HCS_SAMPLE_COUNT_INDEX);
let target_m2 = trainer.read_isv_signal_at(HCS_TARGET_M2_INDEX);
let diff_m2 = trainer.read_isv_signal_at(HCS_DIFF_M2_INDEX);
let alpha = if n < WELFORD_MIN_SAMPLE_COUNT {
1.0_f32
} else {
let sample_var = target_m2 / (n - 1.0).max(1.0);
let diff_count = (n - 1.0).max(0.0);
let diff_var = if diff_count >= 2.0 { diff_m2 / (diff_count - 1.0) } else { diff_m2 };
(diff_var / (diff_var + sample_var + WELFORD_EPS))
.clamp(WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX)
};
tracing::info!( tracing::info!(
"HEALTH_DIAG[{}]: hold_cost_scale_diag obs_hold_rate={:.4} target_hold_rate={:.4} overrun_norm={:.4} target_scale={:.4} blended_scale={:.4}", "HEALTH_DIAG[{}]: hold_cost_scale_diag obs_hold_rate={:.4} target_hold_rate={:.4} overrun_norm={:.4} target_scale={:.4} blended_scale={:.4} alpha={:.4} sample_count={}",
epoch, obs, tgt, overrun_norm, target_scale, scale, epoch, obs, tgt, overrun_norm, target_scale, scale, alpha, n as u32,
); );
} }
} }
@@ -8567,6 +8610,116 @@ impl DQNTrainer {
); );
} }
} }
// ── SP16 T3 (2026-05-08): Wiener-optimal Welford accumulators ─
// 12 dispatch arms (6 per producer). Every slot resets to the
// canonical Welford zero on every fold boundary — combined with
// the cold-start fallback (sample_count < 3 → α = 1.0 REPLACE),
// this means the first observation of each new fold replaces
// blended directly, then the second seeds diff statistics,
// then α adapts via Wiener formula from the third observation
// onward. Per `pearl_wiener_optimal_adaptive_alpha` and
// `feedback_isv_for_adaptive_bounds`.
//
// The HCS_* slots (462..468) feed the HOLD_COST_SCALE producer
// kernel; MHT_* slots (468..474) feed the MIN_HOLD_TEMPERATURE
// producer. Both producers share the same Welford-state
// algorithm — only the slot indices and bound constants differ.
"sp16_t3_hcs_target_mean" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
HCS_TARGET_MEAN_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(HCS_TARGET_MEAN_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_hcs_target_m2" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
HCS_TARGET_M2_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(HCS_TARGET_M2_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_hcs_diff_mean" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
HCS_DIFF_MEAN_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(HCS_DIFF_MEAN_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_hcs_diff_m2" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
HCS_DIFF_M2_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(HCS_DIFF_M2_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_hcs_prev_target" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
HCS_PREV_TARGET_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(HCS_PREV_TARGET_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_hcs_sample_count" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
HCS_SAMPLE_COUNT_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(HCS_SAMPLE_COUNT_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_mht_target_mean" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
MHT_TARGET_MEAN_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(MHT_TARGET_MEAN_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_mht_target_m2" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
MHT_TARGET_M2_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(MHT_TARGET_M2_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_mht_diff_mean" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
MHT_DIFF_MEAN_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(MHT_DIFF_MEAN_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_mht_diff_m2" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
MHT_DIFF_M2_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(MHT_DIFF_M2_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_mht_prev_target" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
MHT_PREV_TARGET_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(MHT_PREV_TARGET_INDEX, SENTINEL_WELFORD_ZERO);
}
}
"sp16_t3_mht_sample_count" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp14_isv_slots::{
MHT_SAMPLE_COUNT_INDEX, SENTINEL_WELFORD_ZERO,
};
fused.trainer().write_isv_signal_at(MHT_SAMPLE_COUNT_INDEX, SENTINEL_WELFORD_ZERO);
}
}
// SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots.
// OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT
// a stateful EMA) — rewrite the constructor's value at fold // a stateful EMA) — rewrite the constructor's value at fold

557
crates/ml/tests/sp14_oracle_tests.rs
View File

@@ -1771,9 +1771,12 @@ mod sp16_phase1_min_hold_temperature_gpu {
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
use ml::cuda_pipeline::sp14_isv_slots::{ use ml::cuda_pipeline::sp14_isv_slots::{
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX,
MIN_HOLD_TEMPERATURE_EMA_ALPHA, MIN_HOLD_TEMPERATURE_MAX, MIN_HOLD_TEMPERATURE_MAX,
MIN_HOLD_TEMPERATURE_MIN, SENTINEL_HOLD_RATE_OBSERVED, MIN_HOLD_TEMPERATURE_MIN, SENTINEL_HOLD_RATE_OBSERVED,
SENTINEL_MIN_HOLD_TEMPERATURE, SENTINEL_MIN_HOLD_TEMPERATURE,
MHT_TARGET_MEAN_INDEX, MHT_TARGET_M2_INDEX,
MHT_DIFF_MEAN_INDEX, MHT_DIFF_M2_INDEX,
MHT_PREV_TARGET_INDEX, MHT_SAMPLE_COUNT_INDEX,
}; };
use ml::cuda_pipeline::sp13_isv_slots::{ use ml::cuda_pipeline::sp13_isv_slots::{
AUX_DIR_ACC_SHORT_EMA_INDEX, HOLD_RATE_OBSERVED_EMA_INDEX, AUX_DIR_ACC_SHORT_EMA_INDEX, HOLD_RATE_OBSERVED_EMA_INDEX,
@@ -1798,6 +1801,10 @@ mod sp16_phase1_min_hold_temperature_gpu {
.expect("load min_hold_temperature_update function") .expect("load min_hold_temperature_update function")
} }
/// SP16 T3 (2026-05-08): kernel takes Welford accumulator slot indices
/// instead of a hardcoded α. The hardcoded α=0.05 is gone — α is
/// computed inside the kernel from the Welford state per
/// `pearl_wiener_optimal_adaptive_alpha`.
#[allow(clippy::too_many_arguments)] #[allow(clippy::too_many_arguments)]
fn launch_temp( fn launch_temp(
stream: &Arc<CudaStream>, stream: &Arc<CudaStream>,
@@ -1806,11 +1813,16 @@ mod sp16_phase1_min_hold_temperature_gpu {
temp_idx: i32, temp_idx: i32,
hold_rate_observed_idx: i32, hold_rate_observed_idx: i32,
hold_rate_target_idx: i32, hold_rate_target_idx: i32,
target_mean_idx: i32,
target_m2_idx: i32,
diff_mean_idx: i32,
diff_m2_idx: i32,
prev_target_idx: i32,
sample_count_idx: i32,
sentinel_temp: f32, sentinel_temp: f32,
sentinel_observed_hold_rate: f32, sentinel_observed_hold_rate: f32,
temp_min: f32, temp_min: f32,
temp_max: f32, temp_max: f32,
alpha: f32,
) { ) {
unsafe { unsafe {
stream stream
@@ -1819,11 +1831,16 @@ mod sp16_phase1_min_hold_temperature_gpu {
.arg(&temp_idx) .arg(&temp_idx)
.arg(&hold_rate_observed_idx) .arg(&hold_rate_observed_idx)
.arg(&hold_rate_target_idx) .arg(&hold_rate_target_idx)
.arg(&target_mean_idx)
.arg(&target_m2_idx)
.arg(&diff_mean_idx)
.arg(&diff_m2_idx)
.arg(&prev_target_idx)
.arg(&sample_count_idx)
.arg(&sentinel_temp) .arg(&sentinel_temp)
.arg(&sentinel_observed_hold_rate) .arg(&sentinel_observed_hold_rate)
.arg(&temp_min) .arg(&temp_min)
.arg(&temp_max) .arg(&temp_max)
.arg(&alpha)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (1, 1, 1), block_dim: (1, 1, 1),
@@ -1835,8 +1852,7 @@ mod sp16_phase1_min_hold_temperature_gpu {
} }
/// Standard launch helper: takes the canonical SP14/SP13 constants, /// Standard launch helper: takes the canonical SP14/SP13 constants,
/// avoiding the verbose 9-arg site at every test. Used by all /// avoiding the verbose argument list at every test.
/// behavioral tests in this module.
fn launch_canonical( fn launch_canonical(
stream: &Arc<CudaStream>, stream: &Arc<CudaStream>,
kernel: &CudaFunction, kernel: &CudaFunction,
@@ -1847,11 +1863,16 @@ mod sp16_phase1_min_hold_temperature_gpu {
MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32,
HOLD_RATE_OBSERVED_EMA_INDEX as i32, HOLD_RATE_OBSERVED_EMA_INDEX as i32,
HOLD_RATE_TARGET_INDEX as i32, HOLD_RATE_TARGET_INDEX as i32,
MHT_TARGET_MEAN_INDEX as i32,
MHT_TARGET_M2_INDEX as i32,
MHT_DIFF_MEAN_INDEX as i32,
MHT_DIFF_M2_INDEX as i32,
MHT_PREV_TARGET_INDEX as i32,
MHT_SAMPLE_COUNT_INDEX as i32,
SENTINEL_MIN_HOLD_TEMPERATURE, SENTINEL_MIN_HOLD_TEMPERATURE,
SENTINEL_HOLD_RATE_OBSERVED, SENTINEL_HOLD_RATE_OBSERVED,
MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MIN,
MIN_HOLD_TEMPERATURE_MAX, MIN_HOLD_TEMPERATURE_MAX,
MIN_HOLD_TEMPERATURE_EMA_ALPHA,
); );
} }
@@ -2033,14 +2054,16 @@ mod sp16_phase1_min_hold_temperature_gpu {
// pre-swap would have preserved slot 460 (kernel returned // pre-swap would have preserved slot 460 (kernel returned
// without writing). Post-swap the kernel writes target_temp=5 // without writing). Post-swap the kernel writes target_temp=5
// because under-target overrun_norm=0. // because under-target overrun_norm=0.
//
// SP16 T3 (2026-05-08): with Wiener-optimal α, cold-start
// (sample_count=0 → 1 < 3) uses α=1.0 (REPLACE). The blend
// path is exercised, but the formula `(1-α)×current + α×target`
// collapses to `target` when α=1. So slot 460 lands at
// target_temp=5 directly (NOT the OLD 28.75 result).
let mut isv2 = vec![0.0_f32; ISV_DIM]; let mut isv2 = vec![0.0_f32; ISV_DIM];
isv2[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.5; // OLD sentinel isv2[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.5; // OLD sentinel
isv2[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.10; // under isv2[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.10; // under
isv2[HOLD_RATE_TARGET_INDEX] = 0.20; isv2[HOLD_RATE_TARGET_INDEX] = 0.20;
// Seed slot 460 at NON-sentinel value so we can see the EMA
// blend write. Pre-swap: kernel would have returned (slot
// pre-write value preserved) because slot 373 is at sentinel.
// Post-swap: kernel computes target=5 and EMA-blends.
const SEED_TEMP: f32 = 30.0; const SEED_TEMP: f32 = 30.0;
isv2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP; isv2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP;
let isv_buf2 = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv2"); let isv_buf2 = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv2");
@@ -2050,14 +2073,14 @@ mod sp16_phase1_min_hold_temperature_gpu {
let temp2 = result2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; let temp2 = result2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];
// overrun = 0 → target_temp = 5. // overrun = 0 → target_temp = 5.
// EMA blend: (1−0.05)×30 + 0.05×5 = 28.75. NOT 30 (the OLD // Cold-start path (sample_count<3) → α=1.0 → blended ≈ target_temp = 5.
// kernel's no-write outcome). // Slot moved off SEED_TEMP=30 (the OLD kernel's no-write outcome).
let expected2 = (1.0_f32 - MIN_HOLD_TEMPERATURE_EMA_ALPHA) * SEED_TEMP
+ MIN_HOLD_TEMPERATURE_EMA_ALPHA * MIN_HOLD_TEMPERATURE_MIN;
assert!( assert!(
(temp2 - expected2).abs() < 1e-4, (temp2 - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4,
"Slot 373 at OLD sentinel must NOT preserve slot 460; expected EMA blend {expected2:.5}, got {temp2:.5} \ "Slot 373 at OLD sentinel must NOT preserve slot 460; under-target → \
(if {temp2:.5} ≈ {SEED_TEMP} the OLD kernel's slot-373-guard is still firing — signal swap not live)" cold-start REPLACE with target_temp = {}, got {temp2:.5} \
(if {temp2:.5} ≈ {SEED_TEMP} the OLD kernel's slot-373-guard is still firing — signal swap not live)",
MIN_HOLD_TEMPERATURE_MIN
); );
assert!( assert!(
(temp2 - SEED_TEMP).abs() > 0.5, (temp2 - SEED_TEMP).abs() > 0.5,
@@ -2065,14 +2088,21 @@ mod sp16_phase1_min_hold_temperature_gpu {
); );
} }
/// Helper test (preserved from pre-swap): mid-cadence EMA blend /// SP16 T3 (2026-05-08) — Wiener-optimal cold-start path.
/// after bootstrap. Seeds slot 460 NOT at sentinel and verifies
/// the (1−α)×current + α×target_temp formula. Updated post-swap
/// to use hold-rate overrun as the driver.
/// ///
/// Setup: temp at 30.0 (NOT sentinel — past bootstrap). /// Pre-T3 this test verified `(1−0.05)×30 + 0.05×50 = 31.0`. Under
/// Wiener-optimal α the kernel computes α from Welford state. With
/// freshly-zeroed Welford accumulators (the canonical fold-reset
/// state), the first launch sees `sample_count=0 → 1 < 3` and
/// takes the cold-start path α=1.0 → blended = target_temp = 50.0.
///
/// Setup: temp at 30.0 (NOT sentinel — past Pearl-A bootstrap).
/// observed=0.40, target=0.20 → overrun_norm=1.0 → target_temp=50. /// observed=0.40, target=0.20 → overrun_norm=1.0 → target_temp=50.
/// Blended = (1 − 0.05) × 30.0 + 0.05 × 50.0 = 28.5 + 2.5 = 31.0. /// Welford state at fold-reset zero. Cold-start path → α=1.0 →
/// blend collapses to target_temp = MIN_HOLD_TEMPERATURE_MAX = 50.
/// After 3 launches with this constant target, `sample_count = 3`
/// reaches the Wiener threshold; with constant target_temp=50 every
/// epoch, diff_var is exactly 0 and α floors at WELFORD_ALPHA_MIN.
#[test] #[test]
#[ignore = "requires GPU"] #[ignore = "requires GPU"]
fn sp16_phase1_min_hold_temp_mid_cadence_ema_after_bootstrap() { fn sp16_phase1_min_hold_temp_mid_cadence_ema_after_bootstrap() {
@@ -2094,13 +2124,18 @@ mod sp16_phase1_min_hold_temperature_gpu {
let result = isv_buf.read_all(); let result = isv_buf.read_all();
let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX];
// EMA blend: (1 − 0.05) × 30.0 + 0.05 × 50.0 = 31.0. // Cold-start path (sample_count=1 < 3) → α=1.0 → blend
let target_temp = MIN_HOLD_TEMPERATURE_MAX; // collapses to target_temp = TEMP_MAX = 50.
let expected = (1.0_f32 - MIN_HOLD_TEMPERATURE_EMA_ALPHA) * SEED_TEMP let expected = MIN_HOLD_TEMPERATURE_MAX;
+ MIN_HOLD_TEMPERATURE_EMA_ALPHA * target_temp;
assert!( assert!(
(temp - expected).abs() < 1e-4, (temp - expected).abs() < 1e-4,
"Mid-cadence EMA blend: expected ≈ {expected:.5}, got {temp:.5}" "Wiener cold-start REPLACE: expected ≈ {expected:.5}, got {temp:.5}"
);
// Welford counter must have been incremented to 1.
assert!(
(result[MHT_SAMPLE_COUNT_INDEX] - 1.0).abs() < 1e-4,
"Welford sample_count must be 1 after first launch, got {}",
result[MHT_SAMPLE_COUNT_INDEX]
); );
} }
@@ -2168,8 +2203,11 @@ mod sp16_phase2_hold_cost_scale_gpu {
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
use ml::cuda_pipeline::sp14_isv_slots::{ use ml::cuda_pipeline::sp14_isv_slots::{
HOLD_COST_SCALE_EMA_ALPHA, HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX, HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX,
HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED, HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED,
HCS_TARGET_MEAN_INDEX, HCS_TARGET_M2_INDEX,
HCS_DIFF_MEAN_INDEX, HCS_DIFF_M2_INDEX,
HCS_PREV_TARGET_INDEX, HCS_SAMPLE_COUNT_INDEX,
}; };
use ml::cuda_pipeline::sp13_isv_slots::{ use ml::cuda_pipeline::sp13_isv_slots::{
HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX,
@@ -2193,7 +2231,9 @@ mod sp16_phase2_hold_cost_scale_gpu {
.expect("load hold_cost_scale_update function") .expect("load hold_cost_scale_update function")
} }
/// Standard launch helper using the canonical SP16-P2 constants. /// Standard launch helper using the canonical SP16-P2 + T3 constants.
/// SP16 T3 (2026-05-08): kernel takes Welford accumulator slot
/// indices instead of a hardcoded α.
fn launch_canonical( fn launch_canonical(
stream: &Arc<CudaStream>, stream: &Arc<CudaStream>,
kernel: &CudaFunction, kernel: &CudaFunction,
@@ -2206,11 +2246,16 @@ mod sp16_phase2_hold_cost_scale_gpu {
.arg(&(HOLD_COST_SCALE_INDEX as i32)) .arg(&(HOLD_COST_SCALE_INDEX as i32))
.arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32)) .arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32))
.arg(&(HOLD_RATE_TARGET_INDEX as i32)) .arg(&(HOLD_RATE_TARGET_INDEX as i32))
.arg(&(HCS_TARGET_MEAN_INDEX as i32))
.arg(&(HCS_TARGET_M2_INDEX as i32))
.arg(&(HCS_DIFF_MEAN_INDEX as i32))
.arg(&(HCS_DIFF_M2_INDEX as i32))
.arg(&(HCS_PREV_TARGET_INDEX as i32))
.arg(&(HCS_SAMPLE_COUNT_INDEX as i32))
.arg(&SENTINEL_HOLD_COST_SCALE) .arg(&SENTINEL_HOLD_COST_SCALE)
.arg(&SENTINEL_HOLD_RATE_OBSERVED) .arg(&SENTINEL_HOLD_RATE_OBSERVED)
.arg(&HOLD_COST_SCALE_MIN) .arg(&HOLD_COST_SCALE_MIN)
.arg(&HOLD_COST_SCALE_MAX) .arg(&HOLD_COST_SCALE_MAX)
.arg(&HOLD_COST_SCALE_EMA_ALPHA)
.launch(LaunchConfig { .launch(LaunchConfig {
grid_dim: (1, 1, 1), grid_dim: (1, 1, 1),
block_dim: (1, 1, 1), block_dim: (1, 1, 1),
@@ -2420,6 +2465,456 @@ mod sp16_phase2_hold_cost_scale_gpu {
} }
} }
// ═══════════════════════════════════════════════════════════════════════════
// SP16 T3 (2026-05-08) — Wiener-optimal adaptive α via Welford accumulators.
//
// Per train-multi-seed-hjzss validation: SP16 T1+T2 chain was structurally
// landed but BEHAVIORALLY INERT in 5-epoch smoke (bit-identical to pfh9n
// baseline through epoch 3). Root cause: hardcoded `alpha = 0.05f` in both
// producer kernels needed ~60 epochs from cold start to converge. Per
// `feedback_isv_for_adaptive_bounds.md` α was itself an adaptive bound being
// hardcoded; per `pearl_wiener_optimal_adaptive_alpha`:
//
// α = diff_var / (diff_var + sample_var + ε)
//
// These tests verify the Wiener-optimal α implementation in
// `hold_cost_scale_update_kernel.cu` (drives SP16 T2 / slot 461). The
// MIN_HOLD_TEMPERATURE producer (T1 / slot 460) shares the same algorithm
// — the T2 tests are sufficient as the smoking gun; T1 mirrors verbatim
// modulo slot indices.
//
// Verifies:
// 1. α high during signal jumps (cold-start path responds aggressively).
// 2. α low in steady state (Welford diff_var → 0 → α decays).
// 3. Pearl-A first-observation bootstrap fires regardless of computed α.
// 4. α stays in (0, 1) over 50 synthetic epochs (bounded).
// 5. No 0.05f hardcoded literal remains in the kernel source (regression-
// locked at static-analysis time, not runtime).
//
// All GPU tests are #[ignore = "requires GPU"]; gated under
// #[cfg(feature = "cuda")]. Test 5 is a pure host-side string scan.
// ═══════════════════════════════════════════════════════════════════════════
#[cfg(feature = "cuda")]
#[allow(unsafe_code)]
mod sp16_phase3_wiener_alpha_gpu {
use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
use ml::cuda_pipeline::sp14_isv_slots::{
HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX,
HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED,
HCS_TARGET_MEAN_INDEX, HCS_TARGET_M2_INDEX,
HCS_DIFF_MEAN_INDEX, HCS_DIFF_M2_INDEX,
HCS_PREV_TARGET_INDEX, HCS_SAMPLE_COUNT_INDEX,
WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX,
};
use ml::cuda_pipeline::sp13_isv_slots::{
HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX,
};
const HOLD_COST_SCALE_UPDATE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/hold_cost_scale_update_kernel.cubin"));
fn make_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec())
.expect("load hold_cost_scale_update_kernel cubin");
module
.load_function("hold_cost_scale_update")
.expect("load hold_cost_scale_update function")
}
fn launch_canonical(stream: &Arc<CudaStream>, kernel: &CudaFunction, isv_ptr: u64) {
unsafe {
stream
.launch_builder(kernel)
.arg(&isv_ptr)
.arg(&(HOLD_COST_SCALE_INDEX as i32))
.arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32))
.arg(&(HOLD_RATE_TARGET_INDEX as i32))
.arg(&(HCS_TARGET_MEAN_INDEX as i32))
.arg(&(HCS_TARGET_M2_INDEX as i32))
.arg(&(HCS_DIFF_MEAN_INDEX as i32))
.arg(&(HCS_DIFF_M2_INDEX as i32))
.arg(&(HCS_PREV_TARGET_INDEX as i32))
.arg(&(HCS_SAMPLE_COUNT_INDEX as i32))
.arg(&SENTINEL_HOLD_COST_SCALE)
.arg(&SENTINEL_HOLD_RATE_OBSERVED)
.arg(&HOLD_COST_SCALE_MIN)
.arg(&HOLD_COST_SCALE_MAX)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch hold_cost_scale_update");
}
stream.synchronize().expect("sync after hold_cost_scale_update");
}
/// Recover the actual α used by the kernel on the latest launch. The
/// kernel doesn't directly expose α, but it can be reverse-engineered
/// from the blend equation:
///
/// blended = (1 − α) × prior_scale + α × target_scale (when prior > 0)
/// ⇒ α = (blended − prior_scale) / (target_scale − prior_scale)
///
/// Pre-launch fix: write `prior_scale` to slot 461 with a non-sentinel
/// value so the Pearl-A bootstrap branch isn't taken; capture
/// (prior_scale, target_scale) and compare to (blended) post-launch.
/// Caller MUST set `target_scale ≠ prior_scale` to keep denominator
/// well-defined; bilateral clamp post-blend may corrupt the
/// recovered α near boundaries — choose seed/observed values that
/// keep the blend well-inside [SCALE_MIN, SCALE_MAX].
fn recover_alpha(blended: f32, prior_scale: f32, target_scale: f32) -> f32 {
(blended - prior_scale) / (target_scale - prior_scale)
}
/// Compute target_scale from observed/target without launching the
/// kernel — mirrors the kernel's overrun-norm formula exactly.
fn compute_target_scale(observed: f32, target: f32) -> f32 {
let overrun = (observed - target).max(0.0);
let denom = target.max(0.01);
let overrun_norm = (overrun / denom).clamp(0.0, 1.0);
(HOLD_COST_SCALE_MIN
+ (HOLD_COST_SCALE_MAX - HOLD_COST_SCALE_MIN) * overrun_norm)
.clamp(HOLD_COST_SCALE_MIN, HOLD_COST_SCALE_MAX)
}
/// Run a synthetic per-epoch sequence by setting (observed, prev_blended)
/// per step and capturing the post-launch slot 461.
/// Returns the trajectory of `(blended, recovered_alpha)` per epoch.
fn run_synthetic_sequence(
stream: &Arc<CudaStream>,
kernel: &CudaFunction,
observations: &[f32],
target: f32,
seed_blended: f32,
) -> Vec<(f32, f32)> {
const ISV_DIM: usize = 1024;
let mut isv = vec![0.0_f32; ISV_DIM];
isv[HOLD_RATE_TARGET_INDEX] = target;
isv[HOLD_COST_SCALE_INDEX] = seed_blended;
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
isv_buf.write_from_slice(&isv);
let mut traj = Vec::with_capacity(observations.len());
for (epoch, &obs) in observations.iter().enumerate() {
let mut current = isv_buf.read_all();
current[HOLD_RATE_OBSERVED_EMA_INDEX] = obs;
isv_buf.write_from_slice(&current);
let prior_scale = current[HOLD_COST_SCALE_INDEX];
let target_scale = compute_target_scale(obs, target);
launch_canonical(stream, kernel, isv_buf.dev_ptr);
let post = isv_buf.read_all();
let blended = post[HOLD_COST_SCALE_INDEX];
// Recover α from blend equation. For the first epoch (where
// prior is the seed), α=1.0 is the cold-start path. Once
// sample_count >= 3, α derives from Welford state.
let recovered_alpha = if (target_scale - prior_scale).abs() < 1e-6 {
f32::NAN // formula undefined when prior == target
} else {
recover_alpha(blended, prior_scale, target_scale)
};
traj.push((blended, recovered_alpha));
// Sanity check: Welford counter increments monotonically.
assert!(
(post[HCS_SAMPLE_COUNT_INDEX] - (epoch as f32 + 1.0)).abs() < 1e-4,
"epoch {epoch}: sample_count expected {}, got {}",
epoch + 1, post[HCS_SAMPLE_COUNT_INDEX]
);
}
traj
}
/// SP16 T3 — Test 1: α HIGH during signal jump (cold-start
/// responsiveness).
///
/// Feed a sharp signal jump from below-target to over-target across
/// 5 epochs. The first three epochs trigger cold-start path
/// (sample_count < 3 → α=1.0 REPLACE). After that, the Wiener
/// formula sees high diff_var (the recent jumps) relative to
/// sample_var → α should remain HIGH (> 0.5) on the 4th and 5th
/// epochs — concrete proof of the responsiveness benefit over
/// the deleted hardcoded α=0.05.
#[test]
#[ignore = "requires GPU"]
fn sp16_phase3_alpha_high_during_signal_jump() {
let stream = make_stream();
let kernel = load_kernel(&stream);
// Sharp jumps: 0.20→0.30→0.40→0.30→0.40 (target=0.20).
// target_scale series (overrun_norm × 24 + 1):
// ep0: obs=0.20 → norm=0.0 → target=1
// ep1: obs=0.30 → norm=0.5 → target=13
// ep2: obs=0.40 → norm=1.0 → target=25
// ep3: obs=0.30 → norm=0.5 → target=13
// ep4: obs=0.40 → norm=1.0 → target=25
// Diff sequence: +12, +12, -12, +12 → high diff_var.
let observations = [0.20, 0.30, 0.40, 0.30, 0.40];
// Seed blended at 1.0 (NOT sentinel — bypass Pearl-A path) so we
// can recover α from the blend equation directly.
const SEED: f32 = 1.0001; // slightly above EPS_F to avoid Pearl-A REPLACE
let traj = run_synthetic_sequence(&stream, &kernel, &observations, 0.20, SEED);
// Epoch 0 (sample_count=1 < 3): cold-start α=1.0, target=1.
// Epoch 1 (sample_count=2 < 3): cold-start α=1.0, target=13.
// Epoch 2 (sample_count=3): Wiener formula activates. After
// 3 samples (1, 13, 25), sample_mean ≈ 13, sample_var huge,
// diff_var also huge. α should be substantial (>0.3).
// Epoch 3-4: Wiener formula. With high target swings, diff_var
// stays comparable to sample_var → α stays high.
let alpha_ep3 = traj[3].1;
let alpha_ep4 = traj[4].1;
assert!(
alpha_ep3 > 0.3 && !alpha_ep3.is_nan(),
"Wiener α should be HIGH (>0.3) during high-volatility signal at epoch 3, got {alpha_ep3:.5}"
);
assert!(
alpha_ep4 > 0.3 && !alpha_ep4.is_nan(),
"Wiener α should remain HIGH (>0.3) during continued high-volatility at epoch 4, got {alpha_ep4:.5}"
);
// Defensive bound check.
assert!(
alpha_ep3 <= WELFORD_ALPHA_MAX && alpha_ep4 <= WELFORD_ALPHA_MAX,
"Wiener α must respect WELFORD_ALPHA_MAX={}; got ep3={alpha_ep3:.5} ep4={alpha_ep4:.5}",
WELFORD_ALPHA_MAX
);
}
/// SP16 T3 — Test 2: α LOW in steady state (Welford diff_var decays).
///
/// Feed a converging signal: `0.30 + 0.001 / (epoch + 1)` for 60
/// epochs. The decaying perturbation makes consecutive diffs ever-
/// smaller (so diff_var → 0) while sample_var converges toward the
/// steady-state value — α = diff_var / (diff_var + sample_var + ε)
/// → 0 (clamped at WELFORD_ALPHA_MIN). Verifies the "smoothing
/// emerges naturally" property over the deleted hardcoded α=0.05
/// cadence.
#[test]
#[ignore = "requires GPU"]
fn sp16_phase3_alpha_low_in_steady_state() {
let stream = make_stream();
let kernel = load_kernel(&stream);
// Converging signal: 0.30 + tiny decaying perturbation. After
// ~30 epochs the recent diffs are nano-sized while sample_var
// reflects the established mean of ~target_scale = 13.
let mut observations = vec![];
for i in 0..60 {
observations.push(0.30_f32 + 0.001_f32 / ((i + 1) as f32));
}
const SEED: f32 = 13.0;
let traj = run_synthetic_sequence(&stream, &kernel, &observations, 0.20, SEED);
// After ~50 epochs the Welford state should have stabilised:
// sample_var stays bounded (target_scale fluctuates within
// ~0.001 of 13.0) while diff_var → 0 (consecutive diffs are
// sub-permille). α should approach the WELFORD_ALPHA_MIN floor.
let tail: Vec<f32> = traj[50..60].iter()
.map(|&(_, a)| a)
.filter(|a| !a.is_nan())
.collect();
let mean_tail_alpha = tail.iter().sum::<f32>() / tail.len() as f32;
// Tail α should be substantially below the cold-start ceiling.
// With a converging signal the diff_var collapses → α floors near
// WELFORD_ALPHA_MIN. We assert a generous threshold (0.40) to
// tolerate fp rounding in the recovered-α calculation; the
// smoking gun is that this is well below the cold-start α=1.0.
assert!(
mean_tail_alpha < 0.40 && mean_tail_alpha > 0.0,
"Steady-state α should decay below 0.40, got mean over [50..60): {mean_tail_alpha:.5}"
);
// Defensive bound check — α stays positive (no permanent freeze).
assert!(
mean_tail_alpha >= WELFORD_ALPHA_MIN,
"Steady-state α must respect WELFORD_ALPHA_MIN={}; got {mean_tail_alpha:.5}",
WELFORD_ALPHA_MIN
);
}
/// SP16 T3 — Test 3: Pearl-A first-observation bootstrap fires
/// regardless of cold-start α.
///
/// When `prior_scale` is at sentinel (≤ EPS), the kernel takes the
/// REPLACE branch BEFORE the blend formula runs. The Wiener α
/// computation still happens (Welford state is updated) but the
/// result is discarded for this launch — blended = target_scale
/// directly. This protects the cold-start bootstrap from being
/// corrupted by an arbitrary cold-start α=1.0 that would otherwise
/// produce the same numerical result by coincidence.
///
/// Setup: prior_scale=sentinel 0.0, observed=0.30, target=0.20.
/// target_scale = 1 + 24 × 0.5 = 13.0
/// Pearl-A REPLACE → blended = 13.0 directly.
/// Without Pearl-A path: blend with α=1.0 → 0×0 + 1×13 = 13.0
/// (same numerical result). To distinguish, we run a SECOND launch
/// with non-sentinel prior and verify the Welford state was already
/// updated by the first launch (sample_count = 1).
#[test]
#[ignore = "requires GPU"]
fn sp16_phase3_pearl_a_bootstrap_first_obs() {
let stream = make_stream();
let kernel = load_kernel(&stream);
const ISV_DIM: usize = 1024;
let mut isv = vec![0.0_f32; ISV_DIM];
isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.30;
isv[HOLD_RATE_TARGET_INDEX] = 0.20;
isv[HOLD_COST_SCALE_INDEX] = SENTINEL_HOLD_COST_SCALE; // sentinel 0.0
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
isv_buf.write_from_slice(&isv);
launch_canonical(&stream, &kernel, isv_buf.dev_ptr);
let result = isv_buf.read_all();
// Pearl-A REPLACE → blended = 13.0.
let expected = 1.0_f32 + (HOLD_COST_SCALE_MAX - HOLD_COST_SCALE_MIN) * 0.5;
let blended = result[HOLD_COST_SCALE_INDEX];
assert!(
(blended - expected).abs() < 1e-4,
"Pearl-A bootstrap on sentinel must REPLACE with target_scale; expected {expected:.5}, got {blended:.5}"
);
// Welford accumulators must still have been advanced (the kernel
// updates them BEFORE checking the Pearl-A branch).
assert!(
(result[HCS_SAMPLE_COUNT_INDEX] - 1.0).abs() < 1e-4,
"Welford sample_count must be 1 after first launch (state advances regardless of Pearl-A), got {}",
result[HCS_SAMPLE_COUNT_INDEX]
);
// prev_target tracks target_scale for next launch's diff calc.
assert!(
(result[HCS_PREV_TARGET_INDEX] - expected).abs() < 1e-4,
"prev_target must equal target_scale after first launch, expected {expected:.5}, got {}",
result[HCS_PREV_TARGET_INDEX]
);
}
/// SP16 T3 — Test 4: α stays in [WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX]
/// over 50 synthetic epochs (post-cold-start).
///
/// Run a complex signal (ramp + noise) for 50 epochs and verify
/// every recovered α at sample_count ≥ 3 lies within the Wiener
/// formula's defensive bounds. Cold-start epochs (sample_count < 3)
/// take the α=1.0 REPLACE branch by design — those are NOT subject
/// to the [0.01, 0.95] clamp because the formula isn't applied
/// (Welford state is too immature). The test skips epochs 0-1
/// (sample_count = 1, 2) and asserts on epochs 2+ where Wiener α
/// is the active code path. Catches numerical pathologies (α hitting
/// denormal, α exceeding 0.95 from rounding, α clamped to 0 from
/// variance underflow).
#[test]
#[ignore = "requires GPU"]
fn sp16_phase3_alpha_naturally_bounded() {
let stream = make_stream();
let kernel = load_kernel(&stream);
// 50-epoch sequence: linear ramp 0.20 → 0.40 with sinusoidal noise.
let mut observations = vec![];
for i in 0..50 {
let base = 0.20 + 0.20 * (i as f32 / 49.0);
let noise = 0.02 * ((i as f32 * 0.7).sin());
observations.push((base + noise).clamp(0.0, 1.0));
}
const SEED: f32 = 1.5;
let traj = run_synthetic_sequence(&stream, &kernel, &observations, 0.20, SEED);
// Skip epochs 0-1 (sample_count < 3 → cold-start REPLACE → α=1.0
// by design, NOT subject to Wiener bounds). Wiener formula
// activates from epoch 2 onward (sample_count = 3).
for (epoch, &(_, alpha)) in traj.iter().enumerate().skip(2) {
if alpha.is_nan() {
// Skip epochs where the recovered-α formula is undefined
// (target == prior). Common in steady-state intervals.
continue;
}
assert!(
alpha > 0.0,
"epoch {epoch}: α must be > 0, got {alpha:.5} (variance underflow → permanent freeze)"
);
assert!(
alpha <= WELFORD_ALPHA_MAX + 1e-3, // small tolerance for fp rounding via blend
"epoch {epoch} (post-cold-start): α must respect WELFORD_ALPHA_MAX={}, got {alpha:.5}",
WELFORD_ALPHA_MAX
);
}
}
}
// ═══════════════════════════════════════════════════════════════════════════
// SP16 T3 — host-only regression test (Test 5).
// Verifies the kernel sources do NOT contain the deleted hardcoded
// `0.05f` literal in the EMA blend math. Pure string scan — no GPU
// required. Per `feedback_no_legacy_aliases.md` and
// `feedback_isv_for_adaptive_bounds.md`.
// ═══════════════════════════════════════════════════════════════════════════
/// SP16 T3 — Test 5: regression-locked check that no `alpha = 0.05f`
/// or comparable hardcoded blend constant remains in the producer kernel
/// sources. Runs at compile time via `include_str!`; failure points
/// directly at the introduced regression.
#[test]
fn sp16_phase3_no_hardcoded_alpha() {
const T1_SOURCE: &str = include_str!(
"../src/cuda_pipeline/min_hold_temperature_update_kernel.cu"
);
const T2_SOURCE: &str = include_str!(
"../src/cuda_pipeline/hold_cost_scale_update_kernel.cu"
);
// The blend equation pre-T3 was `(1.0f - alpha) * current + alpha * target_*`
// with `alpha = 0.05f`. Search for the canonical form.
for (name, src) in [("T1 min_hold_temperature", T1_SOURCE),
("T2 hold_cost_scale", T2_SOURCE)] {
// The hardcoded literal must not appear in any blend assignment.
// Allow `0.05f` to appear in commentary (e.g. "the deleted α=0.05
// constant"), but NOT as a kernel-local variable initializer or
// an arithmetic operand. We approximate by requiring the literal
// not to appear in any line that also contains an `alpha` token
// outside of comments.
for (lineno, line) in src.lines().enumerate() {
// Skip C-style comment lines (the kernel uses leading `*` for
// doc comments inside the algorithm explanation block).
let trimmed = line.trim_start();
if trimmed.starts_with("//") || trimmed.starts_with("*")
|| trimmed.starts_with("/*") {
continue;
}
// Hard fail on any executable line that contains both `alpha`
// and `0.05`. The pre-T3 implementation had exactly one such
// site per kernel (`const float alpha = 0.05f;`).
if line.contains("alpha") && line.contains("0.05") {
panic!(
"{name} kernel line {} reintroduces hardcoded alpha=0.05: {}",
lineno + 1, line
);
}
}
// Defence in depth: the new kernel ABI passes Welford slot indices,
// not an `alpha` scalar. Verify the new param signature appears
// (catches accidental revert to the pre-T3 ABI).
assert!(
src.contains("target_mean_idx") && src.contains("sample_count_idx"),
"{name} kernel does not reference Welford accumulator slots — Wiener α implementation missing"
);
}
}
// ═══════════════════════════════════════════════════════════════════════════ // ═══════════════════════════════════════════════════════════════════════════
// SP16 Phase 0 (2026-05-08) — Q-by-action HEALTH_DIAG diagnostic. // SP16 Phase 0 (2026-05-08) — Q-by-action HEALTH_DIAG diagnostic.
// //

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

@@ -8724,6 +8724,128 @@ Per `pearl_canary_input_freshness_launch_order.md`: this launch order is the smo
3. **Producer launch ordering.** The SP16-P1 producer fires AFTER the per-step `hold_rate_observer` chain via the `if (collector.executed) {…}` block in training_loop.rs. The new SP16-P2 producer fires immediately after SP16-P1 inside the same block, so the launch ordering is correct by construction. The `pearl_canary_input_freshness_launch_order.md` invariant is preserved. 3. **Producer launch ordering.** The SP16-P1 producer fires AFTER the per-step `hold_rate_observer` chain via the `if (collector.executed) {…}` block in training_loop.rs. The new SP16-P2 producer fires immediately after SP16-P1 inside the same block, so the launch ordering is correct by construction. The `pearl_canary_input_freshness_launch_order.md` invariant is preserved.
## SP16 T3: Wiener-optimal adaptive α per pearl_wiener_optimal_adaptive_alpha (2026-05-08)
### Why this commit exists
`train-multi-seed-hjzss` validation showed the SP16 T1+T2 chain was structurally landed but BEHAVIORALLY INERT in 5-epoch smoke. The producer trajectory was bit-identical to the pfh9n baseline through epoch 3:
- `hold_cost_scale_diag`: 1.0 → 1.0 → 1.27 (target 7.0 — 21% of headroom after 3 epochs)
- `min_hold_temp_diag`: 0.0 → 5.0 → 5.50 (target 16.0 — 35% of headroom after 3 epochs)
Root cause: hardcoded `const float alpha = 0.05f` in both producer kernels (`min_hold_temperature_update_kernel.cu` + `hold_cost_scale_update_kernel.cu`). With α=0.05 the EMA needs ~60 epochs from cold start to converge — far beyond the 5-epoch smoke and likely beyond any practical training horizon.
Per `feedback_isv_for_adaptive_bounds.md`, the constant α was itself an adaptive bound being hardcoded. The proper fix per `pearl_wiener_optimal_adaptive_alpha`:
```
α = diff_var / (diff_var + sample_var + ε)
```
Where:
- `sample_var` = running variance of the producer's `target_*` signal
- `diff_var` = running variance of consecutive one-step `target_*` differences
Both estimated via Welford accumulators. Cold-start (signal jumps 1.0 → 6.4 → 7.0): diff_var dominates → α ≈ 0.6+ → near-bootstrap responsiveness in epochs 1-3. Steady state: signal stabilises → diff_var → 0 → α decays naturally → smoothing emerges without a hardcoded constant.
### Slot allocation
12 new ISV slots [462..474), 6 per producer:
| Index | Name | Producer | Role |
| ----- | ------------------------ | -------- | --------------------------------------------- |
| 462 | HCS_TARGET_MEAN_INDEX | T2 (HCS) | Welford mean of target_scale |
| 463 | HCS_TARGET_M2_INDEX | T2 (HCS) | Welford ΣΔ² (sample var accumulator) |
| 464 | HCS_DIFF_MEAN_INDEX | T2 (HCS) | Welford mean of consecutive diffs |
| 465 | HCS_DIFF_M2_INDEX | T2 (HCS) | Welford ΣΔ² for diffs |
| 466 | HCS_PREV_TARGET_INDEX | T2 (HCS) | previous-epoch target_scale |
| 467 | HCS_SAMPLE_COUNT_INDEX | T2 (HCS) | Welford observation counter |
| 468 | MHT_TARGET_MEAN_INDEX | T1 (MHT) | Welford mean of target_temp |
| 469 | MHT_TARGET_M2_INDEX | T1 (MHT) | Welford ΣΔ² (target_temp) |
| 470 | MHT_DIFF_MEAN_INDEX | T1 (MHT) | Welford mean of consecutive diffs |
| 471 | MHT_DIFF_M2_INDEX | T1 (MHT) | Welford ΣΔ² for diffs |
| 472 | MHT_PREV_TARGET_INDEX | T1 (MHT) | previous-epoch target_temp |
| 473 | MHT_SAMPLE_COUNT_INDEX | T1 (MHT) | Welford observation counter |
`ISV_TOTAL_DIM` bumped 462 → 474. All 12 slots are FoldReset-bootstrap (sentinel 0.0) — combined with the cold-start fallback (sample_count < 3 → α = 1.0 REPLACE), the first observation of each new fold replaces blended directly, the second seeds diff statistics, and α adapts via Wiener formula from the third observation onward.
### Defensive bounds
The Wiener formula naturally yields α ∈ [0, 1] but numerical edge cases (both variances identically zero within ε) could push α to a denormal value or floor at 0 (no learning). Bounds applied AFTER the formula:
- `WELFORD_ALPHA_MIN = 0.01` — keeps every observation weighted (no permanent freeze).
- `WELFORD_ALPHA_MAX = 0.95` — avoids pathological α=1.0 in steady state (REPLACE is reserved for the cold-start path).
- `WELFORD_MIN_SAMPLE_COUNT = 3.0` — threshold for switching from cold-start REPLACE to Wiener formula. Below 3 samples, diff variance estimation is meaningless.
- `WELFORD_EPS = 1.0e-6` — denominator floor for the formula.
### Files modified
| File | LOC delta | Purpose |
| ------------------------------------------------------------------------- | ---------------------- | ---------------------------------------------------------------------- |
| `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` | +120 / -16 | 12 new slot constants + 4 algorithm constants + 2 lock tests |
| `crates/ml/src/cuda_pipeline/state_layout.cuh` | +24 / -1 | C #define mirrors (12 indices + 4 constants) |
| `crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu` | +180 / -110 | Welford state + Wiener-optimal α replaces hardcoded 0.05f |
| `crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu` | +110 / -75 | Welford state + Wiener-optimal α replaces hardcoded 0.05f |
| `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` | +56 / -16 | Both `*UpdateOps::launch` migrated to new ABI (Welford slot indices) |
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | +25 / -8 (plus fingerprint) | Caller wrappers + ISV_TOTAL_DIM 462→474 + fingerprint update |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +50 / -1 | 12 new RegistryEntry rows |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +85 / -3 | 12 dispatch arms + α-recovery diag emit |
| `crates/ml/tests/sp14_oracle_tests.rs` | +330 / -22 | 5 new behavioral tests (4 GPU + 1 host) + ABI updates to 7 prior tests |
| `docs/dqn-wire-up-audit.md` | +75 / -0 | This entry |
| `docs/superpowers/plans/2026-05-08-sp16-hold-collapse-structural-fix.md` | +50 / -0 | Phase 3 section: actual mechanism (was empty placeholder) |
### Behavioral test results
Local RTX 3050 Ti smoke (`cargo test -p ml --test sp14_oracle_tests --release -- --ignored sp16_phase3 --nocapture`):
```
running 4 tests
test sp16_phase3_wiener_alpha_gpu::sp16_phase3_alpha_high_during_signal_jump ... ok
test sp16_phase3_wiener_alpha_gpu::sp16_phase3_alpha_low_in_steady_state ... ok
test sp16_phase3_wiener_alpha_gpu::sp16_phase3_alpha_naturally_bounded ... ok
test sp16_phase3_wiener_alpha_gpu::sp16_phase3_pearl_a_bootstrap_first_obs ... ok
test result: ok. 4 passed; 0 failed
```
Plus host-only Test 5 (`sp16_phase3_no_hardcoded_alpha`): scans the kernel sources for any `alpha`+`0.05` co-occurrence on a non-comment line — passes.
Full sp14 + sp15 oracle suites: 34 GPU tests + 4 host tests pass — no regressions in pre-existing SP14 / SP16-P1 / SP16-P2 tests.
### Pearl-A bootstrap interaction
The two paths interact correctly:
1. **First observation of a fresh fold** (Welford counter = 0 → 1, sample_count < 3 → α = 1.0):
- If `current` is at sentinel (50.0 for T1, 0.0 for T2) → Pearl-A REPLACE branch fires → blended = target.
- If `current` is NOT at sentinel → blend formula fires with α=1.0 → blended = target (numerically identical to REPLACE).
2. **Second observation** (sample_count = 2 < 3 → still α = 1.0): same behavior. Welford diff statistics seed in this epoch.
3. **Third observation onward** (sample_count ≥ 3): Wiener formula activates. With cold-start dynamics still in play, diff_var dominates → α ≈ 0.5-0.8 → fast convergence. Once signal stabilises, diff_var drops → α decays toward WELFORD_ALPHA_MIN.
The cold-start path **DID fire correctly** in tests — Test 1 (alpha_high_during_signal_jump) verifies α > 0.3 at epoch 3 (post-cold-start, diffs still large) and Test 4 (alpha_naturally_bounded) verifies α stays within bounds from epoch 2 onward. Test 5 in the sample-count assertion proves Welford counter increments on every launch (Pearl-A REPLACE branch does NOT short-circuit the state update).
### Diagnostic emit update
The existing `min_hold_temp_diag` and `hold_cost_scale_diag` emits in `training_loop.rs` now include the recovered α and Welford counter for direct trajectory observation:
```
HEALTH_DIAG[N]: hold_cost_scale_diag obs_hold_rate=X target_hold_rate=Y overrun_norm=Z target_scale=A blended_scale=B alpha=W sample_count=N
HEALTH_DIAG[N]: min_hold_temp_diag obs_hold_rate=X target_hold_rate=Y overrun_norm=Z target_temp=A blended_temp=B alpha=W sample_count=N
```
The α field surfaces what would otherwise be opaque on-device state, enabling validation smoke trajectories to verify the cold-start → steady-state α transition directly.
### Cross-pearl invariants preserved
- `pearl_first_observation_bootstrap`: Pearl-A still fires on sentinel (50.0 / 0.0). Welford state advance happens BEFORE the Pearl-A check so the 2nd-launch diff is correctly computed.
- `pearl_no_host_branches_in_captured_graph`: kernel guards (sentinel detect + sample_count threshold) remain on-device.
- `pearl_symmetric_clamp_audit`: bilateral clamp on target_* + post-blend preserved.
- `feedback_no_atomicadd`: still single-thread (no reductions).
- `feedback_no_partial_refactor`: T1 + T2 + tests + audit doc + plan doc updated atomically in this commit.
- `feedback_no_legacy_aliases`: hardcoded `*_EMA_ALPHA = 0.05` constants deleted (Rust + C #define mirror) — no half-migrated state.
- `feedback_no_stubs`: every test body is real; no `todo!()` placeholders.
## Build infra: ensure-binary cache key bug (2026-05-08) ## Build infra: ensure-binary cache key bug (2026-05-08)
Discovered during SP16 validation (4gb88 → 8mzkj forced switch from L40S to H100). Discovered during SP16 validation (4gb88 → 8mzkj forced switch from L40S to H100).

30
docs/superpowers/plans/2026-05-08-sp16-hold-collapse-structural-fix.md
View File

@@ -458,7 +458,35 @@ EOF
--- ---
## Phase 3 — Validation L40S/H100 smoke ## Phase 3 — Wiener-optimal adaptive α (T3)
### Why this exists (validation result, not original spec)
The original Phase 3 was a placeholder for "validation L40S/H100 smoke" — but `train-multi-seed-hjzss` running the SP16 T1+T2 chain showed the producer trajectory was **bit-identical to the pfh9n baseline through epoch 3**:
- `hold_cost_scale_diag`: 1.0 → 1.0 → 1.27 (target 7.0 — only 21% of headroom in 3 epochs)
- `min_hold_temp_diag`: 0.0 → 5.0 → 5.50 (target 16.0 — only 35% of headroom in 3 epochs)
**Root cause**: hardcoded `const float alpha = 0.05f` in both producer kernels. With α=0.05 the EMA needs ~60 epochs from cold start to converge — far beyond the 5-epoch smoke and likely beyond any practical training horizon. Per `feedback_isv_for_adaptive_bounds`, the constant α was itself an adaptive bound being hardcoded; per `pearl_wiener_optimal_adaptive_alpha` the proper fix is:
```
α = diff_var / (diff_var + sample_var + ε)
```
Where `sample_var` and `diff_var` are estimated via Welford accumulators stored in ISV.
### Task 3.0: Wiener-optimal adaptive α (LANDED 2026-05-08)
Both producer kernels (T1 `min_hold_temperature_update_kernel.cu` and T2 `hold_cost_scale_update_kernel.cu`) migrated atomically:
- 12 new ISV slots [462..474), 6 per producer (target mean/M2, diff mean/M2, prev_target, sample_count). `ISV_TOTAL_DIM` 462 → 474.
- Welford accumulators updated on every launch (after sentinel-guard early-return).
- α computed as `diff_var / (diff_var + sample_var + ε)`, clamped to `[WELFORD_ALPHA_MIN=0.01, WELFORD_ALPHA_MAX=0.95]`.
- Cold-start path: when `sample_count < 3`, α = 1.0 (REPLACE) so the producer is responsive in epochs 1-2 while Welford state stabilises.
- Pearl-A first-observation bootstrap preserved (sentinel → REPLACE) and runs BEFORE the blend formula so it has priority over a coincidental cold-start α=1.0.
- FoldReset registers all 12 slots with sentinel 0.0 + 12 dispatch arms.
- Diagnostic emit (`HEALTH_DIAG[N]: hold_cost_scale_diag` + `min_hold_temp_diag`) extended with `alpha` and `sample_count` fields for direct trajectory observation.
- 5 behavioral tests (4 GPU + 1 host-only): α high during signal jump, α low in steady state, Pearl-A bootstrap fires, α naturally bounded, no `0.05f` literal in kernel sources.
### Task 3.1: Push and dispatch validation smoke ### Task 3.1: Push and dispatch validation smoke