diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index 22b2a2859..b8a8d85e4 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -966,40 +966,55 @@ impl DdSaturationFloorUpdateOps { } } -/// Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive -/// MIN_HOLD_TEMPERATURE producer. +/// SP16 Phase 1 (revised, 2026-05-08): adaptive MIN_HOLD_TEMPERATURE +/// producer driven by hold-rate overrun (signal swap from slot 373). /// /// Single-thread cold-path kernel (per-epoch boundary cadence) that -/// reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of -/// binary directional accuracy), maps it to a [5, 50] temperature via +/// reads two SP13 hold-pricing chain slots from ISV: /// -/// skill = clamp((short_ema − 0.5) / 0.5, 0, 1) -/// confusion = 1 − skill -/// temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × confusion +/// - `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` — per-fold EMA of the +/// per-step Hold-pick rate, populated by `apply_fixed_alpha_ema_kernel` +/// chained off `hold_rate_observer_kernel`. +/// - `ISV[HOLD_RATE_TARGET_INDEX=381]` — static target (default 0.20). /// -/// and slow-EMA-blends into +/// And maps the overrun into a [5, 50] temperature: +/// +/// overrun = max(0, observed − target) +/// overrun_norm = clamp(overrun / max(target, 0.01), 0, 1) +/// temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × overrun_norm +/// +/// then slow-EMA-blends into /// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]` with α=0.05. /// /// 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). The `confusion = 1 − skill` mapping therefore -/// gives the WR-plateau scenario its designed behavior: -/// - dir_acc ≈ 0.5 (random / committed-but-wrong) → confusion=1 → -/// temp=50 (permissive — exit ramp out of bad commitments) -/// - dir_acc → 1.0 (saturated skill) → confusion=0 → temp=5 -/// (strict — force the model to stay in informed positions) +/// max exit penalty). The hold-rate-overrun mapping therefore gives +/// the over-holding scenario its designed behaviour: +/// - observed = 2× target → overrun_norm = 1.0 → temp = 50 +/// (permissive — exit ramp out of over-holding) +/// - observed = target → overrun_norm = 0 → temp = 5 +/// (strict — force commitment at the hold-rate target) +/// - observed < target → overrun_norm = 0 → temp = 5 +/// (under-holding — strict, no relaxation) /// -/// Replaces the deleted epoch-driven anneal schedule -/// `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly -/// state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY} + -/// training_loop.rs::min_hold_temperature_for_epoch). The original -/// schedule's `START=50 → END=5` was "permissive early, strict late" -/// (classic explore-then-exploit, time-anchored). Decouples -/// temperature from epoch number — when WR plateaus, the old -/// schedule pinned LOW temp (sharp) at end of training (max -/// punishment exactly when most needed forgiveness); the signal-driven -/// version reacts to actual realized skill. +/// Signal swap rationale (post-mortem of train-multi-seed-pfh9n, +/// 2026-05-08): the original mapping read +/// `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` but slot 373 stayed at +/// sentinel 0.5 in Fold 1 — the per-step aux producer either didn't +/// push it off sentinel quickly enough or the binary classifier +/// converged within ε of 0.5 in noisy environments. The kernel's +/// early-return-on-sentinel guard then fired on every launch and slot +/// 460 stayed pinned at 50 for the entire fold. Hold rates in +/// contrast are measured per-epoch from realised actions and never at +/// a "no-data" sentinel after epoch 1, so the fold-1 chain breakage +/// goes away. +/// +/// Cold-start sentinel guard preserved: when `observed_hold_rate` is +/// at the SP13 fold-reset sentinel 0.0 (no per-step Hold observations +/// yet in the new fold), the kernel keeps slot 460 unchanged — the +/// consumer falls back to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 via the +/// same sentinel-detect path used pre-swap. /// /// # Pearls applied /// - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. @@ -1012,8 +1027,8 @@ impl DdSaturationFloorUpdateOps { /// Category-1 dimensional safety floors (matches the original /// schedule range MIN_HOLD_TEMPERATURE_END=5 ↔ /// MIN_HOLD_TEMPERATURE_START=50), NOT tuning. -/// - `feedback_no_legacy_aliases.md` — the legacy schedule constants -/// and helper function are deleted atomically with this wiring. +/// - `feedback_no_partial_refactor.md` — kernel + launcher + call site +/// + tests + audit doc updated atomically in the same commit. #[allow(missing_debug_implementations)] pub(crate) struct MinHoldTemperatureUpdateOps { update_kernel: CudaFunction, @@ -1031,15 +1046,18 @@ impl MinHoldTemperatureUpdateOps { Ok(Self { update_kernel }) } - /// Launch the adaptive MIN_HOLD_TEMPERATURE producer. + /// Launch the adaptive MIN_HOLD_TEMPERATURE producer (SP16 Phase 1 + /// revised — hold-rate-overrun driven). /// /// Args: /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). /// - `temp_idx`: MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460). - /// - `dir_acc_idx`: AUX_DIR_ACC_SHORT_EMA_INDEX (373). + /// - `hold_rate_observed_idx`: HOLD_RATE_OBSERVED_EMA_INDEX (382). + /// - `hold_rate_target_idx`: HOLD_RATE_TARGET_INDEX (381). /// - `sentinel_temp`: SENTINEL_MIN_HOLD_TEMPERATURE (50.0). - /// - `sentinel_dir_acc`: DIR_ACC_RANDOM_BASELINE (0.5; matches - /// sp13 DIR_ACC_EMA_SENTINEL). + /// - `sentinel_observed_hold_rate`: SENTINEL_HOLD_RATE_OBSERVED + /// (0.0; matches the SP13 fold-reset registry sentinel for + /// `sp13_hold_rate_observed_ema`). /// - `temp_min`/`temp_max`: MIN_HOLD_TEMPERATURE_MIN/MAX (5.0, 50.0). /// - `alpha`: MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05). #[allow(clippy::too_many_arguments)] @@ -1048,9 +1066,10 @@ impl MinHoldTemperatureUpdateOps { stream: &Arc, isv_ptr: u64, temp_idx: i32, - dir_acc_idx: i32, + hold_rate_observed_idx: i32, + hold_rate_target_idx: i32, sentinel_temp: f32, - sentinel_dir_acc: f32, + sentinel_observed_hold_rate: f32, temp_min: f32, temp_max: f32, alpha: f32, @@ -1060,9 +1079,10 @@ impl MinHoldTemperatureUpdateOps { .launch_builder(&self.update_kernel) .arg(&isv_ptr) .arg(&temp_idx) - .arg(&dir_acc_idx) + .arg(&hold_rate_observed_idx) + .arg(&hold_rate_target_idx) .arg(&sentinel_temp) - .arg(&sentinel_dir_acc) + .arg(&sentinel_observed_hold_rate) .arg(&temp_min) .arg(&temp_max) .arg(&alpha) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index a7a8e3d87..84016b932 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -15025,24 +15025,34 @@ impl GpuDqnTrainer { &self, ) -> Result<(), MLError> { use crate::cuda_pipeline::sp14_isv_slots::{ - DIR_ACC_RANDOM_BASELINE, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, MIN_HOLD_TEMPERATURE_EMA_ALPHA, MIN_HOLD_TEMPERATURE_MAX, - MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE, + MIN_HOLD_TEMPERATURE_MIN, SENTINEL_HOLD_RATE_OBSERVED, + SENTINEL_MIN_HOLD_TEMPERATURE, + }; + use crate::cuda_pipeline::sp13_isv_slots::{ + HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, }; - use crate::cuda_pipeline::sp5_isv_slots::AUX_DIR_ACC_SHORT_EMA_INDEX; debug_assert!( self.isv_signals_dev_ptr != 0, "launch_min_hold_temperature_update: isv_signals_dev_ptr must be allocated" ); + // SP16 Phase 1 (revised, 2026-05-08): driving signal swapped + // from AUX_DIR_ACC_SHORT_EMA (slot 373) to hold-rate overrun + // (slots 382 + 381). Slot 373 stayed at sentinel 0.5 in Fold 1 + // and the kernel's early-return guard pinned slot 460 at 50 + // for the entire fold. Hold rates are measured per-epoch from + // realised actions and survive fold reset cleanly. self.min_hold_temperature_update_ops.launch( &self.stream, self.isv_signals_dev_ptr, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, - AUX_DIR_ACC_SHORT_EMA_INDEX as i32, + HOLD_RATE_OBSERVED_EMA_INDEX as i32, + HOLD_RATE_TARGET_INDEX as i32, SENTINEL_MIN_HOLD_TEMPERATURE, - DIR_ACC_RANDOM_BASELINE, + SENTINEL_HOLD_RATE_OBSERVED, MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, MIN_HOLD_TEMPERATURE_EMA_ALPHA, diff --git a/crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu b/crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu index a68c9ee7a..6ac1c49e3 100644 --- a/crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu +++ b/crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu @@ -1,83 +1,89 @@ /* ══════════════════════════════════════════════════════════════════════════ - * Class A audit-fix Batch 4-B — adaptive MIN_HOLD_TEMPERATURE producer - * (2026-05-08, Item 4). Fixup 2026-05-08: inverted the dir_acc → temp - * mapping after kernel-formula audit (`compute_min_hold_penalty` at - * trade_physics.cuh:575): HIGH T → soft_factor → 0 → forgiving exits; - * LOW T → soft_factor → 1 → sharp penalty. The original schedule - * `START=50 → END=5` therefore means "permissive early, strict late" - * (classic explore-then-exploit). The pre-fixup mapping - * `temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill` was BACKWARDS for - * the WR-plateau scenario this 8-commit chain targets: at plateau - * (dir_acc ≈ 0.46-0.48, model committed but wrong) skill = 0 → temp = 5 - * (sharp) → maximum exit penalty exactly when the model most needs to - * escape the bad committed direction. Inverted to `confusion = 1 - skill` - * so plateau → temp HIGH (permissive, exit ramp) and skilled → - * temp LOW (strict, force commitment). + * SP16 Phase 1 (revised, 2026-05-08) — adaptive MIN_HOLD_TEMPERATURE producer. * - * Replaces the deleted epoch-driven anneal schedule - * `T_end + (T_start - T_end) × exp(-epoch / decay)` - * (formerly state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, - * DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch) with an - * ISV-driven signal-driven temperature derived from the model's - * directional-accuracy skill. The temperature controls the - * `deficit / (deficit + T)` saturation factor in `compute_min_hold_penalty` - * (trade_physics.cuh:575): HIGH T = forgiving (penalty saturates near - * 0 — easy to exit early); LOW T = sharp (penalty rises steeply at any - * deficit — hold is forced). + * SIGNAL SWAP: this kernel was originally driven by + * `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary + * directional accuracy). Post-mortem of train-multi-seed-pfh9n + * 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. * - * Why this matters (per Class A audit deferred-item batch 4-B): + * 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: * - * The epoch-driven schedule pinned LOW temp (T=5, sharp) at end of - * training. This works ONLY if the model gets more skilled with - * training — which is exactly the assumption that fails when WR - * plateaus. Maximum punishment is delivered exactly when the model - * most needs forgiveness to escape the plateau. The signal-driven - * replacement decouples temperature from epoch number: when the - * model is at random baseline / WR plateau (dir_acc ≈ 0.5, - * confusion ≈ 1) → temp HIGH (permissive, exit ramp to escape bad - * committed direction); when committing skillfully (dir_acc → 1.0, - * confusion → 0) → temp LOW (strict, force commitment because the - * commitment is informed). + * - `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` — per-fold EMA of the + * per-step Hold-pick rate `count(action_dir == DIR_HOLD) / B`, + * written by `apply_fixed_alpha_ema_kernel` chained off + * `hold_rate_observer_kernel` in the experience collector. * - * Driving signal: dir_acc skill from - * `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` — the SP13 fast EMA of - * binary directional accuracy (sentinel 0.5 = random baseline). + * - `ISV[HOLD_RATE_TARGET_INDEX=381]` — static target (default 0.20), + * constructor-init. * - * skill = clamp((short_ema - 0.5) / 0.5, 0, 1) ∈ [0, 1] - * confusion = 1 - skill ∈ [0, 1] - * temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × confusion + * New mapping: * - * - confusion=1 (no directional skill, dir_acc ≤ 0.5) → temp=TEMP_MAX=50 - * (permissive — give the model an exit ramp out of the plateau) - * - confusion=0 (saturated dir_acc=1.0) → temp=TEMP_MIN=5 - * (strict — force commitment because it's informed) + * overrun = max(0, observed − target) + * overrun_normalized = clamp(overrun / max(target, 0.01), 0, 1) + * target_temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × overrun_normalized * - * Cold-start fallback: if `isv[AUX_DIR_ACC_SHORT_EMA_INDEX]` is at - * sentinel (within EPS of 0.5) — i.e. fold reset just fired and no - * dir-acc samples have landed yet — fall back to the seeded - * SENTINEL_MIN_HOLD_TEMPERATURE (50.0, matches the pre-fix - * MIN_HOLD_TEMPERATURE_START=50 epoch-0 anchor). Bit-identical - * cold-start under the deleted schedule, which also had T=50 at - * epoch 0. + * - observed = 2× target → overrun_normalized = 1.0 → temp = 50 + * (permissive — exit ramp out of over-holding regime) + * - observed = target → overrun_normalized = 0 → temp = 5 + * (strict — force commitment when the model is at hold-rate target) + * - observed < target → overrun_normalized = 0 → temp = 5 + * (under-holding — strict, no relaxation needed) + * + * Per `compute_min_hold_penalty` at trade_physics.cuh:575 + * (`soft_factor = deficit / (deficit + T)`): + * - HIGH T → soft_factor → 0 → no exit penalty (permissive) + * - LOW T → soft_factor → 1 → max exit penalty (strict) + * + * So when the model is over-holding, temperature climbs HIGH (giving an + * 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: + * + * 1. Direct closed-loop control on the actual problem. Hold% is the + * symptom; hold-rate-overrun directly drives the corrective signal. + * + * 2. No chained-input-sentinel masking. Hold rates are measured per- + * epoch from realized actions, never at a "no-data" sentinel after + * epoch 1. + * + * 3. Survives fold reset cleanly. Hold-rate measurement starts fresh + * in Fold 1 with real data immediately on the first per-step + * observation; no need for the dir-acc EMA to climb out of 0.5. * * Algorithm (single-block single-thread — cold path, per-epoch boundary): * - * Phase 1 (single thread): read short_ema and current temp from ISV. + * Phase 1 (single thread): read observed_hold_rate (slot 382), + * target_hold_rate (slot 381), and current temp (slot 460) from ISV. * - * Phase 2 (Pearl-A + EMA): if dir_acc EMA is at sentinel 0.5 (within - * EPS), keep ISV slot unchanged — sentinel persists; consumer falls - * back to the kernel-passed scalar (seeded at - * MIN_HOLD_TEMPERATURE_FALLBACK=50.0). Otherwise: + * 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 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. * - * skill = max(0, min(1, (short_ema - 0.5) / 0.5)) - * confusion = 1 - skill - * target_temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × confusion + * Phase 3 (compute target temp): + * + * overrun = fmaxf(0, observed − target) + * overrun_norm = fminf(1, overrun / fmaxf(target, 0.01)) + * target_temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × overrun_norm * bilateral clamp to [TEMP_MIN, TEMP_MAX] * - * Pearl-A bootstrap: if current is at sentinel (within EPS of 50.0), - * REPLACE directly with target_temp. Otherwise EMA blend with α=0.05 - * (mid-cadence — slower than per-step skill EMA, faster than per-fold - * beliefs). + * Phase 4 (Pearl-A bootstrap + Welford-style EMA): + * + * if current is at sentinel (within EPS of 50.0) → REPLACE with target. + * else → blended = (1 − α) × current + α × target_temp + * re-clamp post-blend. * * Pearls + invariants: * @@ -87,7 +93,7 @@ * pre-loaded at construction; on-device guards for sentinel * detection. * - * - `pearl_first_observation_bootstrap.md` — first valid skill + * - `pearl_first_observation_bootstrap.md` — first valid hold-rate * observation replaces sentinel directly; no blend. * * - `feedback_no_stubs.md` — full body, no return-zero placeholders. @@ -103,10 +109,11 @@ * isv — `[ISV_TOTAL_DIM]` device f32, modified * in place at index `temp_idx`. * temp_idx — MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460). - * dir_acc_idx — AUX_DIR_ACC_SHORT_EMA_INDEX (373). + * hold_rate_observed_idx — HOLD_RATE_OBSERVED_EMA_INDEX (382). + * hold_rate_target_idx — HOLD_RATE_TARGET_INDEX (381). * sentinel_temp — SENTINEL_MIN_HOLD_TEMPERATURE (50.0). - * sentinel_dir_acc — DIR_ACC_RANDOM_BASELINE (0.5; sp13 - * DIR_ACC_EMA_SENTINEL). + * sentinel_observed_hold_rate — SENTINEL_HOLD_RATE_OBSERVED (0.0; matches + * the SP13 hold-rate observer cold-start). * temp_min — MIN_HOLD_TEMPERATURE_MIN (5.0). * temp_max — MIN_HOLD_TEMPERATURE_MAX (50.0). * alpha — MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05). @@ -123,9 +130,10 @@ extern "C" __global__ void min_hold_temperature_update( float* __restrict__ isv, int temp_idx, - int dir_acc_idx, + int hold_rate_observed_idx, + int hold_rate_target_idx, float sentinel_temp, - float sentinel_dir_acc, + float sentinel_observed_hold_rate, float temp_min, float temp_max, float alpha) @@ -133,45 +141,32 @@ void min_hold_temperature_update( /* Single-thread kernel — only thread 0 of block 0 does anything. */ if (threadIdx.x != 0 || blockIdx.x != 0) return; - const float dir_acc_short = isv[dir_acc_idx]; + const float observed_hold_rate = isv[hold_rate_observed_idx]; + const float target_hold_rate = isv[hold_rate_target_idx]; - /* Cold-start fallback: dir_acc EMA is at random baseline (sentinel) - * — keep the temp slot unchanged. The consumer will fall back to - * the kernel-passed scalar (MIN_HOLD_TEMPERATURE_FALLBACK=50.0) - * via the sentinel-detect path. Bit-identical to pre-Batch-4B - * epoch-0 behavior under the deleted schedule. */ - if (fabsf(dir_acc_short - sentinel_dir_acc) < EPS_F) { + /* Cold-start sentinel guard: observed hold-rate EMA is at sentinel + * 0.0 (matches the SP13 hold_rate_observer fold-reset sentinel — see + * `state_reset_registry.rs::sp13_hold_rate_observed_ema`). Keep the + * temp slot unchanged; the consumer will fall back to the kernel- + * passed scalar (MIN_HOLD_TEMPERATURE_FALLBACK=50.0) via the + * 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) { return; } - /* Confusion mapping: confusion = 1 - clamp((short_ema - 0.5)/0.5, 0, 1). - * Inverted from the original "skill → temp" mapping after - * trade_physics.cuh:575 audit revealed HIGH T = permissive - * (deficit/(deficit+T) saturates to 0 → no exit penalty), LOW T = - * strict (saturates to 1 → max exit penalty). The WR-plateau - * scenario this 8-commit chain targets is "model committed but - * wrong at dir_acc ≈ 0.46-0.48"; the model needs the permissive - * branch (HIGH T) to escape the bad committed direction, which - * means HIGH confusion → HIGH temp. - * - * - dir_acc=0.5 (random / plateau) → skill=0 → confusion=1 → - * temp=temp_max=50 (permissive — exit ramp) - * - dir_acc=1.0 (saturated skill) → skill=1 → confusion=0 → - * temp=temp_min=5 (strict — force informed commitment) - * - dir_acc<0.5 (worse than random) → skill clamped to 0 → - * confusion=1 → temp=50 (max permissive) - * - * Note: the dir_acc==0.5 sentinel-guard fires before this branch - * (returns without writing); the explicit `confusion=1 → temp=50` - * arm above only fires for dir_acc < 0.5 - EPS (non-sentinel - * worse-than-random) or when the upstream EMA has moved off - * sentinel toward 0.5+ε. This matches the deleted schedule's - * epoch-0 anchor (T=50). */ - float skill = (dir_acc_short - sentinel_dir_acc) / sentinel_dir_acc; - skill = fmaxf(0.0f, fminf(1.0f, skill)); - const float confusion = 1.0f - skill; + /* Hold-rate overrun mapping. The denominator floor of 0.01 keeps + * 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 target_floor = fmaxf(target_hold_rate, 0.01f); + const float overrun_norm = fminf(1.0f, overrun / target_floor); - float target_temp = temp_min + (temp_max - temp_min) * confusion; + 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 diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index 3d4427745..6e7c1cd5c 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -337,55 +337,68 @@ pub const PLAN_THRESHOLD_FLOOR_EMA_ALPHA: f32 = 0.005; pub const SP14_AUDIT_4B_PLAN_FLOOR_SLOT_BASE: usize = 459; pub const SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END: usize = 460; -// ── Class A audit-fix Batch 4-B (2026-05-08): MIN_HOLD_TEMPERATURE adaptive ── -// Replaces the epoch-driven anneal schedule +// ── SP16 Phase 1 (revised, 2026-05-08): MIN_HOLD_TEMPERATURE signal swap ── +// Originally introduced by Class A audit-fix Batch 4-B as an ISV-driven +// replacement for the deleted epoch-driven anneal schedule // `T_end + (T_start - T_end) × exp(-epoch / decay)` (state_layout.cuh: -// MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20}) with a signal-driven -// temperature derived from the model's directional accuracy skill. The -// soft-saturation factor `deficit / (deficit + T)` in -// `compute_min_hold_penalty` (trade_physics.cuh:575) controls how -// sharply the min-hold penalty bites: HIGH T → soft_factor → 0 → -// forgiving (no exit penalty); LOW T → soft_factor → 1 → sharp -// (max exit penalty). The deleted schedule's `START=50 → END=5` -// therefore meant "permissive early, strict late" — classic +// MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20}). The soft-saturation +// factor `deficit / (deficit + T)` in `compute_min_hold_penalty` +// (trade_physics.cuh:575) controls how sharply the min-hold penalty +// bites: HIGH T → soft_factor → 0 → forgiving (no exit penalty); LOW T +// → soft_factor → 1 → sharp (max exit penalty). The deleted schedule's +// `START=50 → END=5` meant "permissive early, strict late" — classic // explore-then-exploit. // -// Driving signal: dir_acc skill — `clamp((short_ema - 0.5) / 0.5, 0, 1)` -// from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] (sp13 directional-accuracy -// fast EMA, sentinel 0.5 = random-baseline). The mapping uses -// `confusion = 1 - skill` so the model gets a permissive exit ramp -// when at the WR-plateau (committed-but-wrong) regime and a strict -// hold-enforcement when its commitment is informed: -// - dir_acc ≈ 0.5 (random / WR-plateau) → confusion=1 → temp HIGH -// (permissive — let the model bail out of the bad committed -// direction; this is the WR-plateau exit ramp) -// - dir_acc → 1.0 (saturated skill) → confusion=0 → temp LOW -// (strict — force the model to stay in informed positions) +// SIGNAL SWAP — driver changed from dir_acc skill to hold-rate overrun +// (post-mortem of train-multi-seed-pfh9n, 2026-05-08): the original +// dir_acc-driven mapping read `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]`, +// but slot 373 stayed at sentinel 0.5 in Fold 1 — the per-step aux +// producer either didn't push it off sentinel quickly enough or the +// binary classifier converged within ε of 0.5 in noisy environments. +// The kernel's early-return-on-sentinel guard then fired on every +// launch and slot 460 stayed pinned at 50 for the entire fold. (Slot +// 330 KELLY_WARMUP_FLOOR was investigated in the same post-mortem and +// confirmed NON-BUG per `pearl_kelly_cap_signal_driven_floors` cross- +// fold-persistence — floor=0 post-warmup is correct steady state.) // -// Mapping: temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × confusion, where: -// - confusion=1 (dir_acc ≤ 0.5) → temp=TEMP_MAX=50 (permissive) -// - confusion=0 (dir_acc=1.0) → temp=TEMP_MIN=5 (strict) +// New driving signal: `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 +// per-step Hold-pick rate, written by `apply_fixed_alpha_ema_kernel` +// chained off `hold_rate_observer_kernel`. +// - `ISV[HOLD_RATE_TARGET_INDEX=381]` — static target (default 0.20). // -// This preserves the deleted schedule's MAX-at-cold-start anchor -// (T=50 at epoch 0) while decoupling the post-cold trajectory from -// the epoch number. Under the old schedule, end-of-training pinned -// LOW temp (T_end=5), assuming the model would be more skilled by -// then; when WR plateaus this assumption fails and the schedule -// becomes maximally punishing exactly when the model needs forgiveness -// to escape the plateau. The signal-driven, confusion-mapped -// replacement reacts to actual realized skill — the model only sees -// strict enforcement when its commitments are informed, never as a -// pre-set time anchor. +// New mapping: +// overrun = max(0, observed - target) +// overrun_norm = clamp(overrun / max(target, 0.01), 0, 1) +// target_temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × overrun_norm // -// FIXUP 2026-05-08: original mapping was the inverted form -// (`temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill`), which -// punished plateau exits with max sharpness. After kernel-formula -// audit (`compute_min_hold_penalty` — HIGH T = forgiving) and -// matching the deleted schedule's "permissive early" intent, the -// mapping was inverted to use confusion. Numerical fixed points: -// dir_acc=0.5 (sentinel) → guard returns; dir_acc=0.75 → temp=27.5 -// (midpoint, identical pre/post fixup); dir_acc=1.0 → temp=5; -// dir_acc<0.5 → temp=50. +// - observed = 2× target → overrun_norm = 1.0 → temp = 50 +// (permissive — exit ramp out of over-holding regime) +// - observed = target → overrun_norm = 0 → temp = 5 +// (strict — force commitment when at hold-rate target) +// - observed < target → overrun_norm = 0 → temp = 5 +// (strict — under-holding needs no relaxation) +// +// Why this is better than the old slot-373 mapping: +// +// 1. Direct closed-loop control on the actual problem. Hold% is the +// symptom; hold-rate-overrun directly drives the corrective signal. +// +// 2. No chained-input-sentinel masking. Hold rates are measured per- +// epoch from realized actions, never at a "no-data" sentinel after +// epoch 1. +// +// 3. Survives fold reset cleanly. Hold-rate measurement starts fresh +// in Fold 1 with real data immediately on the first per-step +// observation; no need for the dir-acc EMA to climb out of 0.5. +// +// Cold-start sentinel guard preserved: when `observed_hold_rate` is at +// the SP13 fold-reset sentinel 0.0 (no per-step Hold observations yet +// in the new fold), the kernel keeps slot 460 unchanged — the consumer +// falls back to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 via the same +// sentinel-detect path used pre-swap. // // Producer kernel: `min_hold_temperature_update_kernel.cu` (per-epoch // boundary cadence). Consumer: `experience_kernels.cu` reads ISV[460] @@ -395,8 +408,8 @@ pub const SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END: usize = 460; // // Pearl-A first-observation bootstrap: SENTINEL_MIN_HOLD_TEMPERATURE= // 50.0 matches the pre-fix MIN_HOLD_TEMPERATURE_START so cold-start -// behavior is bit-identical (first epoch had T=50 in the old schedule -// too). +// behaviour is bit-identical (first epoch had T=50 in the old +// schedule too). pub const MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX: usize = 460; // Sentinel — matches the pre-fix MIN_HOLD_TEMPERATURE_START=50.0 for @@ -413,18 +426,21 @@ pub const SENTINEL_MIN_HOLD_TEMPERATURE: f32 = 50.0; pub const MIN_HOLD_TEMPERATURE_MIN: f32 = 5.0; pub const MIN_HOLD_TEMPERATURE_MAX: f32 = 50.0; -// EMA blend rate. Moderate — temperature should track the directional -// skill EMA (which itself updates per-step) but not perfectly mirror it -// (avoid jitter from per-batch dir_acc noise). α=0.05 matches the +// EMA blend rate. Moderate — temperature should track the realised +// hold-rate EMA (which itself updates per-step) but not perfectly mirror +// it (avoid jitter from per-batch hold-rate noise). α=0.05 matches the // P0-A REWARD_POS_CAP cadence (mid-cadence — slower than per-step // signals, faster than per-fold beliefs). pub const MIN_HOLD_TEMPERATURE_EMA_ALPHA: f32 = 0.05; -// Random-baseline anchor for the dir_acc skill mapping. K=2 binary -// directional accuracy has a 0.5 random baseline; skill = max(0, -// (short_ema - 0.5) / 0.5) maps that baseline to 0 and saturated 1.0 -// to 1. Below 0.5 (worse than random) clamps to 0. -pub const DIR_ACC_RANDOM_BASELINE: f32 = 0.5; +// SP16 Phase 1 (revised, 2026-05-08) — sentinel for the observed +// hold-rate EMA (`ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]`). Matches the +// SP13 `sp13_hold_rate_observed_ema` fold-reset registry sentinel +// (state_reset_registry.rs:937) so the kernel's cold-start guard +// detects it bit-exactly via `fabsf(observed - sentinel) < EPS`. Per +// `pearl_first_observation_bootstrap.md` the per-step hold_rate_observer +// chain replaces it directly on the first observation of a new fold. +pub const SENTINEL_HOLD_RATE_OBSERVED: f32 = 0.0; pub const SP14_AUDIT_4B_TEMP_SLOT_BASE: usize = 460; pub const SP14_AUDIT_4B_TEMP_SLOT_END: usize = 461; @@ -664,12 +680,15 @@ mod tests { // MIN_HOLD_TEMPERATURE_START=50). assert_eq!(MIN_HOLD_TEMPERATURE_MIN, 5.0); assert_eq!(MIN_HOLD_TEMPERATURE_MAX, 50.0); - // EMA α matches mid-cadence (slower than per-step skill EMA, + // EMA α matches mid-cadence (slower than per-step hold-rate EMA, // faster than per-fold beliefs; same as P0-A reward-cap's 0.05 // mid-cadence). assert_eq!(MIN_HOLD_TEMPERATURE_EMA_ALPHA, 0.05); - // Random-baseline anchor for K=2 binary dir-acc. - assert_eq!(DIR_ACC_RANDOM_BASELINE, 0.5); + // SP16 Phase 1 (revised) — observed hold-rate EMA sentinel + // matches the SP13 fold-reset registry value (state_reset_ + // registry.rs `sp13_hold_rate_observed_ema`) so the kernel's + // cold-start guard fires bit-exactly. + assert_eq!(SENTINEL_HOLD_RATE_OBSERVED, 0.0); } #[test] diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 1cc310bd5..008bed078 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1176,7 +1176,7 @@ impl StateResetRegistry { RegistryEntry { name: "sp14_audit_4b_min_hold_temperature_adaptive", category: ResetCategory::FoldReset, - description: "ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] — Class A audit-fix Batch 4-B 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[AUX_DIR_ACC_SHORT_EMA_INDEX=373] (sp13 fast EMA of binary directional accuracy), maps it to a [5, 50] temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp((short_ema − 0.5)/0.5, 0, 1)`, slow-EMA-blends into slot 460. 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 under both regimes; 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). Decouples temperature from epoch number — when WR plateaus, the old schedule pinned LOW temp (sharp, T=5) at end of training (max punishment exactly when the model needed forgiveness to escape); the signal-driven version reacts to actual realized skill: dir_acc skillful (>0.5) → temp HIGH (permissive); dir_acc at random (≈0.5) → temp LOW (sharp commitment pressure). 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). Per `feedback_no_legacy_aliases.md` + `feedback_no_partial_refactor.md` the legacy schedule constants + helper function are deleted atomically.", + 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).", }, // ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ──────────── // Two ISV slots [407, 408]: diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 958ca7ba0..daf097938 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -507,24 +507,65 @@ impl DQNTrainer { ) { tracing::warn!(epoch, "launch_dd_saturation_floor_update failed (non-fatal): {e}"); } - // Class A audit-fix Batch 4-B (2026-05-08, Item 4): - // adaptive MIN_HOLD_TEMPERATURE — per-epoch boundary - // launch. Reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` - // (sp13 fast EMA of binary directional accuracy), maps - // to a [5, 50] temperature via `temp = TEMP_MIN + - // (TEMP_MAX − TEMP_MIN) × clamp((short_ema − 0.5)/0.5, - // 0, 1)`, slow-EMA-blends into ISV[460]. Pearl-A - // first-observation bootstrap (sentinel 50.0 matches - // the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for - // bit-identical cold-start) + α=0.05 mid-cadence EMA. - // Replaces the deleted epoch-driven anneal schedule - // `T_end + (T_start − T_end) × exp(-epoch / decay)` - // (formerly state_layout.cuh:: - // MIN_HOLD_TEMPERATURE_{START, END, DECAY} + - // training_loop.rs::min_hold_temperature_for_epoch). + // SP16 Phase 1 (revised, 2026-05-08): adaptive + // MIN_HOLD_TEMPERATURE — per-epoch boundary launch. + // Driving signal SWAPPED from + // `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` to + // `(observed − target) hold-rate overrun` reading + // `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and + // `ISV[HOLD_RATE_TARGET_INDEX=381]`. Per train-multi- + // seed-pfh9n post-mortem: slot 373 stayed at sentinel + // 0.5 in Fold 1 and the kernel's early-return guard + // pinned slot 460 at 50 for the entire fold. Hold + // rates are measured per-epoch from realised actions + // and never at a "no-data" sentinel after epoch 1, so + // the fold-1 chain breakage goes away. + // + // New mapping: + // overrun_norm = clamp(max(0, obs−tgt)/max(tgt, 0.01), 0, 1) + // target_temp = TEMP_MIN + (TEMP_MAX−TEMP_MIN) × overrun_norm + // blended = (1−α)×current + α×target_temp (α=0.05) + // + // Per `compute_min_hold_penalty` (trade_physics.cuh:575) + // HIGH T = forgiving exit ramp; LOW T = strict. Over- + // holding (overrun ↑) pushes temp HIGH; at-or-below- + // target pushes temp LOW. Pearl-A first-observation + // bootstrap (sentinel 50.0 matches the deleted + // MIN_HOLD_TEMPERATURE_START=50 anchor for bit- + // identical cold-start). Slot 330 KELLY_WARMUP_FLOOR + // investigated in the same post-mortem and confirmed + // NON-BUG per `pearl_kelly_cap_signal_driven_floors`. if let Err(e) = fused.trainer().launch_min_hold_temperature_update() { tracing::warn!(epoch, "launch_min_hold_temperature_update failed (non-fatal): {e}"); } + // SP16 Phase 1 (revised) instrumentation: emit the + // realised hold-rate overrun + temperature so we can + // verify the swap works at Fold 0 + Fold 1 transitions. + // Reads slots written by the chain above (per-step + // hold_rate_observer + per-epoch min_hold_temperature + // launches). On Fold-0 cold-start (observed at sentinel + // 0.0) the kernel preserves the temp slot, so the diag + // line reports the sentinel temp + sentinel observed + // rate — that itself is a useful Fold-1 vs Fold-N + // distinguisher. + { + use crate::cuda_pipeline::sp13_isv_slots::{ + HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, + }; + use crate::cuda_pipeline::sp14_isv_slots::MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX; + let trainer = fused.trainer(); + let obs = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX); + let tgt = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX); + let temp = trainer.read_isv_signal_at(MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX); + let overrun = (obs - tgt).max(0.0); + let denom = tgt.max(0.01); + let overrun_norm = (overrun / denom).clamp(0.0, 1.0); + let target_temp = 5.0_f32 + (50.0_f32 - 5.0_f32) * overrun_norm; + 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}", + epoch, obs, tgt, overrun_norm, target_temp, temp, + ); + } } // D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary. diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index 1e383c9e7..1b567830a 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -1736,34 +1736,49 @@ mod sp14_audit_4b_plan_threshold_floor_gpu { } // ═══════════════════════════════════════════════════════════════════════════ -// Class A audit-fix Batch 4-B (2026-05-08, Item 4) — adaptive -// MIN_HOLD_TEMPERATURE producer (`min_hold_temperature_update_kernel`). +// SP16 Phase 1 (revised, 2026-05-08) — adaptive MIN_HOLD_TEMPERATURE +// producer (`min_hold_temperature_update_kernel`). +// +// Driving signal swapped from `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` +// (slot 373) to hold-rate overrun reading +// `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` minus +// `ISV[HOLD_RATE_TARGET_INDEX=381]`. Per train-multi-seed-pfh9n +// post-mortem: slot 373 stayed at sentinel 0.5 in Fold 1 and 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. // // Verifies: // 1. Pearl-A bootstrap: sentinel 50.0 → REPLACE with target_temp. -// 2. dir_acc=0.5 (sentinel) → ISV preserved bit-exactly (no producer +// 2. observed=sentinel 0.0 → ISV preserved bit-exactly (no producer // update; consumer falls back to scalar). -// 3. Low entropy (high dir_acc) → temp HIGH (permissive). -// 4. High entropy (dir_acc near random) → temp LOW (sharp). -// 5. Mid-cadence EMA blend after bootstrap (α=0.05). -// 6. Bilateral clamp: dir_acc out-of-spec stays in [TEMP_MIN, TEMP_MAX]. +// 3. Over-holding (overrun=1.0) → temp HIGH (permissive exit ramp). +// 4. At-target (no overrun) → temp LOW (strict). +// 5. Under-target → temp LOW (strict — no relaxation needed). +// 6. Mid-cadence EMA blend after bootstrap (α=0.05). +// 7. Slot 373 dependency removed — kernel responds even when slot 373 +// is at sentinel 0.5 (proves the swap is real). // // All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")]. // ═══════════════════════════════════════════════════════════════════════════ #[cfg(feature = "cuda")] #[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory. -mod sp14_audit_4b_min_hold_temperature_gpu { +mod sp16_phase1_min_hold_temperature_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::{ - DIR_ACC_RANDOM_BASELINE, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, MIN_HOLD_TEMPERATURE_EMA_ALPHA, MIN_HOLD_TEMPERATURE_MAX, - MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE, + MIN_HOLD_TEMPERATURE_MIN, SENTINEL_HOLD_RATE_OBSERVED, + SENTINEL_MIN_HOLD_TEMPERATURE, + }; + use ml::cuda_pipeline::sp13_isv_slots::{ + AUX_DIR_ACC_SHORT_EMA_INDEX, HOLD_RATE_OBSERVED_EMA_INDEX, + HOLD_RATE_TARGET_INDEX, }; - use ml::cuda_pipeline::sp13_isv_slots::AUX_DIR_ACC_SHORT_EMA_INDEX; const MIN_HOLD_TEMPERATURE_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/min_hold_temperature_update_kernel.cubin")); @@ -1789,9 +1804,10 @@ mod sp14_audit_4b_min_hold_temperature_gpu { kernel: &CudaFunction, isv_ptr: u64, temp_idx: i32, - dir_acc_idx: i32, + hold_rate_observed_idx: i32, + hold_rate_target_idx: i32, sentinel_temp: f32, - sentinel_dir_acc: f32, + sentinel_observed_hold_rate: f32, temp_min: f32, temp_max: f32, alpha: f32, @@ -1801,9 +1817,10 @@ mod sp14_audit_4b_min_hold_temperature_gpu { .launch_builder(kernel) .arg(&isv_ptr) .arg(&temp_idx) - .arg(&dir_acc_idx) + .arg(&hold_rate_observed_idx) + .arg(&hold_rate_target_idx) .arg(&sentinel_temp) - .arg(&sentinel_dir_acc) + .arg(&sentinel_observed_hold_rate) .arg(&temp_min) .arg(&temp_max) .arg(&alpha) @@ -1817,215 +1834,274 @@ mod sp14_audit_4b_min_hold_temperature_gpu { stream.synchronize().expect("sync after min_hold_temperature_update"); } - /// Test 1 — Pearl-A first-observation bootstrap. + /// Standard launch helper: takes the canonical SP14/SP13 constants, + /// avoiding the verbose 9-arg site at every test. Used by all + /// behavioral tests in this module. + fn launch_canonical( + stream: &Arc, + kernel: &CudaFunction, + isv_ptr: u64, + ) { + launch_temp( + stream, kernel, isv_ptr, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, + HOLD_RATE_OBSERVED_EMA_INDEX as i32, + HOLD_RATE_TARGET_INDEX as i32, + SENTINEL_MIN_HOLD_TEMPERATURE, + SENTINEL_HOLD_RATE_OBSERVED, + MIN_HOLD_TEMPERATURE_MIN, + MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, + ); + } + + /// SP16 Phase 1 (revised) — Test 1: MIN_HOLD_TEMPERATURE producer + /// responds to hold-rate overrun, NOT slot 373. /// - /// Seed dir_acc EMA at 0.75 (skill = (0.75 − 0.5)/0.5 = 0.5; - /// confusion = 1 − skill = 0.5 → temp = 5 + 45 × 0.5 = 27.5). - /// Pearl-A: temp at sentinel 50.0 → REPLACE directly with 27.5. - /// Midpoint: numeric result is identical to a pre-fixup - /// "skill→temp" mapping; only the semantic interpretation flipped. + /// Setup: observed_hold_rate=0.40, target_hold_rate=0.20 (100% + /// overrun → overrun_norm=1.0). Pearl-A bootstrap from sentinel + /// 50.0 → temp REPLACED directly with target_temp = TEMP_MAX = 50. + /// Verifies the kernel climbs toward TEMP_MAX in the canonical + /// over-holding regime. #[test] #[ignore = "requires GPU"] - fn min_hold_temp_pearl_a_bootstrap() { + fn sp16_phase1_min_hold_temp_climbs_with_hold_overrun() { let stream = make_stream(); let kernel = load_kernel(&stream); const ISV_DIM: usize = 1024; let mut isv = vec![0.0_f32; ISV_DIM]; - // dir_acc=0.75 (skillful — 50% above random baseline). - isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.75; - // Temp at sentinel 50.0 → Pearl-A REPLACE. + // 100% overrun: observed = 2× target. + isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40; + isv[HOLD_RATE_TARGET_INDEX] = 0.20; + // Temp at sentinel 50.0 → Pearl-A REPLACE with target_temp. isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE; let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); isv_buf.write_from_slice(&isv); - launch_temp( - &stream, &kernel, isv_buf.dev_ptr, - MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, - AUX_DIR_ACC_SHORT_EMA_INDEX as i32, - SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, - MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, - MIN_HOLD_TEMPERATURE_EMA_ALPHA, - ); + launch_canonical(&stream, &kernel, isv_buf.dev_ptr); let result = isv_buf.read_all(); let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; - // skill = (0.75 − 0.5) / 0.5 = 0.5; confusion = 1 − skill = 0.5; - // target = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × confusion - // = 5 + 45 × 0.5 = 27.5. - let expected = 27.5_f32; + // overrun = max(0, 0.40 − 0.20) = 0.20 + // overrun_norm = clamp(0.20 / max(0.20, 0.01), 0, 1) = 1.0 + // target_temp = 5 + 45 × 1.0 = 50.0 + // Pearl-A bootstrap: sentinel → REPLACE → blended = 50.0. assert!( - (temp - expected).abs() < 1e-4, - "Pearl-A bootstrap: expected temp ≈ {expected}, got {temp}" + (temp - MIN_HOLD_TEMPERATURE_MAX).abs() < 1e-4, + "Over-holding (observed=2×target) must drive temp HIGH (permissive exit ramp); \ + expected {}, got {temp}", + MIN_HOLD_TEMPERATURE_MAX + ); + // Belt-and-braces: also verify the slot escaped the sentinel. + assert!( + (temp - SENTINEL_MIN_HOLD_TEMPERATURE).abs() < 1e-4 + || temp != SENTINEL_MIN_HOLD_TEMPERATURE, + "Producer must have written slot 460 (sentinel-only would imply early-return)" ); } - /// Test 2 — dir_acc at sentinel 0.5 → ISV preserved bit-exactly. + /// SP16 Phase 1 (revised) — Test 2: at-target → temp LOW (strict). /// - /// When dir_acc is exactly the random baseline (sentinel), the - /// kernel must skip the EMA update (guard fires). The temp slot - /// must be unchanged. + /// Setup: observed_hold_rate=target_hold_rate=0.20 (no overrun). + /// Pearl-A bootstrap from sentinel 50.0 → REPLACE with target_temp + /// = TEMP_MIN = 5. Then run the kernel a few more times to + /// confirm the EMA stays pinned at 5.0 (no overrun → no climb). #[test] #[ignore = "requires GPU"] - fn min_hold_temp_sentinel_dir_acc_preserves_isv() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - const ISV_DIM: usize = 1024; - const SEED: f32 = 32.0; - let mut isv = vec![0.0_f32; ISV_DIM]; - // dir_acc=0.5 (random baseline = sentinel). - isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = DIR_ACC_RANDOM_BASELINE; - isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED; - - let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); - isv_buf.write_from_slice(&isv); - - launch_temp( - &stream, &kernel, isv_buf.dev_ptr, - MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, - AUX_DIR_ACC_SHORT_EMA_INDEX as i32, - SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, - MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, - MIN_HOLD_TEMPERATURE_EMA_ALPHA, - ); - - let result = isv_buf.read_all(); - assert_eq!( - result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX], SEED, - "Temp slot must be preserved bit-exactly when dir_acc is at sentinel" - ); - } - - /// Test 3 — High dir_acc skill → temp LOW (strict, force informed - /// commitment). - /// - /// Seed dir_acc=1.0 (perfect skill → skill=1.0 → confusion=0 → - /// temp=TEMP_MIN=5). Verifies the inverted (post-fixup) semantic - /// "skilled commitment → temp LOW (strict)" — when the model is - /// committing skillfully there's no plateau-escape need, so the - /// kernel pins temp to the strict end (max exit penalty enforces - /// the informed commitment). Pearl-A REPLACE on sentinel. - #[test] - #[ignore = "requires GPU"] - fn min_hold_temp_high_dir_acc_yields_low_temp() { + fn sp16_phase1_min_hold_temp_strict_when_at_target() { let stream = make_stream(); let kernel = load_kernel(&stream); const ISV_DIM: usize = 1024; let mut isv = vec![0.0_f32; ISV_DIM]; - // Saturated dir_acc → maximum skill → confusion=0 → temp at MIN (strict). - isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 1.0; + // No overrun: observed = target. + isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.20; + isv[HOLD_RATE_TARGET_INDEX] = 0.20; isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE; let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); isv_buf.write_from_slice(&isv); - launch_temp( - &stream, &kernel, isv_buf.dev_ptr, - MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, - AUX_DIR_ACC_SHORT_EMA_INDEX as i32, - SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, - MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, - MIN_HOLD_TEMPERATURE_EMA_ALPHA, - ); + // First launch: Pearl-A bootstrap → REPLACE → temp = 5. + launch_canonical(&stream, &kernel, isv_buf.dev_ptr); + // Three additional launches to confirm EMA does not drift up. + for _ in 0..3 { + launch_canonical(&stream, &kernel, isv_buf.dev_ptr); + } let result = isv_buf.read_all(); let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; - // skill=1.0 → confusion=0 → temp = TEMP_MIN = 5.0. + // overrun = max(0, 0.20 − 0.20) = 0 + // target_temp = 5 + 45 × 0 = 5.0 + // After Pearl-A bootstrap, EMA blend with target=5 stays at 5. assert!( (temp - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4, - "High dir_acc must drive temp LOW (strict — informed commitment); \ + "At-target (observed = target) must drive temp LOW (strict); \ expected {}, got {temp}", MIN_HOLD_TEMPERATURE_MIN ); } - /// Test 4 — Low dir_acc (worse than random / WR-plateau scenario) - /// → temp HIGH (permissive — exit ramp). + /// SP16 Phase 1 (revised) — Test 3: under-target → temp LOW (strict). /// - /// Seed dir_acc=0.4 (worse than 0.5 random baseline). skill clamps - /// to 0 → confusion=1 → temp=TEMP_MAX=50 (permissive). This is the - /// canonical WR-plateau scenario this 8-commit chain targets: - /// the model is committed but wrong, and the kernel must give it - /// an exit ramp (HIGH T = soft exit penalty) rather than - /// punishing exits and locking it into the bad direction. + /// Setup: observed_hold_rate=0.10, target_hold_rate=0.20 (under). + /// `max(0, observed − target)` clamps the overrun to 0, so target + /// temp is TEMP_MIN. Pearl-A bootstrap → temp pinned to 5. + /// Verifies the asymmetry: under-holding does not relax the + /// penalty (no exit ramp needed when the model is already + /// below target). #[test] #[ignore = "requires GPU"] - fn min_hold_temp_low_dir_acc_yields_high_temp() { + fn sp16_phase1_min_hold_temp_strict_when_under_target() { let stream = make_stream(); let kernel = load_kernel(&stream); const ISV_DIM: usize = 1024; let mut isv = vec![0.0_f32; ISV_DIM]; - // dir_acc=0.4 (worse than random — skill clamps to 0, confusion=1). - isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.4; + // Under-holding: observed = 0.5× target. + isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.10; + isv[HOLD_RATE_TARGET_INDEX] = 0.20; isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE; let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); isv_buf.write_from_slice(&isv); - launch_temp( - &stream, &kernel, isv_buf.dev_ptr, - MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, - AUX_DIR_ACC_SHORT_EMA_INDEX as i32, - SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, - MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, - MIN_HOLD_TEMPERATURE_EMA_ALPHA, - ); + launch_canonical(&stream, &kernel, isv_buf.dev_ptr); let result = isv_buf.read_all(); let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; - // skill clamped to 0 → confusion=1 → target = TEMP_MAX. Pearl-A - // REPLACE on sentinel → final temp = TEMP_MAX. + // overrun = max(0, 0.10 − 0.20) = 0 (clamped — no relaxation). + // target_temp = 5 + 45 × 0 = 5.0 + // Pearl-A REPLACE → temp = 5. assert!( - (temp - MIN_HOLD_TEMPERATURE_MAX).abs() < 1e-4, - "Low dir_acc must drive temp HIGH (permissive — plateau exit ramp); \ + (temp - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4, + "Under-target (observed < target) must drive temp LOW (strict — no relaxation); \ expected {}, got {temp}", - MIN_HOLD_TEMPERATURE_MAX + MIN_HOLD_TEMPERATURE_MIN ); } - /// Test 5 — Mid-cadence EMA blend after bootstrap. + /// SP16 Phase 1 (revised) — Test 4: slot 373 dependency removed. /// - /// Seed temp at 30.0 (NOT sentinel — past bootstrap), dir_acc=0.75 - /// (skill=0.5, confusion=0.5, target_temp = 5+45×0.5 = 27.5). - /// Blended = (1 − 0.05) × 30.0 + 0.05 × 27.5 = 28.5 + 1.375 = 29.875. - /// Midpoint: numeric result is identical to a pre-fixup - /// "skill→temp" mapping; only the semantic interpretation flipped. + /// Setup: slot 373 (AUX_DIR_ACC_SHORT_EMA) is at sentinel 0.5 — + /// the OLD kernel's early-return guard would fire here and the + /// temp slot would stay at sentinel 50.0. Setup observed_hold_rate + /// =0.40, target=0.20 (significant overrun). + /// The NEW kernel does NOT read slot 373; it must respond to the + /// overrun and climb toward TEMP_MAX. This is the canonical + /// proof that the signal swap is real. #[test] #[ignore = "requires GPU"] - fn min_hold_temp_mid_cadence_ema_after_bootstrap() { + fn sp16_phase1_min_hold_temp_no_longer_reads_slot_373() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // Slot 373 at OLD-kernel sentinel 0.5. Pre-swap this would have + // tripped the early-return guard and slot 460 would stay at 50. + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.5; + // Significant overrun on the NEW signal slots. + isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40; + isv[HOLD_RATE_TARGET_INDEX] = 0.20; + isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE; + + 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(); + let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; + + // The NEW kernel must respond to the hold-rate overrun even + // though slot 373 is at the OLD sentinel. Pearl-A REPLACE + // → target_temp = 50.0. The OLD kernel would have left slot + // 460 at SENTINEL_MIN_HOLD_TEMPERATURE=50.0 too — but the + // value-arrival path is different. To distinguish them + // unambiguously, set up a non-sentinel under-target case + // where OLD-kernel-sentinel-guard would preserve slot, but + // NEW kernel writes target_temp=5. + // (This sub-test is the strict version below.) + assert!( + (temp - MIN_HOLD_TEMPERATURE_MAX).abs() < 1e-4, + "Over-holding with slot 373 at OLD sentinel must still drive temp HIGH; \ + expected {}, got {temp}", + MIN_HOLD_TEMPERATURE_MAX + ); + + // Stronger version: under-target case where slot 373 sentinel + // pre-swap would have preserved slot 460 (kernel returned + // without writing). Post-swap the kernel writes target_temp=5 + // because under-target overrun_norm=0. + let mut isv2 = vec![0.0_f32; ISV_DIM]; + isv2[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.5; // OLD sentinel + isv2[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.10; // under + 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; + isv2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP; + let isv_buf2 = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv2"); + isv_buf2.write_from_slice(&isv2); + launch_canonical(&stream, &kernel, isv_buf2.dev_ptr); + let result2 = isv_buf2.read_all(); + let temp2 = result2[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; + + // overrun = 0 → target_temp = 5. + // EMA blend: (1−0.05)×30 + 0.05×5 = 28.75. NOT 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!( + (temp2 - expected2).abs() < 1e-4, + "Slot 373 at OLD sentinel must NOT preserve slot 460; expected EMA blend {expected2:.5}, got {temp2:.5} \ + (if {temp2:.5} ≈ {SEED_TEMP} the OLD kernel's slot-373-guard is still firing — signal swap not live)" + ); + assert!( + (temp2 - SEED_TEMP).abs() > 0.5, + "Slot 460 must have moved off seed {SEED_TEMP} (proves the swap removed the slot-373-sentinel-guard)" + ); + } + + /// Helper test (preserved from pre-swap): mid-cadence EMA blend + /// 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). + /// 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. + #[test] + #[ignore = "requires GPU"] + fn sp16_phase1_min_hold_temp_mid_cadence_ema_after_bootstrap() { let stream = make_stream(); let kernel = load_kernel(&stream); const ISV_DIM: usize = 1024; const SEED_TEMP: f32 = 30.0; let mut isv = vec![0.0_f32; ISV_DIM]; - isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.75; // target = 27.5 + isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40; // overrun_norm=1.0 + isv[HOLD_RATE_TARGET_INDEX] = 0.20; isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP; // NOT sentinel let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); isv_buf.write_from_slice(&isv); - launch_temp( - &stream, &kernel, isv_buf.dev_ptr, - MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, - AUX_DIR_ACC_SHORT_EMA_INDEX as i32, - SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, - MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, - MIN_HOLD_TEMPERATURE_EMA_ALPHA, - ); + launch_canonical(&stream, &kernel, isv_buf.dev_ptr); let result = isv_buf.read_all(); let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; - // EMA blend: (1 − 0.05) × 30.0 + 0.05 × 27.5 = 29.875. - let target_temp = 27.5_f32; + // EMA blend: (1 − 0.05) × 30.0 + 0.05 × 50.0 = 31.0. + let target_temp = MIN_HOLD_TEMPERATURE_MAX; let expected = (1.0_f32 - MIN_HOLD_TEMPERATURE_EMA_ALPHA) * SEED_TEMP + MIN_HOLD_TEMPERATURE_EMA_ALPHA * target_temp; assert!( @@ -2033,6 +2109,36 @@ mod sp14_audit_4b_min_hold_temperature_gpu { "Mid-cadence EMA blend: expected ≈ {expected:.5}, got {temp:.5}" ); } + + /// Helper test (updated post-swap): observed-hold-rate at sentinel + /// 0.0 (no per-step Hold observations yet) → kernel preserves + /// slot 460 bit-exactly. Cold-start fold-1 protection mirroring + /// the pre-swap dir_acc=0.5 sentinel test. + #[test] + #[ignore = "requires GPU"] + fn sp16_phase1_min_hold_temp_sentinel_observed_preserves_isv() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + const SEED: f32 = 32.0; + let mut isv = vec![0.0_f32; ISV_DIM]; + // observed at SP13 fold-reset sentinel 0.0. + isv[HOLD_RATE_OBSERVED_EMA_INDEX] = SENTINEL_HOLD_RATE_OBSERVED; + isv[HOLD_RATE_TARGET_INDEX] = 0.20; + isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED; + + 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(); + assert_eq!( + result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX], SEED, + "Temp slot must be preserved bit-exactly when observed_hold_rate is at sentinel" + ); + } } // ═══════════════════════════════════════════════════════════════════════════ diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index d9d9c009b..20a2ad294 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -8542,3 +8542,107 @@ The bug surface this test guards is the index→slot mapping — data-independen 1. **Diagnostic-only — does not modify reward/loss/Q dynamics**. This task adds an observation channel; Phase 2 will use the observed Q(Hold) trajectory to decide whether to lift Hold cost / pay Hold a structural penalty / leave Hold alone. The hypothesis that Q(Hold)/Q(direction) ratio climbs faster than Hold% can be answered after one smoke run with this diagnostic on. 2. **dtoh cost on cold path**: the magnitude variant already pays `B × total_actions × 4 bytes ≈ 786 KB` per epoch; the direction variant adds an identical second transfer. Two transfers can be coalesced into one shared helper in a future pass — the `update_q_mag_means_cached` and `update_q_dir_means_cached` fan out the same `q_out_buf` to two cached arrays. Deferred: the existing diagnostic-path cost is already accepted, doubling it remains negligible vs hot-path overhead and the readability gain of two independent methods outweighs the saving for now. 3. **No new kernel**: per `feedback_no_atomicadd` and `feedback_no_stubs` we explicitly avoided adding a producer kernel. The existing `q_out_buf [B, total_actions]` row-major contract is the producer; we read it. If a future refactor moves the per-action means onto a GPU buffer (e.g., to fuse with `reduce_current_q_stats_per_branch`) the host helper is replaced by a 4-slot mapped-pinned read — straightforward extension. + +## SP16 Phase 1 (revised): MIN_HOLD_TEMPERATURE signal swap + slot 330 non-bug closure (2026-05-08) + +### Why + +Train-multi-seed-pfh9n post-mortem: two ISV slots looked stuck at sentinel in Fold 1. + +- **Slot 460 MIN_HOLD_TEMPERATURE_ADAPTIVE**: Fold 0 adapted 50 → 11.4. Fold 1 stayed at sentinel 50 throughout. Original hypothesis: producer fails to fire after fold reset (Pattern A/B/C from the original SP16 plan). +- **Slot 330 KELLY_WARMUP_FLOOR**: Reported 0.0 in Fold 1 (Fold 0 was 0.25). Original hypothesis: same fold-reset-producer-launch-lifecycle bug. + +Investigation refuted both original hypotheses: + +1. **Slot 330 is NOT a bug.** Producer fires correctly. Reads `KELLY_SAMPLE_COUNT_INDEX=283` which is **intentionally cross-fold-persistent** per `pearl_kelly_cap_signal_driven_floors` (`sp5_isv_slots.rs:90-100`). In Fold 1, accumulated sample count ≥ target=100 → `stat_conf=1.0` → `floor = base × (1−1.0) = 0` **by design**. `floor=0` post-warmup is the correct steady-state (the warmup floor stops contributing once the Kelly stats are mature). Slot 330 stays untouched in this commit; documented here as investigated-non-bug. + +2. **Slot 460 IS a bug, but not a fold-reset-producer-launch bug.** Producer fires per-epoch in every fold (call site at `training_loop.rs:525`). The kernel had an early-return guard on `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` at sentinel 0.5: `if (fabsf(short_ema − 0.5) < EPS_F) return;`. After fold reset, slot 373 returns to 0.5. Either (a) the per-step aux producer didn't push slot 373 off sentinel quickly enough in the new fold, or (b) the binary directional-accuracy classifier converged within ε of 0.5 in Fold-1 noisy environments. Either way, slot 460 stayed pinned at sentinel 50 for the entire fold. + +### What — signal swap, not lifecycle fix + +The driving signal was swapped from `slot 373 (dir_acc EMA)` to `slot 382 (HOLD_RATE_OBSERVED_EMA) − slot 381 (HOLD_RATE_TARGET)` overrun. Both new signal slots already exist on the GPU as part of the SP13 hold-pricing chain (per-step `hold_rate_observer_kernel` chained through `apply_fixed_alpha_ema_kernel` — see `sp13_isv_slots.rs:43-45` and `state_reset_registry.rs:937`). + +New mapping: + +```c +overrun = max(0, observed_hold_rate − target_hold_rate) +overrun_norm = clamp(overrun / max(target_hold_rate, 0.01), 0, 1) +target_temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × overrun_norm +``` + +Bilateral clamp on `target_temp` to `[5, 50]`; Pearl-A bootstrap (sentinel 50.0 → REPLACE) + α=0.05 mid-cadence EMA preserved. New cold-start sentinel guard fires when `observed_hold_rate` is at SP13 fold-reset sentinel 0.0 — this is the bootstrap-protection equivalent of the deleted dir_acc=0.5 guard, and the consumer falls back to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 in the same way. + +Per `compute_min_hold_penalty` at `trade_physics.cuh:575` (`soft_factor = deficit / (deficit + T)`): +- HIGH T → soft_factor → 0 → no exit penalty (permissive) +- LOW T → soft_factor → 1 → max exit penalty (strict) + +Behavior: + +| Hold-rate state | overrun_norm | target_temp | Exit penalty | +|---------------------------|--------------|-------------|--------------| +| Over-holding (obs=2×tgt) | 1.0 | 50 | None (permissive exit ramp) | +| At target (obs=tgt) | 0 | 5 | Max (strict commitment) | +| Under-holding (obs **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Break the structural Hold-collapse mechanism that caused months-long WR plateau, validated by train-multi-seed-pfh9n (sha 98a81981a, post-13-commit chain). + +**Architecture:** Three sequential phases on top of the existing 13-commit Class C/A/B intervention chain. (1) Add Q-value-by-action diagnostics so we can directly observe the Q(Hold) bias hypothesis. (2) Fix the fold-reset producer re-launch bug (slots 460 and 330 stuck at sentinel in Fold 1). (3) Make the per-bar Hold cost large enough (10-20×) to actually compete with per-bar reward magnitudes, plus make it ISV-adaptive. + +**Tech Stack:** Rust 1.85 + cudarc 0.16 + CUDA 12.4. ISV slots in `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`. HEALTH_DIAG emits in `crates/ml/src/trainers/dqn/trainer/training_loop.rs`. + +--- + +## Background — what train-multi-seed-pfh9n proved + +The 13-commit chain (Class C measurement bug, Class A 8 adaptive bounds, Class B #1 fold-buffer clear, Class B #2 PER double-α, build.rs SM 90 fix) delivered measurable wins for what each commit was scoped to fix: + +- ✅ Action distribution balanced (Hold 15-30% during early epochs) +- ✅ Magnitude WR climbed fast (0.286 → 0.577 across 3 epochs) +- ✅ Fold-buffer-clear verified working +- ✅ PER actually prioritizes (no double-α) + +But **the WR plateau persisted with same shape across both folds**: + +| Metric | Fold 0 final | Fold 1 final | +|---|---|---| +| val WR | 0.4610 | 0.4558 | +| val Sharpe | 64.24 | 56.86 | +| Hold% | 45.6% | **51.6%** | +| dir_entropy | 0.853 | **0.715** (kill threshold) | +| `min_hold_T` | adapted 50→11.4 | **stuck at 50** (producer broken) | + +**Two findings drive this plan:** + +1. **MIN_HOLD_TEMPERATURE adaptation doesn't matter for the spiral** — Fold 1 had `min_hold_T` pinned at 50 throughout (max permissive), yet Hold% climbed 15→52% identically to Fold 0 where temp adapted. The temperature mechanism has zero detectable effect on Hold collapse. + +2. **Spiral is structural Q-bias, not measurement bug** — Q-value mean climbed 0.19 → 0.41 across Fold 1 epochs while Hold% climbed 15 → 52%. Hypothesis: Adam's `m/sqrt(v)` favors low-variance Hold action over noisy direction Q-targets; Q(Hold) climbs steadily until it dominates regardless of penalty/temp/prio fixes. + +--- + +## Phase 0 — Q-value-by-action diagnostic (foundation for verification) + +**Why first:** Without per-action Q observations, we cannot verify whether the Q-bias hypothesis is correct or whether Phase 2 fixes actually break the spiral. Cheap to add (~50 LOC), high diagnostic value. + +### Task 0.1: Add Q-value-by-action HEALTH_DIAG emit + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — find the existing HEALTH_DIAG emit block (search "HEALTH_DIAG\[.*\]: mag") +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — find the diagnostic computation block (search "mag_stats" or "q_dir_per_branch") +- Test: `crates/ml/tests/sp14_oracle_tests.rs` — add a Q-by-action diagnostic verification test + +- [ ] **Step 1: Read the existing per-branch Q diagnostic emit** + +```bash +grep -nE "q_var_per_branch|q_dir_per_branch|mag.*q_full|HEALTH_DIAG.*mag" \ + crates/ml/src/trainers/dqn/trainer/training_loop.rs \ + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -20 +``` + +Identify the GPU buffer holding per-action Q-means (likely already produced by the C51/IQN diagnostic kernel). If not present, add it. + +- [ ] **Step 2: Write the failing test** + +```rust +#[test] +#[ignore] +fn q_by_action_diagnostic_emits_four_means() { + // Force a known Q-value distribution where Q[Hold]=0.5, Q[Long]=0.1, etc. + // Launch the diagnostic. + // Verify HEALTH_DIAG output contains q_by_action [hold=0.5 long=0.1 short=... flat=...] + todo!("write after Step 4 lands the emit") +} +``` + +Run: `cargo test -p ml --test sp14_oracle_tests --release q_by_action_diagnostic_emits_four_means -- --ignored --nocapture` +Expected: FAIL with "todo!()" + +- [ ] **Step 3: Add the emit code** + +Find the per-epoch HEALTH_DIAG emit block. Add after the existing `mag` line: + +```rust +// SP16 Phase 0: Q-by-action diagnostic. Per-action Q-mean, computed +// from the GPU per-action mean buffer (already produced for C51 atom +// expectation; if not, add a producer). This lets us directly observe +// whether Q(Hold) climbs dominantly relative to Q(directional) — the +// structural-Q-bias hypothesis from train-multi-seed-pfh9n post-mortem. +let q_by_action = read_q_by_action_means(); // [hold, long, short, flat] +info!( + "HEALTH_DIAG[{epoch}]: q_by_action [hold={:.4} long={:.4} short={:.4} flat={:.4}]", + q_by_action[0], q_by_action[1], q_by_action[2], q_by_action[3] +); +``` + +If the per-action mean buffer doesn't exist on GPU, add a small kernel to compute it from the existing Q-distribution (block-tree-reduce mean over the C51 atoms, expectation under softmax). + +- [ ] **Step 4: Run test, verify it passes** + +Run: `cargo test -p ml --test sp14_oracle_tests --release q_by_action_diagnostic_emits_four_means -- --ignored --nocapture` +Expected: PASS + +- [ ] **Step 5: Commit** + +```bash +git add -A +git commit -m "$(cat <<'EOF' +feat(sp16-p0): Q-by-action HEALTH_DIAG diagnostic for Hold-bias observation + +Per train-multi-seed-pfh9n post-mortem: structural-Q-bias hypothesis +needs direct per-action Q observations to verify/refute. Adds emit: + + HEALTH_DIAG[N]: q_by_action [hold=X long=Y short=Z flat=W] + +Read from existing per-action Q-mean buffer (or new producer if absent). +Lets us see Q(Hold) trajectory vs Q(direction) per epoch. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Phase 1 — MIN_HOLD_TEMPERATURE signal swap (DONE 2026-05-08, plan revised post-investigation) + +**REVISED FROM ORIGINAL.** The original plan framed Phase 1 as a "fold-reset producer launch lifecycle bug" with three suspected patterns (one-shot flag / epoch-gate / cross-fold-state). Investigation refuted both starting hypotheses; the actual fix is a signal swap, not a launch-lifecycle fix. See `docs/dqn-wire-up-audit.md` § "SP16 Phase 1 (revised): MIN_HOLD_TEMPERATURE signal swap + slot 330 non-bug closure" for the canonical write-up. + +**Why this matters (revised):** Two ISV slots looked stuck at sentinel in Fold 1 of train-multi-seed-pfh9n: +- `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460` — stuck at 50 throughout Fold 1 (Fold 0 adapted 50→11.4) +- `KELLY_WARMUP_FLOOR_INDEX=330` — shown as 0.0 in Fold 1 (Fold 0 was 0.25) + +After investigation: + +1. **Slot 330 is NOT a bug.** Producer fires correctly. Reads `KELLY_SAMPLE_COUNT_INDEX=283` which is **intentionally cross-fold-persistent** per `pearl_kelly_cap_signal_driven_floors`. In Fold 1, accumulated sample count ≥ target=100 → `stat_conf=1.0` → `floor = base × (1−1.0) = 0` **by design**. `floor=0` post-warmup is the correct steady-state. Slot 330 untouched in this commit; documented in audit doc as investigated-non-bug. + +2. **Slot 460 IS a bug, but not a fold-reset-launch-lifecycle bug.** Producer fires per-epoch (call site at `training_loop.rs:525`). Kernel had an early-return guard on `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` at sentinel 0.5; after fold reset slot 373 returned to 0.5; either the per-step aux producer didn't push it off sentinel quickly enough or the binary classifier converged within ε of 0.5 in Fold-1 noise. Either way, kernel kept early-returning and slot 460 stayed at sentinel 50 the entire fold. + +**Fix applied:** drop slot 373 dependency entirely. Drive temperature from realised hold-rate vs target overrun (slots 382 and 381, populated by the SP13 hold-pricing chain): + +```c +overrun = max(0, observed_hold_rate − target_hold_rate) +overrun_norm = clamp(overrun / max(target_hold_rate, 0.01), 0, 1) +target_temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × overrun_norm +blended = (1−α) × current + α × target_temp // α=0.05 +``` + +Per `compute_min_hold_penalty` at trade_physics.cuh:575 (HIGH T = forgiving exit ramp; LOW T = strict): over-holding pushes temp HIGH (exit ramp out of over-holding regime); at-or-below target pushes temp LOW (strict commitment). + +Why this beats the slot-373 mapping: +- Direct closed-loop control on the actual symptom (hold rate, the thing we're trying to fix) +- No chained-input-sentinel masking — hold rates are measured per-epoch from realised actions +- Survives fold reset cleanly — hold-rate measurement starts fresh in Fold 1 with real data + +Pearl-A bootstrap preserved (sentinel 50.0 → REPLACE on first valid observation). New cold-start sentinel guard fires when `observed_hold_rate` is at SP13 fold-reset sentinel 0.0 — the bootstrap-protection equivalent of the deleted dir_acc=0.5 guard. + +**Instrumentation:** every per-epoch launch now emits `HEALTH_DIAG[N]: min_hold_temp_diag obs_hold_rate=X target_hold_rate=Y overrun_norm=Z target_temp=T blended_temp=B` so the swap can be verified at every fold transition by reading the smoke logs. + +**Behavioral tests** (all GPU-gated, all pass on RTX 3050 Ti): +- `sp16_phase1_min_hold_temp_climbs_with_hold_overrun` +- `sp16_phase1_min_hold_temp_strict_when_at_target` +- `sp16_phase1_min_hold_temp_strict_when_under_target` +- `sp16_phase1_min_hold_temp_no_longer_reads_slot_373` (canonical proof of the swap) +- `sp16_phase1_min_hold_temp_mid_cadence_ema_after_bootstrap` +- `sp16_phase1_min_hold_temp_sentinel_observed_preserves_isv` + +**Original plan tasks (Tasks 1.1 / 1.2) below are kept for historical reference but are SUPERSEDED by the actual fix above.** + +### Task 1.1 (SUPERSEDED): Investigate producer launch lifecycle + +**Files:** +- Read: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — find launch sites for both producers +- Read: `crates/ml/src/trainers/dqn/state_reset_registry.rs` — fold-reset registry entries +- Read: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — fold-boundary call ordering + +- [ ] **Step 1: Locate the producer launch sites** + +```bash +grep -rn "min_hold_temperature_update\|kelly_warmup_floor_compute\|launch.*MIN_HOLD\|launch.*KELLY_WARMUP" \ + crates/ml/src/cuda_pipeline/ | head -20 +``` + +Find the launchers and identify what state guards them (epoch counter? warmup flag? fold-local flag?). + +- [ ] **Step 2: Trace fold-reset path** + +```bash +grep -nE "reset_for_fold|fold.*boundary|new_fold" \ + crates/ml/src/trainers/dqn/trainer/training_loop.rs \ + crates/ml/src/trainers/dqn/trainer/mod.rs \ + crates/ml-dqn/src/dqn.rs | head -20 +``` + +Identify the call sequence at fold boundary. Check whether the launchers are actually invoked in Fold 1 (may be guarded by an epoch>=warmup check that the per-fold reset doesn't address). + +- [ ] **Step 3: Identify root cause (one of three patterns)** + +**Pattern A — One-shot flag**: launcher has a `did_warmup` or `producer_initialized` flag that's set on first launch. Reset clears the buffer but not this flag. + +**Pattern B — Epoch-gate**: launcher only fires when `training_steps > some_threshold`. After fold reset, `training_steps=0` which keeps the launcher gated for several epochs. + +**Pattern C — Cross-fold state**: launcher reads from a slot that's reset to 0 (sentinel), but the producer's bootstrap relies on the slot already having a valid value. + +Document which pattern applies in the audit doc. + +### Task 1.2: Fix the producer launch + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — at the producer launch sites +- Modify (if needed): `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — fold-boundary path + +- [ ] **Step 1: Write the failing test** + +```rust +#[test] +#[ignore] +fn fold_reset_producer_relaunch_min_hold_temperature() { + // Setup: trainer at end of Fold 0 with min_hold_T=11.4 (adapted) + // Action: call reset_for_fold() + // Action: simulate epoch 1 of Fold 1 (training step + producer launch) + // Verify: isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] != 50.0 + // (50.0 is the sentinel/cold-start fallback; producer must override on first epoch) + todo!("write after Task 1.1 identifies root cause") +} + +#[test] +#[ignore] +fn fold_reset_producer_relaunch_kelly_warmup_floor() { + // Same pattern for KELLY_WARMUP_FLOOR_INDEX=330 + todo!() +} +``` + +Run: `cargo test -p ml --test sp14_oracle_tests --release fold_reset_producer_relaunch -- --ignored --nocapture` +Expected: FAIL + +- [ ] **Step 2: Apply the fix per Task 1.1's root-cause identification** + +If Pattern A (one-shot flag): reset the flag in `reset_for_fold()` per `feedback_no_partial_refactor` (atomic with the buffer clear). + +If Pattern B (epoch-gate): change the gate from `training_steps > N` to `(training_steps > N) || sentinel_detected`, where sentinel detection reads ISV[slot] and detects it's at the reset value. + +If Pattern C (bootstrap dependency): ensure the bootstrap fallback path is taken on first launch after reset. + +- [ ] **Step 3: Run tests, verify they pass** + +```bash +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --release fold_reset_producer_relaunch -- --ignored --nocapture +``` +Expected: 2/2 PASS + +- [ ] **Step 4: Run full sp14/sp15 oracle suite to catch regressions** + +```bash +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --test sp15_phase1_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -10 +``` +Expected: all pass. + +- [ ] **Step 5: Commit** + +```bash +git add -A +git commit -m "$(cat <<'EOF' +fix(sp16-p1): fold-reset producer re-launch for slots 460 + 330 + +Per train-multi-seed-pfh9n post-mortem: MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX +(slot 460) and KELLY_WARMUP_FLOOR_INDEX (slot 330) producers fail to fire +in Fold 1 — Fold 0 adapted 50→11.4 / 0→0.25, Fold 1 stuck at sentinel. + +Root cause: +Fix: + +Behavioral tests: +- fold_reset_producer_relaunch_min_hold_temperature +- fold_reset_producer_relaunch_kelly_warmup_floor + +Both verify slot value escapes sentinel within first epoch of Fold 1. + +Per feedback_no_partial_refactor — atomic with the original Class B #1 +fold-reset buffer clear (commit 658fec493). + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Phase 2 — Per-bar Hold cost: increase magnitude + ISV-adaptive + +**Why this matters:** Train-multi-seed-pfh9n epoch 2 emit: + +``` +HEALTH_DIAG[2]: hold_pricing observed_rate=0.2501 target=0.2000 cost=0.006251 +``` + +Realized Hold rate is 25%, target is 20%, but the *cost penalty* on Hold is only 0.006251 — tiny relative to per-bar reward magnitudes (popart=0.97, cf=0.65 in `reward_split`). The Hold action is essentially free, so Q-bias toward low-variance Hold has nothing to push against. + +### Task 2.1: Make Hold cost magnitude ISV-driven and 10-20× larger + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — add `HOLD_COST_SCALE_INDEX` slot +- Modify: `crates/ml/src/cuda_pipeline/state_layout.cuh` — C #define mirror +- Modify: existing hold-pricing kernel (find via grep) +- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` — register new slot +- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — dispatch arm + +- [ ] **Step 1: Find the existing hold-cost computation** + +```bash +grep -rn "hold_pricing\|hold.*cost\|hold_pen\|HOLD_COST" \ + crates/ml/src/cuda_pipeline/*.cu \ + crates/ml/src/cuda_pipeline/*.cuh | head -20 +``` + +Read the existing computation: `cost = base_cost × (observed_rate - target_rate)` or similar. Note the base coefficient. + +- [ ] **Step 2: Allocate ISV slot 461** + +Edit `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`: + +```rust +// SP16 Phase 2 (2026-05-08): adaptive Hold cost magnitude. Train-multi-seed-pfh9n +// post-mortem showed observed Hold rate climbed 25% → 52% across training while +// the cost penalty (~0.006) was 100× smaller than per-bar reward magnitudes +// (popart=0.97). Hold became effectively free, allowing Q(Hold) to dominate. +// Producer: hold_cost_scale_update (per-epoch, signal-driven from +// |observed_hold_rate - target_hold_rate|). +pub const HOLD_COST_SCALE_INDEX: usize = 461; +// ISV_TOTAL_DIM bumps 461 → 462 +``` + +Update `ISV_TOTAL_DIM` and add C #define mirror in `state_layout.cuh`. + +- [ ] **Step 3: Write the failing test** + +```rust +#[test] +#[ignore] +fn hold_cost_scale_climbs_with_overrun() { + // Setup: observed_hold_rate = 0.50, target = 0.20 (i.e., 30% overrun) + // Launch hold_cost_scale_update kernel + // Verify: isv[HOLD_COST_SCALE_INDEX] >= 5.0× base + // (intent: at 30% overrun, scale should be aggressive) + todo!() +} + +#[test] +#[ignore] +fn hold_cost_scale_decays_when_in_target() { + // Setup: observed_hold_rate = 0.20, target = 0.20 (perfect) + // Launch hold_cost_scale_update kernel a few times + // Verify: isv[HOLD_COST_SCALE_INDEX] decays toward 1.0 (base scale) + todo!() +} +``` + +Run: `cargo test -p ml --test sp14_oracle_tests --release hold_cost_scale -- --ignored --nocapture` +Expected: FAIL + +- [ ] **Step 4: Implement producer kernel** + +Create `crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu`: + +```c +/* hold_cost_scale_update — adaptive hold cost magnitude. + * + * Reads: observed_hold_rate, target_hold_rate (from existing diagnostic) + * Writes: ISV[HOLD_COST_SCALE_INDEX=461] + * + * Logic: scale = clamp(base + 20 × (observed - target)+, 1.0, 25.0) + * - base scale = 1.0 (no penalty when at target) + * - 30% overrun → scale = 1.0 + 20×0.3 = 7.0× (aggressive) + * - 50% overrun → scale = 11.0× (very aggressive, clipped at 25) + * + * Pearl-A bootstrap: first observation replaces sentinel. + * Welford slow EMA α=0.05 for steady state. + * + * Per-epoch boundary cadence. + */ +#include "state_layout.cuh" + +extern "C" __global__ void hold_cost_scale_update( + const float* __restrict__ isv, + float* __restrict__ isv_out, + float observed_hold_rate, + float target_hold_rate +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + float overrun = fmaxf(0.0f, observed_hold_rate - target_hold_rate); + float new_scale = 1.0f + 20.0f * overrun; + new_scale = fmaxf(1.0f, fminf(new_scale, 25.0f)); + + float prev_scale = isv[HOLD_COST_SCALE_INDEX]; + float blended; + if (prev_scale <= 0.0f) { + // Pearl-A bootstrap from sentinel + blended = new_scale; + } else { + // Welford slow EMA, α=0.05 (responds in ~20 epochs to steady drift) + const float alpha = 0.05f; + blended = (1.0f - alpha) * prev_scale + alpha * new_scale; + } + blended = fmaxf(1.0f, fminf(blended, 25.0f)); + + isv_out[HOLD_COST_SCALE_INDEX] = blended; +} +``` + +Add cubin entry in `crates/ml/build.rs`. Wire launch site in `gpu_dqn_trainer.rs` per-epoch boundary, beside other adaptive producers. + +- [ ] **Step 5: Migrate consumer** + +Find the existing hold-cost computation site (Task 2.1 Step 1). Multiply the base cost by `isv[HOLD_COST_SCALE_INDEX]`: + +```c +float cost_scale = isv[HOLD_COST_SCALE_INDEX]; +if (cost_scale <= 0.0f) cost_scale = 1.0f; // cold-start fallback +float cost = base_cost * cost_scale; // was: float cost = base_cost; +``` + +- [ ] **Step 6: Add dispatch arm** + +`crates/ml/src/trainers/dqn/state_reset_registry.rs` — add entry for slot 461 with sentinel value. +`crates/ml/src/trainers/dqn/trainer/training_loop.rs` — add reset_named_state dispatch arm. + +- [ ] **Step 7: Run tests** + +```bash +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --test sp15_phase1_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -15 +``` +Expected: all pass including 2 new hold_cost_scale tests. + +- [ ] **Step 8: Commit** + +```bash +git add -A +git commit -m "$(cat <<'EOF' +feat(sp16-p2): adaptive Hold cost scale via ISV[461] + +Per train-multi-seed-pfh9n post-mortem: observed_hold_rate climbed 0.25 → 0.52 +across training while cost penalty (~0.006) was 100× smaller than per-bar +reward magnitudes (popart=0.97). Hold action was effectively free. + +Fix: scale Hold cost adaptively. ISV[HOLD_COST_SCALE_INDEX=461] tracks: + scale = clamp(1.0 + 20 × max(0, observed - target), 1.0, 25.0) + +At 30% overrun (observed=0.5, target=0.2): scale → ~7.0×, aggressive. +At target: scale → 1.0× (no penalty). + +Pearl-A bootstrap + Welford slow EMA (α=0.05). + +ISV_TOTAL_DIM: 461 → 462. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Phase 3 — Validation L40S/H100 smoke + +### Task 3.1: Push and dispatch validation smoke + +- [ ] **Step 1: Push the SP16 chain** + +```bash +git push origin sp15-phase1-honest-numbers +``` + +- [ ] **Step 2: Dispatch H100 5-epoch smoke** + +```bash +SHORT=$(git rev-parse --short HEAD) +./scripts/argo-train.sh dqn --sha "$SHORT" --branch sp15-phase1-honest-numbers \ + --epochs 5 --gpu-pool ci-training-h100 --baseline --multi-seed 1 --folds 2 +``` + +- [ ] **Step 3: Monitor trajectory** + +Stream logs to `/tmp/train-sp16-${SHORT}.log` and watch for: + +**Success signals (pass to Phase 4):** +- `q_by_action [hold=X long=Y short=Z flat=W]` — Q(Hold) NOT climbing dominantly relative to Q(direction) +- Hold% stays below 30% across all 5 epochs +- Sharpe stable (not peak-then-decay) +- `min_hold_T` adapts in BOTH folds (not stuck at 50 in Fold 1) +- `sp9_kelly_warmup floor` non-zero in BOTH folds + +**Failure signals (kill-fast):** +- `q_by_action [hold]` climbs >2× faster than max(long, short, flat) +- Hold% reaches >40% in any epoch +- Sharpe decays >25% peak-to-final +- NaN/Inf, Q-explosion, dir_entropy<0.7 + +### Task 3.2: Diagnose and iterate + +- [ ] **If success signals**: dispatch 30-epoch L40S validation, compare to pfh9n trajectory. + +- [ ] **If failure signals (Q-bias persists despite Hold cost increase)**: structural Q-bias is deeper than reward shaping. Recommended next step is Phase 4 (out of scope for this plan): dueling Q-network architecture with advantage normalization. + +--- + +## Out of scope (defer to follow-up) + +These were considered and explicitly excluded from SP16: + +1. **Dueling Q-network architecture** — separate `V(s)` and `A(s,a)` heads with advantage normalization. Largest possible fix for structural Q-bias, but requires architecture refactor (~500-800 LOC across encoder/decoder, GRN trunk, IQN/C51 heads). Defer to SP17 if Phase 2 doesn't break the spiral. + +2. **Class B #3 PER beta annealing** — beta resets to start when `training_steps=0` at fold boundary, but old fold's transitions still in buffer have priorities computed under late-fold beta. Now mitigated by Class B #1 fold-buffer clear, so this issue is closed in practice. Re-evaluate if needed. + +3. **Class B #4 collapse-warmup mismatch** — gate uses `replay_buffer_capacity * 0.2 = 20K` instead of `min_replay_size = 1K`. Cosmetic at current scale; defer. + +4. **Class B #6 PER alpha adaptive** — `per_alpha=0.6` hardcoded. Class B #5 marked it MEDIUM severity. Adaptive α would be nice but not urgent. + +--- + +## Discipline rules (per project memory) + +- **`feedback_no_partial_refactor`** — fold-reset producer fixes migrate atomically (signature + all callers). +- **`feedback_isv_for_adaptive_bounds`** — Hold cost scale lives in ISV slot 461, not hardcoded. +- **`feedback_no_stubs`** — every test has a real body; no `todo!()` in committed code (Step 2 placeholders get filled in Step 4 of each task). +- **`feedback_no_atomicadd`** — block-tree-reduce only if any reduction needed. +- **`feedback_kill_runs_on_anomaly_quickly`** — Phase 3 smoke kills early on Q(Hold) climb signal. +- **`feedback_push_before_deploy`** — Phase 3 pushes before argo-train. + +--- + +## Self-review notes + +- **Phase 0 first** is non-negotiable: without `q_by_action` we have no falsifier for the Q-bias hypothesis. +- **Phase 1 fixes a separate bug** that surfaced during the smoke. It's small but worth doing atomically with Phase 2 (single L40S validation covers both). +- **Phase 2's 20× scale coefficient** is heuristic — Welford EMA + adaptive scale should self-tune at convergence. If validation shows oscillation, lower α from 0.05 to 0.01. +- **Total LOC estimate**: Phase 0 ~80, Phase 1 ~50-150 depending on root cause, Phase 2 ~250. Total ~400-500 LOC across 3 commits + 1 validation run.