fix(sp16-p1): MIN_HOLD_TEMPERATURE signal swap to hold-rate overrun + slot 330 non-bug closure
Per train-multi-seed-pfh9n post-mortem follow-up: slot 460 stuck at 50 in Fold 1 was NOT a launch-lifecycle bug. Producer fires per-epoch but kernel had early-return guard on AUX_DIR_ACC_SHORT_EMA (slot 373) at sentinel 0.5. Slot 373 reset on fold boundary; aux dir-acc EMA either didn't fire or settled within ε of 0.5 → kernel kept early-returning. Fix: drop slot 373 dependency entirely. Drive temperature from observed hold-rate vs target overrun: overrun = max(0, observed_hold_rate - target_hold_rate) overrun_norm = clamp(overrun / max(target, 0.01), 0, 1) new_temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × overrun_norm blended_temp = Welford EMA α=0.05 with Pearl-A bootstrap When over-holding: temp HIGH → exit ramp permissive (matches design intent). When at/under target: temp LOW → exit penalty strict. Survives fold reset: hold-rate measurement starts fresh with real data immediately, no chained-input-sentinel masking. Slot 330 (KELLY_WARMUP_FLOOR) investigated and confirmed NON-BUG: producer behaves correctly per pearl_kelly_cap_signal_driven_floors cross-fold- persistence. floor=0 post-warmup is correct steady state. Behavioral tests: - 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 Instrumentation: HEALTH_DIAG[N]: min_hold_temp_diag obs/target/overrun/temp. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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<CudaStream>,
|
||||
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)
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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]
|
||||
|
||||
@@ -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]:
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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<CudaStream>,
|
||||
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"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
// ═══════════════════════════════════════════════════════════════════════════
|
||||
|
||||
@@ -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<tgt) | 0 (clamped) | 5 | Max (no relaxation needed) |
|
||||
|
||||
### How — atomic single commit
|
||||
|
||||
- Kernel rewrite: `crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu` — drop slot 373 read; add `hold_rate_observed_idx` + `hold_rate_target_idx` parameters; new mapping + new sentinel guard.
|
||||
- Launcher signature change: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs::MinHoldTemperatureUpdateOps::launch` — replace `dir_acc_idx` + `sentinel_dir_acc` with `hold_rate_observed_idx` + `hold_rate_target_idx` + `sentinel_observed_hold_rate`.
|
||||
- Call site update: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs::launch_min_hold_temperature_update` — read SP13 hold-rate slots, pass to launcher.
|
||||
- ISV slot constants: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — replace `DIR_ACC_RANDOM_BASELINE` with `SENTINEL_HOLD_RATE_OBSERVED=0.0` (matches SP13 `sp13_hold_rate_observed_ema` fold-reset sentinel).
|
||||
- Registry description: `crates/ml/src/trainers/dqn/state_reset_registry.rs::sp14_audit_4b_min_hold_temperature_adaptive` — updated to reflect the new driving signal.
|
||||
- Instrumentation: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — emit a `min_hold_temp_diag` HEALTH_DIAG line after each per-epoch launch, exposing the realised hold-rate inputs + the computed target_temp + the post-EMA blended slot value. Lets us verify the swap works at every fold transition by reading the smoke logs.
|
||||
|
||||
```
|
||||
HEALTH_DIAG[N]: min_hold_temp_diag obs_hold_rate=0.2501 target_hold_rate=0.2000 overrun_norm=0.2505 target_temp=16.27 blended_temp=49.16
|
||||
```
|
||||
|
||||
### Why this is better than the slot-373 mapping
|
||||
|
||||
1. **Direct closed-loop control on the actual symptom.** Hold% is the visible problem; hold-rate-overrun is the corrective signal that directly addresses it.
|
||||
2. **No chained-input-sentinel masking.** Hold rates are measured per-epoch from realised 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.
|
||||
|
||||
### Why slot 330 stays untouched
|
||||
|
||||
Per `pearl_kelly_cap_signal_driven_floors`: Kelly warmup floor is intentionally cross-fold-persistent. The warmup floor's job is to relax once accumulated Kelly statistics mature — `floor = base × (1 − stat_conf)` where `stat_conf = min(1, count / target_count)`. Once `count ≥ target_count`, `floor → 0` is the correct end state, not a bug.
|
||||
|
||||
The original SP16 Phase 1 plan presumed the Fold-1 zero meant "producer didn't fire". Re-reading the producer kernel and slot 283's reset registry entry showed it was correct steady-state. **Documented here so a future investigator doesn't re-open the same dead end.**
|
||||
|
||||
### Files touched
|
||||
|
||||
| File | Δ | Notes |
|
||||
|------|---|-------|
|
||||
| `crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu` | +175 / -198 | Full rewrite — new signal source + new mapping + new sentinel guard |
|
||||
| `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` | +27 / -34 | Launcher signature change + doc-comment refresh |
|
||||
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | +12 / -5 | Call-site update — read SP13 hold-rate slots, drop slot 373 import |
|
||||
| `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` | +18 / -10 | Replace `DIR_ACC_RANDOM_BASELINE` with `SENTINEL_HOLD_RATE_OBSERVED`; layout-locked test updated |
|
||||
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +1 / -1 | Updated description for `sp14_audit_4b_min_hold_temperature_adaptive` |
|
||||
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +44 / -16 | Replace stale comment block + emit `min_hold_temp_diag` HEALTH_DIAG line |
|
||||
| `crates/ml/tests/sp14_oracle_tests.rs` | +280 / -213 | Replace 5 `sp14_audit_4b_min_hold_temperature_gpu` tests with 6 SP16 Phase 1 tests (including the slot-373-no-longer-read test) |
|
||||
|
||||
### Behavioral tests
|
||||
|
||||
`sp16_phase1_min_hold_temperature_gpu` (all #[ignore = "requires GPU"]):
|
||||
|
||||
1. **`sp16_phase1_min_hold_temp_climbs_with_hold_overrun`** — observed=0.40, target=0.20 → Pearl-A bootstrap from sentinel → temp = TEMP_MAX = 50.
|
||||
2. **`sp16_phase1_min_hold_temp_strict_when_at_target`** — observed=target=0.20, four-launch convergence → temp = TEMP_MIN = 5.
|
||||
3. **`sp16_phase1_min_hold_temp_strict_when_under_target`** — observed=0.10 < target=0.20, max(0, …) clamps overrun to 0 → temp = TEMP_MIN = 5.
|
||||
4. **`sp16_phase1_min_hold_temp_no_longer_reads_slot_373`** — slot 373 = 0.5 (OLD sentinel), kernel must still respond to hold-rate signals: (a) over-holding case verifies temp climbs to MAX even with slot 373 at OLD sentinel; (b) under-target case with slot 460 seeded at non-sentinel value verifies the EMA blend writes (vs the OLD kernel's no-write outcome which would have left slot 460 at the seed value).
|
||||
5. **`sp16_phase1_min_hold_temp_mid_cadence_ema_after_bootstrap`** — slot 460 seeded NOT at sentinel, observed=2×target → EMA blend 30 → 31.0.
|
||||
6. **`sp16_phase1_min_hold_temp_sentinel_observed_preserves_isv`** — observed=sentinel 0.0 → slot 460 preserved bit-exactly (cold-start fold-1 protection).
|
||||
|
||||
### Verification
|
||||
|
||||
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets` — clean (warnings unchanged).
|
||||
- `cargo test -p ml --test sp14_oracle_tests --release sp16_phase1 -- --ignored --nocapture` — 6/6 PASS on RTX 3050 Ti.
|
||||
- `cargo test -p ml --test sp14_oracle_tests --release -- --ignored` — 25/25 PASS.
|
||||
- `cargo test -p ml --test sp15_phase1_oracle_tests --release -- --ignored` — 36/36 PASS.
|
||||
|
||||
### Concerns
|
||||
|
||||
1. **Hold-rate observed slot is per-step + per-fold-EMA — survives fold reset, but lags the change.** When a new fold starts, the per-step `hold_rate_observer_kernel` re-populates slot 382 from the first batch's realised actions, then `apply_fixed_alpha_ema_kernel` blends with α=`HOLD_RATE_EMA_ALPHA`. So the temperature responds to actual realised over-holding rather than chained-input sentinels — but with one EMA-time-constant of lag. This lag is the design point of SP13's hold-pricing controller (which is the source of the slots we now read), so the kernel inherits the lag rather than introducing it.
|
||||
|
||||
2. **Per-epoch cadence vs hold-rate observation cadence.** The MIN_HOLD_TEMPERATURE producer fires per-epoch boundary; the hold-rate EMA updates per-step. Fold-1 epoch-1 sees the EMA after only one epoch's worth of per-step updates — typically enough to push the EMA off sentinel 0.0, but with little time-history. The instrumentation will reveal the exact magnitude of the inputs at every transition, so we can falsify "the kernel responds correctly" by reading the smoke log.
|
||||
|
||||
3. **Phase 2 dependency check.** The original SP16 plan's Phase 2 increases the magnitude of the per-bar hold cost. That's orthogonal to this signal swap — the temperature controls when the min-hold *exit* penalty bites; the cost controls how much sitting in Hold itself costs. Both are independently active in production.
|
||||
|
||||
@@ -0,0 +1,534 @@
|
||||
# SP16 — Hold-Collapse Structural Fix Implementation Plan
|
||||
|
||||
> **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) <noreply@anthropic.com>
|
||||
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: <Pattern A/B/C from Task 1.1>
|
||||
Fix: <description>
|
||||
|
||||
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) <noreply@anthropic.com>
|
||||
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) <noreply@anthropic.com>
|
||||
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.
|
||||
Reference in New Issue
Block a user