diff --git a/crates/ml/build.rs b/crates/ml/build.rs index f5eab187c..1ef7bbb82 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -867,18 +867,8 @@ fn main() { // see docs/dqn-wire-up-audit.md § "SP16 Phase 1 (revised)" for // rationale. "min_hold_temperature_update_kernel.cu", - // hold_cost_scale_update_kernel.cu — SP16 Phase 2 adaptive Hold - // cost scale producer. Reads ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] - // and ISV[HOLD_RATE_TARGET_INDEX=381] (same input chain as SP16-P1 - // MIN_HOLD_TEMPERATURE), maps overrun - // `clamp(max(0, observed - target) / max(target, 0.01), 0, 1)` - // to a [SCALE_MIN=1, SCALE_MAX=25] multiplier, blends via Pearl-A - // bootstrap + Welford EMA α=0.05, writes to - // ISV[HOLD_COST_SCALE_INDEX=461]. Consumer: per-bar Hold cost - // subtraction sites in experience_kernels.cu multiply - // `isv[ISV_HOLD_COST_IDX=380]` by `isv[461]`. See - // docs/dqn-wire-up-audit.md § "SP16 Phase 2" for rationale. - "hold_cost_scale_update_kernel.cu", + // hold_cost_scale_update_kernel.cu DELETED by SP18 Phase 1 + // (2026-05-09) — see docs/sp18-wireup-audit.md. // SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward // decomposition diagnostic. Reads `reward_components_per_sample` // ([n_samples × 6]) + `actions` ([n_samples], factored-action diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index ddc524bc9..9c273b4da 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -3150,21 +3150,9 @@ extern "C" __global__ void experience_env_step( * this branch covers that edge. Per-bar (non-segment_complete) * Hold-cost subtraction is in the per-bar branches below. * - * SP16 Phase 2 (2026-05-08): adaptive Hold cost scale via - * ISV[HOLD_COST_SCALE_INDEX=461]. The static base cost (~0.006) - * was ~100× smaller than per-bar reward magnitudes (popart=0.97, - * cf=0.65) post-pfh9n — the producer in - * `hold_cost_scale_update_kernel` reads hold-rate overrun and - * outputs a [1, 25] multiplier into slot 461. Cold-start - * fallback: when slot 461 is at sentinel 0.0 OR outside [1, 25], - * use scale=1.0 (bit-identical pre-Phase-2 cost magnitude). */ - if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL) { - float hold_cost_scale = isv_signals_ptr[ISV_HOLD_COST_SCALE_IDX]; - if (hold_cost_scale < 1.0f || hold_cost_scale > 25.0f) { - hold_cost_scale = 1.0f; /* cold-start / OOB fallback */ - } - base_reward -= isv_signals_ptr[ISV_HOLD_COST_IDX] * hold_cost_scale; - } + * SP18 D-leg replacement: per-bar Hold cost subtraction removed. + * Replaced with compute_sp18_hold_opportunity_cost in Phase 2. + * (No reward write here in interim; Phase 2 lands the new device fn.) */ /* SP12 v3 fix (2026-05-04): asymmetric bounded cap (loss aversion). * SP11 (commit 35db31089) made the cap symmetric ±10 to fix slot-63 * PopArt EMA inflation. That fix preserved stability but erased the @@ -3627,20 +3615,9 @@ extern "C" __global__ void experience_env_step( * required. Spec §"Change 1: Price Hold (replaces v2's Eliminate * Hold)". * - * SP16 Phase 2 (2026-05-08): adaptive Hold cost scale via - * ISV[HOLD_COST_SCALE_INDEX=461]. Same fallback semantics as the - * segment_complete branch above — when slot 461 is at sentinel - * 0.0 OR out-of-bounds, use scale=1.0 (bit-identical pre-Phase-2 - * cost). At 100% overrun the scale climbs to 25× → effective - * cost ~0.15/bar, competitive with per-bar reward magnitudes - * but still below the SP12 cap [-10, +5]. */ - if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL) { - float hold_cost_scale = isv_signals_ptr[ISV_HOLD_COST_SCALE_IDX]; - if (hold_cost_scale < 1.0f || hold_cost_scale > 25.0f) { - hold_cost_scale = 1.0f; /* cold-start / OOB fallback */ - } - r_micro -= isv_signals_ptr[ISV_HOLD_COST_IDX] * hold_cost_scale; - } + * SP18 D-leg replacement: per-bar Hold cost subtraction removed. + * Replaced with compute_sp18_hold_opportunity_cost in Phase 2. + * (No reward write here in interim; Phase 2 lands the new device fn.) */ reward = r_micro; micro_reward_per_sample[out_off] = r_micro; reward_components_per_sample[out_off * 6 + 3] = r_micro; @@ -3703,17 +3680,9 @@ extern "C" __global__ void experience_env_step( * structurally bounded by HOLD_COST_BASE × CEIL_RATIO = 0.005 * per bar, well below the SP12 asymmetric cap [-10, +5]. * - * SP16 Phase 2 (2026-05-08): adaptive Hold cost scale via - * ISV[HOLD_COST_SCALE_INDEX=461]. Same fallback semantics as - * the other 2 sites — when slot 461 is at sentinel 0.0 OR - * out-of-bounds, use scale=1.0 (bit-identical pre-Phase-2). */ - if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL) { - float hold_cost_scale = isv_signals_ptr[ISV_HOLD_COST_SCALE_IDX]; - if (hold_cost_scale < 1.0f || hold_cost_scale > 25.0f) { - hold_cost_scale = 1.0f; /* cold-start / OOB fallback */ - } - r_opp_cost -= isv_signals_ptr[ISV_HOLD_COST_IDX] * hold_cost_scale; - } + * SP18 D-leg replacement: per-bar Hold cost subtraction removed. + * Replaced with compute_sp18_hold_opportunity_cost in Phase 2. + * (No reward write here in interim; Phase 2 lands the new device fn.) */ reward = r_opp_cost; reward_components_per_sample[out_off * 6 + 4] = r_opp_cost; } diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index 4a658d724..be1a4a811 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -53,7 +53,7 @@ use crate::MLError; use super::gpu_dqn_trainer::{ AUX_HORIZON_UPDATE_CUBIN, AUX_TRUNK_BACKWARD_CUBIN, AUX_TRUNK_FORWARD_CUBIN, AVG_WIN_HOLD_TIME_UPDATE_CUBIN, DD_SATURATION_FLOOR_UPDATE_CUBIN, - HOLD_COST_SCALE_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, + H_S2_AUX_RMS_EMA_CUBIN, KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, MIN_HOLD_TEMPERATURE_UPDATE_CUBIN, REWARD_CAP_UPDATE_CUBIN, }; @@ -1201,123 +1201,10 @@ impl HS2AuxRmsEmaOps { } } -/// SP16 Phase 2 (2026-05-08): adaptive Hold cost scale producer ops. -/// -/// Mirrors `MinHoldTemperatureUpdateOps` (same input chain — hold-rate -/// overrun — same per-epoch cadence — same Pearl-A bootstrap pattern). -/// The kernel is a single-block single-thread cold-path producer that -/// reads `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and -/// `ISV[HOLD_RATE_TARGET_INDEX=381]`, maps overrun to a [1, 25] -/// multiplier, and Pearl-A-bootstrap + α=0.05 EMA blends into -/// `ISV[HOLD_COST_SCALE_INDEX=461]`. -/// -/// # Why this struct exists separately from MinHoldTemperatureUpdateOps -/// -/// Both producers read the SAME inputs (slots 382 + 381) and run at the -/// SAME cadence (per-epoch boundary), but write to DIFFERENT outputs -/// (slot 460 — temperature in [5, 50] vs slot 461 — scale in [1, 25]). -/// Combining them into a single fused kernel was rejected because: -/// 1. The two outputs are downstream of two semantically-distinct -/// consumer paths (min-hold soft saturation vs Hold cost -/// magnitude); fusing them couples kernels that should be -/// independently auditable. -/// 2. Per-epoch launch overhead is negligible (single-thread cold -/// kernel), so the fusion buys no measurable perf. -/// 3. Mirrors the SP14-P0-A and SP14-P1-Producer split — ONE producer -/// kernel per ISV slot, ONE wiring point per consumer. -/// -/// # Pearls applied -/// - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. -/// - `pearl_first_observation_bootstrap.md` — sentinel 0.0 → REPLACE. -/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` -/// pre-loaded at construction; on-device guards. -/// - `pearl_symmetric_clamp_audit.md` — bilateral -/// `fmaxf(lo, fminf(x, hi))` clamp on target_scale + post-blend. -/// - `feedback_isv_for_adaptive_bounds.md` — bounds [1, 25] are -/// Category-1 dimensional safety floors (NOT tuning). -/// - `feedback_no_partial_refactor.md` — kernel + launcher + 3 consumer -/// sites + tests + audit doc updated atomically in the same commit. -#[allow(missing_debug_implementations)] -pub(crate) struct HoldCostScaleUpdateOps { - update_kernel: CudaFunction, -} - -impl HoldCostScaleUpdateOps { - pub(crate) fn new(stream: &Arc) -> Result { - let context = stream.context(); - let module = context - .load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec()) - .map_err(|e| MLError::ModelError(format!("hold_cost_scale_update cubin load: {e}")))?; - let update_kernel = module - .load_function("hold_cost_scale_update") - .map_err(|e| MLError::ModelError(format!("hold_cost_scale_update load: {e}")))?; - Ok(Self { update_kernel }) - } - - /// Launch the adaptive Hold cost scale producer (SP16 Phase 2 + - /// T3 Wiener-optimal α). - /// - /// Args: - /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). - /// - `scale_idx`: HOLD_COST_SCALE_INDEX (461). - /// - `hold_rate_observed_idx`: HOLD_RATE_OBSERVED_EMA_INDEX (382). - /// - `hold_rate_target_idx`: HOLD_RATE_TARGET_INDEX (381). - /// - Welford accumulator slot indices (T3): HCS_TARGET_MEAN_INDEX - /// (462), HCS_TARGET_M2_INDEX (463), HCS_DIFF_MEAN_INDEX (464), - /// HCS_DIFF_M2_INDEX (465), HCS_PREV_TARGET_INDEX (466), - /// HCS_SAMPLE_COUNT_INDEX (467). - /// - `sentinel_scale`: SENTINEL_HOLD_COST_SCALE (0.0). - /// - `sentinel_observed_hold_rate`: SENTINEL_HOLD_RATE_OBSERVED - /// (0.0; matches the SP13 fold-reset registry sentinel — shared - /// by design with SP16-P1). - /// - `scale_min`/`scale_max`: HOLD_COST_SCALE_MIN/MAX (1.0, 25.0). - /// - /// SP16 T3: the `alpha` argument is gone — α is computed inside the - /// kernel from the Welford accumulators (Wiener-optimal: α = diff_var - /// / (diff_var + sample_var + ε)). - #[allow(clippy::too_many_arguments)] - pub(crate) fn launch( - &self, - stream: &Arc, - isv_ptr: u64, - scale_idx: i32, - hold_rate_observed_idx: i32, - hold_rate_target_idx: i32, - target_mean_idx: i32, - target_m2_idx: i32, - diff_mean_idx: i32, - diff_m2_idx: i32, - prev_target_idx: i32, - sample_count_idx: i32, - sentinel_scale: f32, - sentinel_observed_hold_rate: f32, - scale_min: f32, - scale_max: f32, - ) -> Result<(), MLError> { - unsafe { - stream - .launch_builder(&self.update_kernel) - .arg(&isv_ptr) - .arg(&scale_idx) - .arg(&hold_rate_observed_idx) - .arg(&hold_rate_target_idx) - .arg(&target_mean_idx) - .arg(&target_m2_idx) - .arg(&diff_mean_idx) - .arg(&diff_m2_idx) - .arg(&prev_target_idx) - .arg(&sample_count_idx) - .arg(&sentinel_scale) - .arg(&sentinel_observed_hold_rate) - .arg(&scale_min) - .arg(&scale_max) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), - shared_mem_bytes: 0, - }) - .map_err(|e| MLError::ModelError(format!("hold_cost_scale_update: {e}")))?; - } - Ok(()) - } -} +// HoldCostScaleUpdateOps DELETED by SP18 Phase 1 (2026-05-09) — the +// SP16 Phase 2 / T3 reactive Hold-cost-scale chain (kernel + ops + 3 +// consumer sites in experience_kernels.cu) is replaced by the SP18 +// D-leg structural Hold opportunity-cost reward (Phase 2 lands the +// `compute_sp18_hold_opportunity_cost` device fn at the SP18-bound caps +// in slots [483..493)). See docs/sp18-wireup-audit.md (A1) for the +// atomic-deletion record. diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index e9bf7a3f9..8e1c046bf 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2229,32 +2229,8 @@ pub(crate) static DD_SATURATION_FLOOR_UPDATE_CUBIN: &[u8] = include_bytes!(conca /// launch (same cadence as P0-A REWARD_POS_CAP). pub(crate) static MIN_HOLD_TEMPERATURE_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/min_hold_temperature_update_kernel.cubin")); -/// SP16 Phase 2 (2026-05-08): adaptive Hold cost scale producer cubin. -/// Single-thread cold-path kernel mirroring the SP16-P1 -/// `min_hold_temperature_update_kernel` driver pattern. Reads -/// `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and -/// `ISV[HOLD_RATE_TARGET_INDEX=381]` (same input chain as SP16-P1) and -/// maps the overrun into a [1, 25] multiplier: -/// -/// overrun = max(0, observed − target) -/// overrun_norm = clamp(overrun / max(target, 0.01), 0, 1) -/// target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_norm -/// -/// then Pearl-A-bootstrap + Welford-α=0.05 EMA blends into -/// `ISV[HOLD_COST_SCALE_INDEX=461]`. The 3 consumer sites in -/// `experience_kernels.cu` multiply `isv[ISV_HOLD_COST_IDX=380]` by -/// `isv[461]` to scale the per-bar Hold cost when over-holding is -/// observed. Cold-start fallback: when slot 461 is at sentinel 0.0 OR -/// outside [1, 25], the consumer uses scale=1.0 (bit-identical -/// pre-Phase-2 behaviour). Per-mortem of train-multi-seed-pfh9n showed -/// realised Hold-rate climbing 0.25 → 0.52 with cost magnitude ~0.006 -/// (~100× smaller than per-bar reward magnitudes); the adaptive scale -/// pushes effective cost to ~0.15/bar at 100% overrun, competitive with -/// per-bar reward magnitudes ~0.01-0.1 but well below SP12 cap [-10, +5]. -/// Pearl-A first-observation bootstrap (sentinel 0.0 → REPLACE). α=0.05 -/// mid-cadence EMA. Loaded by `gpu_aux_trunk::HoldCostScaleUpdateOps`. -/// Per-epoch boundary launch (same cadence as SP16-P1 producer). -pub(crate) static HOLD_COST_SCALE_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/hold_cost_scale_update_kernel.cubin")); +// HOLD_COST_SCALE_UPDATE_CUBIN DELETED by SP18 Phase 1 (2026-05-09). +// See docs/sp18-wireup-audit.md. /// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer cubin. /// Single-block 256-thread kernel computing `sqrt(mean(h_s2_aux²))` over @@ -3661,9 +3637,7 @@ const fn layout_fingerprint_seed() -> &'static [u8] { KELLY_PRIOR_WINS=454;KELLY_PRIOR_LOSSES=455;KELLY_PRIOR_SUM_WINS=456;KELLY_PRIOR_SUM_LOSSES=457;\ DD_SATURATION_FLOOR_ADAPTIVE=458;\ PLAN_THRESHOLD_FLOOR_ADAPTIVE=459;MIN_HOLD_TEMPERATURE_ADAPTIVE=460;\ - HOLD_COST_SCALE=461;\ - HCS_TARGET_MEAN=462;HCS_TARGET_M2=463;HCS_DIFF_MEAN=464;\ - HCS_DIFF_M2=465;HCS_PREV_TARGET=466;HCS_SAMPLE_COUNT=467;\ + RESERVED_GAP_461_TO_468=SP18_P1_RETIRED;\ MHT_TARGET_MEAN=468;MHT_TARGET_M2=469;MHT_DIFF_MEAN=470;\ MHT_DIFF_M2=471;MHT_PREV_TARGET=472;MHT_SAMPLE_COUNT=473;\ SLOT_474_A_VAR_EMA_DIR=474;\ @@ -7845,12 +7819,7 @@ pub struct GpuDqnTrainer { /// boundary launch (same cadence as P0-A REWARD_POS_CAP). min_hold_temperature_update_ops: super::gpu_aux_trunk::MinHoldTemperatureUpdateOps, - /// SP16 Phase 2 (2026-05-08): per-epoch adaptive Hold cost scale - /// producer. Pre-loads the `CudaFunction` handle at construction - /// per `pearl_no_host_branches_in_captured_graph.md`. Per-epoch - /// boundary launch (same cadence as SP16-P1 - /// MIN_HOLD_TEMPERATURE — same input chain). - hold_cost_scale_update_ops: super::gpu_aux_trunk::HoldCostScaleUpdateOps, + // hold_cost_scale_update_ops field DELETED by SP18 Phase 1 (2026-05-09). // ── Speculative inference cache ── /// Cached trunk output [B, SH2] from between-bar speculative forward. @@ -15793,89 +15762,9 @@ impl GpuDqnTrainer { Ok(()) } - /// SP16 Phase 2 (2026-05-08): launch GPU-only adaptive Hold cost - /// scale producer. - /// - /// Reads `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and - /// `ISV[HOLD_RATE_TARGET_INDEX=381]` (same input chain as SP16-P1 - /// MIN_HOLD_TEMPERATURE), maps overrun into a [1, 25] multiplier, - /// and Pearl-A-bootstrap + α=0.05 EMA blends into - /// `ISV[HOLD_COST_SCALE_INDEX=461]`. Consumed by 3 sites in - /// `experience_kernels.cu` (segment_complete + per-bar positioned- - /// Hold + per-bar flat-Hold) which multiply - /// `isv[ISV_HOLD_COST_IDX=380]` by `isv[461]` to scale the - /// per-bar Hold cost when over-holding is observed. - /// - /// Why this kernel exists (post-mortem of train-multi-seed-pfh9n): - /// realised Hold rate climbed 0.25 → 0.52 across training while the - /// cost penalty stayed at ~0.006 — ~100× smaller than per-bar - /// reward magnitudes (popart=0.97, cf=0.65). Hold became - /// effectively free. The adaptive scale pushes effective cost to - /// ~0.15/bar at 100% overrun (observed=2× target), competitive - /// with per-bar reward magnitudes (~0.01-0.1) but well below the - /// SP12 asymmetric cap [-10, +5] so no clamp interaction. - /// - /// Per `feedback_isv_for_adaptive_bounds.md` the bounds [1, 25] are - /// Category-1 dimensional safety floors (NOT tuning). - /// - /// Per `feedback_no_partial_refactor.md` the kernel + launcher + 3 - /// consumer sites + tests + audit doc land in the same atomic commit. - /// - /// Per `pearl_canary_input_freshness_launch_order.md` this launch - /// MUST fire AFTER the per-step `hold_rate_observer` chain (which - /// updates the input slot 382 every step) — same launch ordering - /// constraint as the sibling `launch_min_hold_temperature_update`. - /// - /// # Pearls applied - /// - `pearl_first_observation_bootstrap.md` (sentinel 0.0 → REPLACE) - /// - `pearl_no_host_branches_in_captured_graph.md` - /// - `pearl_symmetric_clamp_audit.md` (bilateral clamp on - /// target_scale + post-blend) - /// - `feedback_isv_for_adaptive_bounds.md` (bounds [1, 25] are - /// dimensional safety floors, not tuning) - /// - `feedback_no_atomicadd.md` (single-thread kernel) - pub(crate) fn launch_hold_cost_scale_update( - &self, - ) -> Result<(), MLError> { - use crate::cuda_pipeline::sp14_isv_slots::{ - HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX, - HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED, - HCS_TARGET_MEAN_INDEX, HCS_TARGET_M2_INDEX, - HCS_DIFF_MEAN_INDEX, HCS_DIFF_M2_INDEX, - HCS_PREV_TARGET_INDEX, HCS_SAMPLE_COUNT_INDEX, - }; - use crate::cuda_pipeline::sp13_isv_slots::{ - HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, - }; - - debug_assert!( - self.isv_signals_dev_ptr != 0, - "launch_hold_cost_scale_update: isv_signals_dev_ptr must be allocated" - ); - - // SP16 T3 (2026-05-08): hardcoded `α=0.05` removed — α is now - // Wiener-optimal, derived from Welford accumulators in slots - // [462..468). Per `pearl_wiener_optimal_adaptive_alpha`. - self.hold_cost_scale_update_ops.launch( - &self.stream, - self.isv_signals_dev_ptr, - HOLD_COST_SCALE_INDEX as i32, - HOLD_RATE_OBSERVED_EMA_INDEX as i32, - HOLD_RATE_TARGET_INDEX as i32, - HCS_TARGET_MEAN_INDEX as i32, - HCS_TARGET_M2_INDEX as i32, - HCS_DIFF_MEAN_INDEX as i32, - HCS_DIFF_M2_INDEX as i32, - HCS_PREV_TARGET_INDEX as i32, - HCS_SAMPLE_COUNT_INDEX as i32, - SENTINEL_HOLD_COST_SCALE, - SENTINEL_HOLD_RATE_OBSERVED, - HOLD_COST_SCALE_MIN, - HOLD_COST_SCALE_MAX, - )?; - - Ok(()) - } + // launch_hold_cost_scale_update DELETED by SP18 Phase 1 (2026-05-09). + // The SP16 P2/T3 reactive Hold-cost-scale chain is replaced by the + // SP18 D-leg structural Hold opportunity-cost reward (Phase 2). /// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary /// producer. @@ -24424,14 +24313,8 @@ impl GpuDqnTrainer { let min_hold_temperature_update_ops = super::gpu_aux_trunk::MinHoldTemperatureUpdateOps::new(&stream)?; - // SP16 Phase 2 (2026-05-08): adaptive Hold cost scale producer. - // Pre-loads the `CudaFunction` handle at construction per - // `pearl_no_host_branches_in_captured_graph.md`. Per-epoch - // boundary launch (same cadence and input chain as SP16-P1 - // MIN_HOLD_TEMPERATURE — they share slots 382 + 381 as input - // and write to slots 460 + 461 respectively). - let hold_cost_scale_update_ops = - super::gpu_aux_trunk::HoldCostScaleUpdateOps::new(&stream)?; + // hold_cost_scale_update_ops constructor DELETED by SP18 Phase 1 + // (2026-05-09). // ── Speculative inference cache ── let speculative_h_s2 = stream.alloc_zeros::(b * sh2) @@ -25292,8 +25175,7 @@ impl GpuDqnTrainer { dd_saturation_floor_update_ops, // Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive MIN_HOLD_TEMPERATURE producer. min_hold_temperature_update_ops, - // SP16 Phase 2 (2026-05-08): adaptive Hold cost scale producer. - hold_cost_scale_update_ops, + // hold_cost_scale_update_ops field-init DELETED by SP18 Phase 1. speculative_h_s2, speculative_features, speculative_valid: false, diff --git a/crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu b/crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu deleted file mode 100644 index a433a6918..000000000 --- a/crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu +++ /dev/null @@ -1,236 +0,0 @@ -/* ══════════════════════════════════════════════════════════════════════════ - * SP16 Phase 2 (2026-05-08) + T3 (Wiener-optimal α, 2026-05-08) — adaptive - * Hold cost scale producer. - * - * Mirrors the SP16 Phase 1 (revised) MIN_HOLD_TEMPERATURE producer - * driver pattern (`min_hold_temperature_update_kernel.cu`) — same - * input signals (HOLD_RATE_OBSERVED_EMA + HOLD_RATE_TARGET), same - * overrun-norm shape, same cadence (per-epoch boundary), same Pearl-A - * bootstrap, same Wiener-optimal α (T3) — but maps the overrun into - * the Hold-cost-scale dimension instead of the min-hold-temperature - * dimension. - * - * Mapping (unchanged from T2): - * - * overrun = max(0, observed − target) - * overrun_normalized = clamp(overrun / max(target, 0.01), 0, 1) - * target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_normalized - * - * SP16 T3 — Wiener-optimal adaptive α (2026-05-08): - * - * The pre-T3 implementation hardcoded `const float alpha = 0.05f` for - * the EMA blend. train-multi-seed-hjzss validation showed this α was - * far too low to converge in 5-epoch smoke. Per - * `pearl_wiener_optimal_adaptive_alpha`: - * - * α = diff_var / (diff_var + sample_var + ε) - * - * Where `sample_var` = running variance of target_scale and - * `diff_var` = running variance of consecutive one-step differences. - * Welford accumulators stored in ISV slots [462..468): - * - * - HCS_TARGET_MEAN_INDEX (462) — running mean of target_scale - * - HCS_TARGET_M2_INDEX (463) — Welford ΣΔ² (target_scale) - * - HCS_DIFF_MEAN_INDEX (464) — running mean of consecutive diffs - * - HCS_DIFF_M2_INDEX (465) — Welford ΣΔ² (diff) - * - HCS_PREV_TARGET_INDEX (466) — previous-epoch target_scale - * - HCS_SAMPLE_COUNT_INDEX (467) — Welford counter - * - * Cold-start path: when `sample_count < 3`, α = 1.0 → REPLACE. After - * the third observation, sample_var and diff_var are well-defined and - * α emerges from the data. - * - * Bounds [WELFORD_ALPHA_MIN=0.4, WELFORD_ALPHA_MAX=0.95] applied to - * the Wiener-derived α before the blend. The MIN floor is a - * structural-control constraint (NOT a tuned value): Wiener-optimal - * α is MSE-optimal only for *stationary* signals, and this loop - * tracks a hold-rate overrun driven by a co-adapting policy. Once - * Welford accumulates 3+ samples and the policy converges toward - * target, diff_var → 0 and α_wiener → 0, leaving the controller - * unable to catch up when the policy subsequently drifts back above - * target (empirical anchor: train-multi-seed-b5gmp sha 641aa0dfd - * Fold 1 collapse — α decayed to 0.248 by HD4, blended scale lagged - * climbing target by HD5, Hold% 30.6% → 52.6% → Sharpe 55.55 vs - * 88.65 baseline). The 0.4 floor preserves ~2.5-step catch-up - * bandwidth. The MAX 0.95 prevents pathological α=1.0 (no - * smoothing). Per `pearl_wiener_alpha_floor_for_nonstationary`. - * - * Algorithm (single-block single-thread — cold path, per-epoch boundary): - * Phases 1-7 mirror the MHT producer; only the slot indices and - * bound constants differ. See `min_hold_temperature_update_kernel.cu` - * for the canonical algorithm comment. - * - * Pearls + invariants: - * - * - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. - * - * - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` - * pre-loaded at construction; on-device guards. - * - * - `pearl_first_observation_bootstrap.md` — first valid hold-rate - * observation replaces sentinel 0.0 directly; no blend. - * - * - `pearl_wiener_optimal_adaptive_alpha.md` (NEW T3) — α derived - * from Welford variance accumulators, not hardcoded. - * - * - `feedback_no_stubs.md` — full body, no return-zero placeholders. - * - * - `feedback_isv_for_adaptive_bounds.md` — bounds [1, 25] are - * dimensional safety floors; α is no longer hardcoded. - * - * - `pearl_symmetric_clamp_audit.md` — bilateral - * `fmaxf(lo, fminf(x, hi))` clamp on target_scale before writing - * and post-blend. - * - * - `feedback_no_partial_refactor.md` — kernel + launcher + 3 - * consumer sites + tests + audit doc updated atomically; both - * T1 and T2 producers migrate together to Wiener-optimal α. - * - * Args: - * isv — `[ISV_TOTAL_DIM]` device f32, modified - * in place at index `scale_idx` and 6 - * Welford slots [462..468). - * scale_idx — HOLD_COST_SCALE_INDEX (461). - * hold_rate_observed_idx — HOLD_RATE_OBSERVED_EMA_INDEX (382). - * hold_rate_target_idx — HOLD_RATE_TARGET_INDEX (381). - * target_mean_idx — HCS_TARGET_MEAN_INDEX (462). - * target_M2_idx — HCS_TARGET_M2_INDEX (463). - * diff_mean_idx — HCS_DIFF_MEAN_INDEX (464). - * diff_M2_idx — HCS_DIFF_M2_INDEX (465). - * prev_target_idx — HCS_PREV_TARGET_INDEX (466). - * sample_count_idx — HCS_SAMPLE_COUNT_INDEX (467). - * sentinel_scale — SENTINEL_HOLD_COST_SCALE (0.0). - * sentinel_observed_hold_rate — SENTINEL_HOLD_RATE_OBSERVED (0.0). - * scale_min — HOLD_COST_SCALE_MIN (1.0). - * scale_max — HOLD_COST_SCALE_MAX (25.0). - * - * Launch: grid=(1, 1, 1), block=(1, 1, 1). - * No shared memory — single-thread cold-path kernel. - * ══════════════════════════════════════════════════════════════════════════ */ - -#include - -#define EPS_F 1e-6f -// WELFORD_ALPHA_MIN_F = 0.4f — non-stationary control-loop floor. -// Wiener-α is MSE-optimal for stationary signals; this controller -// tracks a hold-rate overrun driven by a co-adapting policy, so α -// must never decay below the bandwidth needed to catch up when the -// policy drifts. Cold-start (sample_count < WELFORD_MIN_SAMPLE_COUNT_F) -// uses WELFORD_COLD_START_ALPHA_F=1.0 REPLACE per Pearl A, bypassing -// this floor. See `pearl_wiener_alpha_floor_for_nonstationary`. -#define WELFORD_ALPHA_MIN_F 0.4f -#define WELFORD_ALPHA_MAX_F 0.95f -#define WELFORD_MIN_SAMPLE_COUNT_F 3.0f -#define WELFORD_COLD_START_ALPHA_F 1.0f - -extern "C" __global__ -void hold_cost_scale_update( - float* __restrict__ isv, - int scale_idx, - int hold_rate_observed_idx, - int hold_rate_target_idx, - int target_mean_idx, - int target_M2_idx, - int diff_mean_idx, - int diff_M2_idx, - int prev_target_idx, - int sample_count_idx, - float sentinel_scale, - float sentinel_observed_hold_rate, - float scale_min, - float scale_max) -{ - /* Single-thread kernel — only thread 0 of block 0 does anything. */ - if (threadIdx.x != 0 || blockIdx.x != 0) return; - - const float observed_hold_rate = isv[hold_rate_observed_idx]; - const float target_hold_rate = isv[hold_rate_target_idx]; - - /* Cold-start sentinel guard: observed hold-rate EMA is at sentinel - * 0.0 — keep the scale slot and every Welford accumulator - * unchanged. The consumer falls back to scale=1.0 when slot is at - * sentinel 0.0 (bit-identical pre-Phase-2 cost magnitude). */ - if (fabsf(observed_hold_rate - sentinel_observed_hold_rate) < EPS_F) { - return; - } - - /* Hold-rate overrun mapping. */ - 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_scale = scale_min + (scale_max - scale_min) * overrun_norm; - target_scale = fmaxf(scale_min, fminf(target_scale, scale_max)); - - /* Read Welford accumulator state. */ - float sample_count = isv[sample_count_idx]; - float target_mean = isv[target_mean_idx]; - float target_M2 = isv[target_M2_idx]; - float diff_mean = isv[diff_mean_idx]; - float diff_M2 = isv[diff_M2_idx]; - const float prev_target = isv[prev_target_idx]; - - /* Update Welford on target_scale. */ - sample_count += 1.0f; - const float t_delta = target_scale - target_mean; - target_mean += t_delta / sample_count; - const float t_delta2 = target_scale - target_mean; - target_M2 += t_delta * t_delta2; - - /* Update Welford on consecutive diff (need ≥ 2 samples). */ - const float diff_count = sample_count - 1.0f; - if (sample_count >= 2.0f) { - const float diff = target_scale - prev_target; - const float d_delta = diff - diff_mean; - diff_mean += d_delta / diff_count; - const float d_delta2 = diff - diff_mean; - diff_M2 += d_delta * d_delta2; - } - - /* Compute Wiener-optimal α. */ - float alpha; - if (sample_count < WELFORD_MIN_SAMPLE_COUNT_F) { - /* Cold-start REPLACE — keeps producer responsive in epochs - * 1-2 while Welford state stabilises by epoch 3. */ - alpha = WELFORD_COLD_START_ALPHA_F; - } else { - const float sample_var = target_M2 / (sample_count - 1.0f); - const float diff_var = (diff_count >= 2.0f) - ? diff_M2 / (diff_count - 1.0f) - : diff_M2; - alpha = diff_var / (diff_var + sample_var + EPS_F); - /* Floor at WELFORD_ALPHA_MIN_F=0.4: Wiener-α is MSE-optimal - * for stationary signals only; this control loop tracks an - * overrun driven by the co-adapting policy. Without the floor - * α decays toward 0 once observed_hold_rate converges toward - * target, leaving the controller unable to react when the - * policy drifts (canonical regression: train-multi-seed-b5gmp - * sha 641aa0dfd Fold 1, α=0.248 at HD4 → Hold% 52.6% at HD5). - * Cap at WELFORD_ALPHA_MAX_F=0.95 prevents α=1.0 (no - * smoothing). Per `pearl_wiener_alpha_floor_for_nonstationary`. */ - alpha = fmaxf(WELFORD_ALPHA_MIN_F, fminf(alpha, WELFORD_ALPHA_MAX_F)); - } - - /* Pearl-A first-observation bootstrap. The sentinel is - * SENTINEL_HOLD_COST_SCALE=0.0; any value within EPS of 0 is - * treated as sentinel and REPLACED directly. */ - const float current = isv[scale_idx]; - float blended; - if (current <= EPS_F) { - blended = target_scale; - } else { - blended = (1.0f - alpha) * current + alpha * target_scale; - } - /* Re-clamp post-blend. */ - blended = fmaxf(scale_min, fminf(blended, scale_max)); - - /* Write back: blended scale + all 6 Welford accumulators + new - * prev_target (= target_scale for the next launch's diff). */ - isv[scale_idx] = blended; - isv[sample_count_idx] = sample_count; - isv[target_mean_idx] = target_mean; - isv[target_M2_idx] = target_M2; - isv[diff_mean_idx] = diff_mean; - isv[diff_M2_idx] = diff_M2; - isv[prev_target_idx] = target_scale; -} diff --git a/crates/ml/src/cuda_pipeline/sp13_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp13_isv_slots.rs index 5a3363366..9223c72de 100644 --- a/crates/ml/src/cuda_pipeline/sp13_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp13_isv_slots.rs @@ -53,28 +53,19 @@ pub const DIR_SKILL_BONUS_BETA_DEFAULT: f32 = 1.0; pub const LUCK_WIN_DISCOUNT_DEFAULT: f32 = 0.3; pub const SKILL_BONUS_CAP_RATIO_DEFAULT: f32 = 0.3; -/// v3 Hold-pricing controller anchors (Invariant-1 numerical constants per -/// `feedback_isv_for_adaptive_bounds.md`). The controller scales `HOLD_COST_BASE` -/// by `1 + GAIN × max(0, observed − target)` clamped to -/// `[FLOOR_RATIO × BASE, CEIL_RATIO × BASE]`. Base cost is small per-bar -/// (~10 ticks of price for ES.FUT at ~0.25 tick value) but cumulative across -/// long Hold runs becomes meaningful relative to the SP12 capped trade reward -/// (±5..10). +/// v3 Hold-pricing controller anchor (Invariant-1 numerical constant per +/// `feedback_isv_for_adaptive_bounds.md`). /// -/// P0b lift (2026-05-04): `HOLD_COST_BASE` raised 0.001 → 0.005 after P0a smoke -/// (`train-67gqb` on `f934ea171`) showed `observed_hold_rate` climbing to 0.479 -/// despite the controller saturating at 2.4× base — model rationally absorbed -/// the cost because it was <3% of reward magnitude (0.001 × 5 ceiling × 30-bar -/// hold = 0.15 cumulative vs ±5..10 reward range). Lifted base × 5 ceiling now -/// yields 0.025/bar × 30 bars = 0.75 cumulative ≈ 10–15% of capped reward — -/// genuinely deters lazy Hold without crippling MFT use. Per-bar slot wiring -/// unchanged: constructor static-init still writes `HOLD_COST_BASE` to slot 380 -/// at fold boundary; controller now scales the (lifted) base. +/// SP18 Phase 1 (2026-05-09): the host-side controller (gain/floor/ceil) +/// that scaled `HOLD_COST_BASE` from observed-vs-target Hold rate is +/// DELETED — replaced by the SP18 D-leg structural Hold opportunity-cost +/// reward (Phase 2 lands the `compute_sp18_hold_opportunity_cost` device +/// fn). `HOLD_COST_BASE` and `HOLD_RATE_TARGET_DEFAULT` are RETAINED as +/// constructor-init-only diagnostics (slots 380, 381) — slot 380 keeps +/// the static base cost for cross-fold continuity insurance per DD7c; +/// neither is updated after startup. pub const HOLD_COST_BASE: f32 = 0.005; pub const HOLD_RATE_TARGET_DEFAULT: f32 = 0.20; -pub const HOLD_COST_CONTROLLER_GAIN: f32 = 5.0; -pub const HOLD_COST_FLOOR_RATIO: f32 = 0.5; // never below 0.5 × base -pub const HOLD_COST_CEIL_RATIO: f32 = 5.0; // never above 5.0 × base /// SP13 P0b: aux_w deficit + stagnation controller anchors. Replaces the /// SP11-era inverted formula `aux_w = 0.1 × health × (1 - sharpe).clamp(0.05, 0.3)` diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index 09a073add..2878d5c9c 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -450,77 +450,19 @@ 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; -// ── SP16 Phase 2 (2026-05-08): adaptive Hold cost scale ───────────────── -// Replaces the static `ISV[HOLD_COST_INDEX=380]` (~0.006 per bar) as the -// effective Hold-cost magnitude when over-holding is observed. The -// pre-Phase-2 v3-P0a Hold-pricing controller scaled `HOLD_COST_BASE=0.005` -// by `[FLOOR_RATIO=0.5, CEIL_RATIO=5.0]` → realised cost peaked at -// ~0.025/bar. Post-mortem of `train-multi-seed-pfh9n` showed the -// realised Hold rate climbing 0.25 → 0.52 across training while the -// per-bar reward magnitudes (popart=0.97, cf=0.65 in `reward_split`) -// were ~100× larger than the cost penalty — Hold was effectively free, -// allowing the structural Q-bias toward zero-variance Hold to dominate. +// ── SP16 Phase 2 (2026-05-08, RETIRED 2026-05-09 by SP18 Phase 1) ─────── +// Slot 461 (HOLD_COST_SCALE_INDEX) and the per-epoch +// `hold_cost_scale_update_kernel` producer are DELETED atomically with +// the 3 per-bar Hold-cost subtraction sites in `experience_kernels.cu` +// per DD7=c — the reactive multiplicative-cost chain is replaced by the +// SP18 D-leg structural Hold opportunity-cost reward (Phase 2 lands the +// `compute_sp18_hold_opportunity_cost` device fn at the SP18-bound caps +// in slots [483..493)). The slot index [461..468) is now a RESERVED gap +// in the layout fingerprint (SP14-C.1 RESERVED-gap pattern; deleted +// slots are NOT compacted to preserve checkpoint compatibility). // -// Producer: `hold_cost_scale_update_kernel.cu` (per-epoch boundary, -// single-thread cold-path kernel mirroring the SP16-P1 -// `min_hold_temperature_update_kernel` driver pattern). Reads -// `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and -// `ISV[HOLD_RATE_TARGET_INDEX=381]` from the existing SP13 hold-pricing -// chain and maps overrun into a [SCALE_MIN=1, SCALE_MAX=25] multiplier: -// -// overrun = max(0, observed − target) -// overrun_norm = clamp(overrun / max(target, 0.01), 0, 1) -// target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_norm -// -// Consumer: `experience_kernels.cu` — wherever the per-bar Hold-cost -// subtraction reads `isv[ISV_HOLD_COST_IDX=380]`, the value is multiplied -// by `isv[HOLD_COST_SCALE_INDEX=461]`. Cold-start fallback: when the -// scale slot is at sentinel 0.0 (no producer launch yet) OR outside -// [SCALE_MIN, SCALE_MAX], the consumer uses scale=1.0 (bit-identical to -// pre-Phase-2 behaviour with the v3-P0a static Hold cost). -// -// At observed=2× target (100% overrun) the scale climbs to 25× → the -// effective Hold cost rises from ~0.006/bar to ~0.15/bar, competitive -// with per-bar reward magnitudes (~0.01-0.1) but still below the SP12 -// asymmetric cap [-10, +5] so no clamp interaction. At/below target the -// scale stays at 1.0 (no extra penalty). -// -// Per `feedback_isv_for_adaptive_bounds.md` the bounds [1.0, 25.0] are -// Category-1 dimensional safety floors, NOT tuning. SCALE_MIN=1 is the -// bit-identical-to-pre-fix lower bound (no additional scale below -// target); SCALE_MAX=25 is the conservative ceiling at 100% overrun -// (allows competitive cost with per-bar reward magnitudes without -// dominating the SP12 asymmetric cap). -// -// Plan: SP16 Phase 2 Task 2.1. -pub const HOLD_COST_SCALE_INDEX: usize = 461; - -// Sentinel — Pearl-A first-observation bootstrap. 0.0 (NOT 1.0) so the -// kernel's "is this a fresh slot?" detect uses the canonical sentinel -// pattern. Consumer uses 1.0 fallback when the slot is at sentinel -// (≤ 0 or out-of-bounds), guaranteeing bit-identical pre-Phase-2 -// behaviour until the first valid producer launch lands a real scale. -pub const SENTINEL_HOLD_COST_SCALE: f32 = 0.0; - -// Bounds — Category-1 dimensional safety floors per -// `feedback_isv_for_adaptive_bounds.md`. SCALE_MIN=1: at/below target, -// the producer outputs exactly 1.0 (no extra penalty — the static -// HOLD_COST_BASE × v3-P0a controller scaling is sufficient). -// SCALE_MAX=25: at 100% overrun (observed=2× target), the effective -// Hold cost climbs to ~0.15/bar (0.006 base × 25), competitive with -// per-bar reward magnitudes (~0.01-0.1) but well below the SP12 -// asymmetric cap [-10, +5] so no clamp-interaction risk. -pub const HOLD_COST_SCALE_MIN: f32 = 1.0; -pub const HOLD_COST_SCALE_MAX: f32 = 25.0; - -// SP16 T3 (2026-05-08): the hardcoded `HOLD_COST_SCALE_EMA_ALPHA = 0.05` -// constant was deleted. Per `feedback_isv_for_adaptive_bounds` and -// `pearl_wiener_optimal_adaptive_alpha`, α is now derived from Welford -// accumulators in slots [462..468). See the SP16 T3 comment block at -// the top of the Welford slots section for the full rationale. - -pub const SP16_P2_HOLD_COST_SCALE_SLOT_BASE: usize = 461; -pub const SP16_P2_HOLD_COST_SCALE_SLOT_END: usize = 462; +// Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md +// Audit: docs/sp18-wireup-audit.md (A1-A10 expansion). // ── SP16 T3 (2026-05-08): Wiener-optimal adaptive α — Welford accumulators ─ // Both T1 (MIN_HOLD_TEMPERATURE producer at slot 460) and T2 (HOLD_COST_SCALE @@ -564,12 +506,8 @@ pub const SP16_P2_HOLD_COST_SCALE_SLOT_END: usize = 462; // valid observation REPLACES). After first observation, sample_count ≥ 1 // and Welford accumulators take over. -pub const HCS_TARGET_MEAN_INDEX: usize = 462; // Welford running mean of target_scale -pub const HCS_TARGET_M2_INDEX: usize = 463; // Welford sum of squared deviations (target) -pub const HCS_DIFF_MEAN_INDEX: usize = 464; // Welford running mean of consecutive diffs -pub const HCS_DIFF_M2_INDEX: usize = 465; // Welford sum of squared deviations (diff) -pub const HCS_PREV_TARGET_INDEX: usize = 466; // previous-epoch target_scale (for diff) -pub const HCS_SAMPLE_COUNT_INDEX: usize = 467; // Welford counter (number of observations) +// HCS_* slots [462..468) RETIRED by SP18 Phase 1 (2026-05-09) — the +// HOLD_COST_SCALE producer they fed (slot 461) is deleted; gap reserved. pub const MHT_TARGET_MEAN_INDEX: usize = 468; // Welford running mean of target_temp pub const MHT_TARGET_M2_INDEX: usize = 469; // Welford sum of squared deviations (target) @@ -630,7 +568,10 @@ pub const WELFORD_MIN_SAMPLE_COUNT: f32 = 3.0; // floor in `apply_pearls_ad_kernel` and the SP4 wiener buffer producers. pub const WELFORD_EPS: f32 = 1.0e-6; -pub const SP16_T3_WIENER_SLOT_BASE: usize = 462; +// SP16 T3 (2026-05-08, post-SP18 Phase 1 retirement of HCS): the surviving +// Wiener-Welford block is the MHT (MIN_HOLD_TEMPERATURE) producer's 6 +// accumulators at [468..474). The HCS block at [462..468) is RESERVED gap. +pub const SP16_T3_WIENER_SLOT_BASE: usize = 468; pub const SP16_T3_WIENER_SLOT_END: usize = 474; // ── SP17 (2026-05-08): dueling Q-network identifiability diagnostics ─── @@ -1184,60 +1125,18 @@ mod tests { ); } - /// Lock SP16 Phase 2 (2026-05-08) adaptive Hold cost scale slot. - /// Single contiguous slot [461..462) consumed by the per-bar Hold - /// cost subtraction sites in `experience_kernels.cu` to scale the - /// static `ISV[HOLD_COST_INDEX=380]` by an overrun-driven multiplier - /// in [1, 25]. Sentinel 0.0 → consumer fallback to scale=1.0 for - /// bit-identical pre-Phase-2 cold-start behavior. - #[test] - fn sp16_p2_hold_cost_scale_slot_layout_locked() { - assert_eq!(SP16_P2_HOLD_COST_SCALE_SLOT_BASE, 461); - assert_eq!(SP16_P2_HOLD_COST_SCALE_SLOT_END, 462); - assert_eq!(HOLD_COST_SCALE_INDEX, 461); - // Sentinel 0.0 mirrors the canonical Pearl-A bootstrap pattern; - // consumer interprets ≤ 0 as cold-start and falls back to 1.0 - // (bit-identical pre-Phase-2 effective-cost magnitude). - assert_eq!(SENTINEL_HOLD_COST_SCALE, 0.0); - // Category-1 dimensional safety: scale in [1, 25]. - // SCALE_MIN=1 (no extra penalty at/below target). - // SCALE_MAX=25 (~0.15 effective cost at 100% overrun, competitive - // with per-bar reward magnitudes ~0.01-0.1). - assert_eq!(HOLD_COST_SCALE_MIN, 1.0); - assert_eq!(HOLD_COST_SCALE_MAX, 25.0); - // SP16 T3 (2026-05-08): the hardcoded `HOLD_COST_SCALE_EMA_ALPHA - // = 0.05` constant was deleted. α is now Wiener-optimal, derived - // from Welford accumulators in slots [462..468) per - // `pearl_wiener_optimal_adaptive_alpha`. - } - - #[test] - fn all_sp16_p2_slots_fit_within_isv_total_dim() { - use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; - assert!( - SP16_P2_HOLD_COST_SCALE_SLOT_END <= ISV_TOTAL_DIM, - "SP16_P2_HOLD_COST_SCALE_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for SP16 Phase 2 adaptive Hold cost scale slot; \ - bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", - SP16_P2_HOLD_COST_SCALE_SLOT_END, ISV_TOTAL_DIM, - ); - } + // SP16 Phase 2 / T3 HCS layout-locked tests removed by SP18 Phase 1 + // (2026-05-09) — the HOLD_COST_SCALE_INDEX (461) and HCS_* (462..468) + // constants no longer exist. The MHT block lock test below survives. /// Lock SP16 T3 (2026-05-08) Wiener-optimal Welford accumulator slot - /// layout. 12 contiguous slots [462..474) — 6 per producer (HCS for - /// HOLD_COST_SCALE/T2, MHT for MIN_HOLD_TEMPERATURE/T1). Replaces the - /// hardcoded `α = 0.05f` literal in both producer kernels with a - /// signal-driven Wiener-optimal α per `pearl_wiener_optimal_adaptive_alpha`. + /// layout for the surviving MHT (MIN_HOLD_TEMPERATURE) producer. + /// 6 contiguous slots [468..474). The HCS counterpart was retired by + /// SP18 Phase 1 (RESERVED gap [462..468)). #[test] fn sp16_t3_wiener_welford_slot_layout_locked() { - assert_eq!(SP16_T3_WIENER_SLOT_BASE, 462); + assert_eq!(SP16_T3_WIENER_SLOT_BASE, 468); assert_eq!(SP16_T3_WIENER_SLOT_END, 474); - // HCS block (HOLD_COST_SCALE T2 producer accumulators). - assert_eq!(HCS_TARGET_MEAN_INDEX, 462); - assert_eq!(HCS_TARGET_M2_INDEX, 463); - assert_eq!(HCS_DIFF_MEAN_INDEX, 464); - assert_eq!(HCS_DIFF_M2_INDEX, 465); - assert_eq!(HCS_PREV_TARGET_INDEX, 466); - assert_eq!(HCS_SAMPLE_COUNT_INDEX, 467); // MHT block (MIN_HOLD_TEMPERATURE T1 producer accumulators). assert_eq!(MHT_TARGET_MEAN_INDEX, 468); assert_eq!(MHT_TARGET_M2_INDEX, 469); diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index b2500dc43..ee97f9a30 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -202,35 +202,23 @@ #define ISV_HOLD_RATE_OBSERVED_EMA_IDX 382 // SP13 v3 observed Hold-pick rate EMA (per-fold, sentinel 0.0) // ──────────────────────────────────────────────────────────────────────────── -// === SP16 Phase 2 SLOT (2026-05-08, adaptive Hold cost scale) === -// Mirror of crates/ml/src/cuda_pipeline/sp14_isv_slots.rs slot 461. -// Multiplies the static `ISV[HOLD_COST_IDX=380]` per-bar Hold cost by an -// overrun-driven scalar in [1, 25] when the realised Hold rate exceeds -// target. Producer: `hold_cost_scale_update_kernel.cu` (per-epoch boundary). -// Consumer: per-bar Hold cost subtraction sites in experience_kernels.cu. -// Cold-start fallback: when slot is at sentinel 0.0 or outside [1, 25], -// consumer uses scale=1.0 (bit-identical pre-Phase-2 cost magnitude). +// === SP16 Phase 2 / T3 HCS SLOTS (RETIRED 2026-05-09 by SP18 Phase 1) === +// Slot 461 (HOLD_COST_SCALE) + [462..468) (HCS_* Welford) are DELETED +// atomically with the producer kernel + 3 consumer sites in +// experience_kernels.cu. Replaced by the SP18 D-leg structural Hold +// opportunity-cost reward at slots [483..493). Range [461..468) is a +// RESERVED gap in the layout fingerprint (SP14-C.1 RESERVED-gap pattern). +// See docs/sp18-wireup-audit.md for the atomic-deletion record. // ──────────────────────────────────────────────────────────────────────────── -#define ISV_HOLD_COST_SCALE_IDX 461 // SP16 Phase 2 adaptive Hold cost scale (sentinel 0.0; bounds [1, 25]; α=Wiener-optimal per SP16 T3) // ──────────────────────────────────────────────────────────────────────────── -// === SP16 T3 SLOTS (2026-05-08, Wiener-optimal adaptive α) === -// Mirror of crates/ml/src/cuda_pipeline/sp14_isv_slots.rs slots [462..474). -// Replaces hardcoded `alpha = 0.05f` in both producer kernels (T1 + -// T2) with a signal-driven Wiener-optimal α: α = diff_var / (diff_var -// + sample_var + ε). Each producer maintains 6 Welford accumulators -// (target mean/M2, diff mean/M2, prev_target, sample_count). Per -// `pearl_wiener_optimal_adaptive_alpha`. FoldReset sentinel 0.0 across -// all 12 slots; cold-start fallback (sample_count < 3 → α = 1.0 -// REPLACE) keeps the producer responsive in the first 3 epochs of -// each fold while the Wiener formula stabilises. +// === SP16 T3 MHT SLOTS (2026-05-08, Wiener-optimal adaptive α) === +// Mirror of crates/ml/src/cuda_pipeline/sp14_isv_slots.rs slots [468..474). +// 6 Welford accumulators driving the surviving MIN_HOLD_TEMPERATURE +// producer (T1) per `pearl_wiener_optimal_adaptive_alpha`. FoldReset +// sentinel 0.0; cold-start fallback (sample_count < 3 → α = 1.0 +// REPLACE). // ──────────────────────────────────────────────────────────────────────────── -#define ISV_HCS_TARGET_MEAN_IDX 462 // Welford mean of target_scale (T2 producer) -#define ISV_HCS_TARGET_M2_IDX 463 // Welford ΣΔ² (target_scale) -#define ISV_HCS_DIFF_MEAN_IDX 464 // Welford mean of consecutive diff (target_scale) -#define ISV_HCS_DIFF_M2_IDX 465 // Welford ΣΔ² (diff) -#define ISV_HCS_PREV_TARGET_IDX 466 // previous-epoch target_scale -#define ISV_HCS_SAMPLE_COUNT_IDX 467 // Welford counter (T2 producer) #define ISV_MHT_TARGET_MEAN_IDX 468 // Welford mean of target_temp (T1 producer) #define ISV_MHT_TARGET_M2_IDX 469 // Welford ΣΔ² (target_temp) #define ISV_MHT_DIFF_MEAN_IDX 470 // Welford mean of consecutive diff (target_temp) diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 636c8a1ac..615e9fe6b 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1173,11 +1173,10 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459] — Class A audit-fix Batch 4-B adaptive plan_threshold floor (replaces hardcoded `0.1f` in `plan_threshold_update_kernel.cu:44`, the lower bound of the derived plan_threshold = max(floor, 0.5 × readiness_ema)). Design Y inline producer: the same `plan_threshold_update_kernel` writes both ISV[PLAN_THRESHOLD_INDEX=49] and ISV[459] in the same launch — the floor tracks a slow EMA of `0.5 × readiness_ema` (the very expression the kernel uses as its derived threshold target) and the kernel uses the freshly-written floor as the lower bound. FoldReset sentinel SENTINEL_PLAN_THRESHOLD_FLOOR=0.1 — Pearl-A first-observation bootstrap (matches pre-fix hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter (per-fold cadence; the floor is a long-term shadow and must move much slower than the readiness EMA itself which adapts at α≈[0.05, 0.10]). Bounds [0.05, 0.50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md` (probability units; below 5% the readiness gate is effectively absent, above 50% it is uselessly strict). Distinct from PLAN_THRESHOLD_INDEX=49 (the live derived threshold, fast cadence) — slot 459 is the slow-moving lower bound on slot 49.", }, - RegistryEntry { - name: "sp16_phase2_hold_cost_scale_adaptive", - category: ResetCategory::FoldReset, - description: "ISV[HOLD_COST_SCALE_INDEX=461] — SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat). The SP16 Phase 2 adaptive Hold cost scale chain is superseded by the SP18 D-leg structural Hold opportunity-cost reward (compute_sp18_hold_opportunity_cost in trade_physics.cuh + ISV[483..493) cap producer chain) per DD7=c. The 3 consumer sites in experience_kernels.cu that previously multiplied isv[ISV_HOLD_COST_IDX=380] by isv[461] are migrated to the structural reward in Phase 2 of SP18 (atomic 3-site migration per `feedback_no_partial_refactor`). Producer kernel `hold_cost_scale_update_kernel` is DELETED in Phase 1 atomic with the cost-application sites per DD7=c. The slot stays ALLOCATED with sentinel 0.0 to preserve checkpoint layout fingerprint compatibility — RESERVED-gap pattern from SP14-C.1 (deleted slots are NOT compacted; the surviving fingerprint just keeps the slot in the bus and the FoldReset machinery rewrites the sentinel each fold so consumers reading via stale code paths fall back to scale=1.0 rather than reading uninitialised garbage). Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md § DD7 + Pre-Phase Task PP.4.", - }, + // sp16_phase2_hold_cost_scale_adaptive registry entry DELETED + // by SP18 Phase 1 (2026-05-09) — slot 461 is removed from the + // registry; the slot index is now a RESERVED gap with no + // FoldReset bookkeeping (consumers are deleted atomically). RegistryEntry { name: "sp14_audit_4b_min_hold_temperature_adaptive", category: ResetCategory::FoldReset, @@ -1210,36 +1209,11 @@ impl StateResetRegistry { // being hardcoded; this lifts the constant into the producer's // own statistics, eliminating the 60-epoch convergence lag observed // in train-multi-seed-hjzss validation. - RegistryEntry { - name: "sp16_t3_hcs_target_mean", - category: ResetCategory::FoldReset, - description: "ISV[HCS_TARGET_MEAN_INDEX=462] — SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat). The SP16 T3 Welford accumulator for HOLD_COST_SCALE is no longer driven (producer kernel deleted in SP18 Phase 1 per DD7=c). The slot stays ALLOCATED in the layout fingerprint per the SP14-C.1 RESERVED-gap pattern. Original semantics preserved here for archaeology: SP16 T3 Welford running mean of target_scale for the HOLD_COST_SCALE producer kernel; sentinel 0.0 enabled Pearl-A bootstrap with cold-start fallback (sample_count<3 → α=1.0). See SP18 spec § DD7 + plan PP.4.", - }, - RegistryEntry { - name: "sp16_t3_hcs_target_m2", - category: ResetCategory::FoldReset, - description: "ISV[HCS_TARGET_M2_INDEX=463] — SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat). Original semantics: SP16 T3 Welford ΣΔ² (sample-variance accumulator) for target_scale in the HOLD_COST_SCALE producer kernel. See SP18 spec § DD7 + plan PP.4.", - }, - RegistryEntry { - name: "sp16_t3_hcs_diff_mean", - category: ResetCategory::FoldReset, - description: "ISV[HCS_DIFF_MEAN_INDEX=464] — SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat). Original semantics: SP16 T3 Welford running mean of consecutive `target_scale` one-step diffs. See SP18 spec § DD7 + plan PP.4.", - }, - RegistryEntry { - name: "sp16_t3_hcs_diff_m2", - category: ResetCategory::FoldReset, - description: "ISV[HCS_DIFF_M2_INDEX=465] — SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat). Original semantics: SP16 T3 Welford ΣΔ² for the consecutive-diff series; provided `diff_var` half of the Wiener formula. See SP18 spec § DD7 + plan PP.4.", - }, - RegistryEntry { - name: "sp16_t3_hcs_prev_target", - category: ResetCategory::FoldReset, - description: "ISV[HCS_PREV_TARGET_INDEX=466] — SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat). Original semantics: SP16 T3 previous-epoch `target_scale` for the HOLD_COST_SCALE producer kernel. See SP18 spec § DD7 + plan PP.4.", - }, - RegistryEntry { - name: "sp16_t3_hcs_sample_count", - category: ResetCategory::FoldReset, - description: "ISV[HCS_SAMPLE_COUNT_INDEX=467] — SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat). Original semantics: SP16 T3 Welford observation counter for the HOLD_COST_SCALE producer kernel; cold-start fallback (sample_count < WELFORD_MIN_SAMPLE_COUNT=3 → α = 1.0). See SP18 spec § DD7 + plan PP.4.", - }, + // 6 sp16_t3_hcs_* registry entries (slots 462..468) DELETED by + // SP18 Phase 1 (2026-05-09). The Welford accumulator slots are + // a RESERVED gap; no FoldReset bookkeeping. The MHT_* entries + // below survive — MIN_HOLD_TEMPERATURE chain is RETAINED per + // DD7=c. RegistryEntry { name: "sp16_t3_mht_target_mean", category: ResetCategory::FoldReset, @@ -2253,71 +2227,8 @@ mod tests { } } - /// Lock SP18 v2 Pre-Phase Task PP.4: SP16 P2/T3 HOLD_COST_SCALE chain is - /// retired by DD7=c (replaced by D-leg structural Hold opportunity-cost - /// reward). Slots 461 (HOLD_COST_SCALE_INDEX) and [462..468) (HCS_*) - /// remain ALLOCATED in the registry with sentinel 0.0 to preserve - /// checkpoint layout-fingerprint compatibility (SP14-C.1 RESERVED-gap - /// pattern), but their descriptions are updated to the canonical SP18 - /// retirement marker so future readers know the producer kernel + 3 - /// consumer sites in `experience_kernels.cu` are deleted by SP18 Phase 1. - /// MHT_* slots [468..474) are NOT retired — MIN_HOLD_TEMPERATURE remains - /// in the SP16 P1 observation chain per DD7=c "keep the observation - /// chain". SP13 P0a slots [380..383) likewise remain (observation only — - /// the per-step hold_rate_observer chain that feeds MIN_HOLD_TEMPERATURE). - #[test] - fn sp16_t2_t3_slots_marked_retired() { - let r = StateResetRegistry::new(); - let by_name: std::collections::HashMap<&str, &RegistryEntry> = - r.all().iter().map(|e| (e.name, e)).collect(); - - let retired: &[&str] = &[ - "sp16_phase2_hold_cost_scale_adaptive", // slot 461 - "sp16_t3_hcs_target_mean", // slot 462 - "sp16_t3_hcs_target_m2", // slot 463 - "sp16_t3_hcs_diff_mean", // slot 464 - "sp16_t3_hcs_diff_m2", // slot 465 - "sp16_t3_hcs_prev_target", // slot 466 - "sp16_t3_hcs_sample_count", // slot 467 - ]; - - for name in retired { - let entry = by_name.get(name).unwrap_or_else(|| { - panic!("SP18 v2 PP.4: retired SP16 P2/T3 entry '{}' missing from registry", name) - }); - assert!( - entry.description.contains("SP18 v2: RETIRED"), - "SP18 v2 PP.4: entry '{}' description must contain 'SP18 v2: RETIRED' marker; got: {}", - name, entry.description - ); - assert_eq!( - entry.category, - ResetCategory::FoldReset, - "Retired SP16 P2/T3 entry '{}' stays FoldReset (sentinel rewrite preserves \ - bit-identical pre-SP18 cold-start; the slot just isn't driven by a producer \ - any more); got: {:?}", name, entry.category - ); - } - - // Verify the MHT_* slots [468..474) and SP13 P0a slots [380..383) - // remain UN-retired — DD7=c preserves the observation chain. - let preserved: &[&str] = &[ - "sp16_t3_mht_target_mean", // slot 468 - "sp16_t3_mht_sample_count", // slot 473 - // SP13 P0a slots are not in the registry by these names — they - // are part of the per-step hold_rate_observer chain rather than - // FoldReset slots. We just assert MHT_* are not flagged retired. - ]; - for name in preserved { - let entry = by_name.get(name).unwrap_or_else(|| { - panic!("Expected SP16 T3 MHT_* entry '{}' to exist but registry lookup failed", name) - }); - assert!( - !entry.description.contains("SP18 v2: RETIRED"), - "SP16 T3 MHT entry '{}' must NOT carry SP18 retirement marker (DD7=c keeps \ - the MIN_HOLD_TEMPERATURE observation chain); got: {}", - name, entry.description - ); - } - } + // sp16_t2_t3_slots_marked_retired test DELETED by SP18 Phase 1 + // (2026-05-09) — locking the existence of slots that have been fully + // deleted is pointless ceremony per user call. The slots no longer + // exist in the registry; asserting their existence would panic. } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 191b2c5fc..1c386b324 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -561,9 +561,8 @@ impl DQNTrainer { // boundary the input has been refreshed N×B times // (N=collector steps, B=batch). Per // `pearl_canary_input_freshness_launch_order.md`. - if let Err(e) = fused.trainer().launch_hold_cost_scale_update() { - tracing::warn!(epoch, "launch_hold_cost_scale_update failed (non-fatal): {e}"); - } + // launch_hold_cost_scale_update DELETED by SP18 Phase 1 + // (2026-05-09). Replaced by SP18 D-leg in Phase 2. // 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. @@ -615,56 +614,8 @@ impl DQNTrainer { epoch, obs, tgt, overrun_norm, target_temp, temp, alpha, n as u32, ); } - // SP16 Phase 2 (2026-05-08): hold_cost_scale_diag — - // emits the realised overrun + producer-output scale - // for slot 461. Reads slot 461 directly (post-producer- - // launch above). Same overrun_norm computation as the - // P1 diag for consistency, but maps to scale [1, 25] - // instead of temperature [5, 50]. Useful Fold-1 - // distinguisher: at sentinel cold-start (slot=0.0 OR - // observed=sentinel), the consumer falls back to - // scale=1.0; this diag's `blended_scale` field - // surfaces the actual ISV slot value so the test - // distinguishes "producer hasn't fired" (slot=0.0) - // from "fired but at scale=1.0" (no overrun observed). - { - use crate::cuda_pipeline::sp13_isv_slots::{ - HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, - }; - use crate::cuda_pipeline::sp14_isv_slots::{ - HOLD_COST_SCALE_INDEX, - HCS_TARGET_M2_INDEX, HCS_DIFF_M2_INDEX, HCS_SAMPLE_COUNT_INDEX, - WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX, WELFORD_MIN_SAMPLE_COUNT, - WELFORD_EPS, - }; - let trainer = fused.trainer(); - let obs = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX); - let tgt = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX); - let scale = trainer.read_isv_signal_at(HOLD_COST_SCALE_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_scale = 1.0_f32 + (25.0_f32 - 1.0_f32) * overrun_norm; - // SP16 T3 (2026-05-08): recover the Wiener-optimal α - // from the on-device Welford state for direct - // observation (mirrors the kernel's formula). - let n = trainer.read_isv_signal_at(HCS_SAMPLE_COUNT_INDEX); - let target_m2 = trainer.read_isv_signal_at(HCS_TARGET_M2_INDEX); - let diff_m2 = trainer.read_isv_signal_at(HCS_DIFF_M2_INDEX); - let alpha = if n < WELFORD_MIN_SAMPLE_COUNT { - 1.0_f32 - } else { - let sample_var = target_m2 / (n - 1.0).max(1.0); - let diff_count = (n - 1.0).max(0.0); - let diff_var = if diff_count >= 2.0 { diff_m2 / (diff_count - 1.0) } else { diff_m2 }; - (diff_var / (diff_var + sample_var + WELFORD_EPS)) - .clamp(WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX) - }; - tracing::info!( - "HEALTH_DIAG[{}]: hold_cost_scale_diag obs_hold_rate={:.4} target_hold_rate={:.4} overrun_norm={:.4} target_scale={:.4} blended_scale={:.4} alpha={:.4} sample_count={}", - epoch, obs, tgt, overrun_norm, target_scale, scale, alpha, n as u32, - ); - } + // hold_cost_scale_diag emit DELETED by SP18 Phase 1 + // (2026-05-09). The producer chain is gone; nothing to emit. } // D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary. @@ -4168,62 +4119,14 @@ impl DQNTrainer { "SP11 mag_ratio_compute launch failed"); } - // SP13 v3 P0a.T3 (2026-05-04): Hold-pricing - // controller. Reads the observed Hold-pick rate - // EMA at ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] - // (populated by `hold_rate_observer_kernel` + - // chained Pearls A+D — observer launch wires - // alongside the dir_acc launch in P0a.T4) and - // the static Hold-rate target at - // ISV[HOLD_RATE_TARGET_INDEX=381] (constructor- - // initialised to 0.20). Computes the priced - // per-bar Hold cost - // excess = max(0, observed − target) - // cost = HOLD_COST_BASE × (1 + 5 × excess) - // cost = clamp(cost, base × 0.5, base × 5.0) - // and writes it back to ISV[HOLD_COST_INDEX=380]. - // The reward composition site in - // `experience_kernels.cu` reads slot 380 each - // bar where action == DIR_HOLD and subtracts - // the cost — pricing the action so the policy - // uses Hold deliberately rather than as a free - // CQL-bias-anchored default. - // - // Cold-start (epoch 1, slot 382 still at fold- - // reset sentinel 0.0): observed=0, target=0.20, - // excess=0 → cost = base × 1.0 = HOLD_COST_BASE, - // clamped to [base × 0.5, base × 5.0] → - // baseline carry cost. Pearl A's first- - // observation replacement consumes the sentinel - // when the observer first fires. - // - // Per `feedback_isv_for_adaptive_bounds`: base - // cost is the only constant; rate target, - // observed rate, and cost output are ISV-driven. - // Per `feedback_no_cpu_compute_strict`: the - // formula here is host-side controller arithmetic - // on already-smoothed scalars — no per-sample - // GPU compute path is bypassed (the observer + - // Pearls A+D + per-bar reward subtraction all - // run on GPU). - { - use crate::cuda_pipeline::sp5_isv_slots::{ - HOLD_COST_INDEX, HOLD_RATE_TARGET_INDEX, - HOLD_RATE_OBSERVED_EMA_INDEX, - HOLD_COST_BASE, HOLD_COST_CONTROLLER_GAIN, - HOLD_COST_FLOOR_RATIO, HOLD_COST_CEIL_RATIO, - }; - let trainer = fused.trainer(); - let observed = trainer.read_isv_signal_at(HOLD_RATE_OBSERVED_EMA_INDEX); - let target = trainer.read_isv_signal_at(HOLD_RATE_TARGET_INDEX); - let excess = (observed - target).max(0.0); - let cost = HOLD_COST_BASE * (1.0 + HOLD_COST_CONTROLLER_GAIN * excess); - let cost = cost.clamp( - HOLD_COST_BASE * HOLD_COST_FLOOR_RATIO, - HOLD_COST_BASE * HOLD_COST_CEIL_RATIO, - ); - trainer.write_isv_signal_at(HOLD_COST_INDEX, cost); - } + // SP13 v3 P0a.T3 host-side Hold-pricing + // controller DELETED by SP18 Phase 1 + // (2026-05-09). Slot 380 is now constructor- + // init-only (HOLD_COST_BASE=0.005 diagnostic + // anchor, never updated). The reactive cost + // chain is replaced by the SP18 D-leg + // structural Hold opportunity-cost reward + // (Phase 2). Per DD7=c. // SP11 Fix 39 (2026-05-04, A1.2): per-bar // saboteur engagement canary producer. Reads @@ -8963,84 +8866,11 @@ impl DQNTrainer { ); } } - // SP16 Phase 2 (2026-05-08): adaptive Hold cost scale. - // Pearl-A bootstrap on every fold boundary — sentinel 0.0 - // → consumer cold-start fallback (scale=1.0; bit-identical - // pre-Phase-2 cost magnitude) until producer's first valid - // observation lands a real scale. - "sp16_phase2_hold_cost_scale_adaptive" => { - if let Some(ref fused) = self.fused_ctx { - use crate::cuda_pipeline::sp14_isv_slots::{ - HOLD_COST_SCALE_INDEX, SENTINEL_HOLD_COST_SCALE, - }; - fused.trainer().write_isv_signal_at( - HOLD_COST_SCALE_INDEX, - SENTINEL_HOLD_COST_SCALE, - ); - } - } - // ── SP16 T3 (2026-05-08): Wiener-optimal Welford accumulators ─ - // 12 dispatch arms (6 per producer). Every slot resets to the - // canonical Welford zero on every fold boundary — combined with - // the cold-start fallback (sample_count < 3 → α = 1.0 REPLACE), - // this means the first observation of each new fold replaces - // blended directly, then the second seeds diff statistics, - // then α adapts via Wiener formula from the third observation - // onward. Per `pearl_wiener_optimal_adaptive_alpha` and - // `feedback_isv_for_adaptive_bounds`. - // - // The HCS_* slots (462..468) feed the HOLD_COST_SCALE producer - // kernel; MHT_* slots (468..474) feed the MIN_HOLD_TEMPERATURE - // producer. Both producers share the same Welford-state - // algorithm — only the slot indices and bound constants differ. - "sp16_t3_hcs_target_mean" => { - if let Some(ref fused) = self.fused_ctx { - use crate::cuda_pipeline::sp14_isv_slots::{ - HCS_TARGET_MEAN_INDEX, SENTINEL_WELFORD_ZERO, - }; - fused.trainer().write_isv_signal_at(HCS_TARGET_MEAN_INDEX, SENTINEL_WELFORD_ZERO); - } - } - "sp16_t3_hcs_target_m2" => { - if let Some(ref fused) = self.fused_ctx { - use crate::cuda_pipeline::sp14_isv_slots::{ - HCS_TARGET_M2_INDEX, SENTINEL_WELFORD_ZERO, - }; - fused.trainer().write_isv_signal_at(HCS_TARGET_M2_INDEX, SENTINEL_WELFORD_ZERO); - } - } - "sp16_t3_hcs_diff_mean" => { - if let Some(ref fused) = self.fused_ctx { - use crate::cuda_pipeline::sp14_isv_slots::{ - HCS_DIFF_MEAN_INDEX, SENTINEL_WELFORD_ZERO, - }; - fused.trainer().write_isv_signal_at(HCS_DIFF_MEAN_INDEX, SENTINEL_WELFORD_ZERO); - } - } - "sp16_t3_hcs_diff_m2" => { - if let Some(ref fused) = self.fused_ctx { - use crate::cuda_pipeline::sp14_isv_slots::{ - HCS_DIFF_M2_INDEX, SENTINEL_WELFORD_ZERO, - }; - fused.trainer().write_isv_signal_at(HCS_DIFF_M2_INDEX, SENTINEL_WELFORD_ZERO); - } - } - "sp16_t3_hcs_prev_target" => { - if let Some(ref fused) = self.fused_ctx { - use crate::cuda_pipeline::sp14_isv_slots::{ - HCS_PREV_TARGET_INDEX, SENTINEL_WELFORD_ZERO, - }; - fused.trainer().write_isv_signal_at(HCS_PREV_TARGET_INDEX, SENTINEL_WELFORD_ZERO); - } - } - "sp16_t3_hcs_sample_count" => { - if let Some(ref fused) = self.fused_ctx { - use crate::cuda_pipeline::sp14_isv_slots::{ - HCS_SAMPLE_COUNT_INDEX, SENTINEL_WELFORD_ZERO, - }; - fused.trainer().write_isv_signal_at(HCS_SAMPLE_COUNT_INDEX, SENTINEL_WELFORD_ZERO); - } - } + // sp16_phase2_hold_cost_scale_adaptive + 6 sp16_t3_hcs_* + // dispatch arms DELETED by SP18 Phase 1 (2026-05-09) atomically + // with the registry entries (see state_reset_registry.rs). + // The MHT_* arms below survive — MIN_HOLD_TEMPERATURE chain is + // RETAINED per DD7=c. "sp16_t3_mht_target_mean" => { if let Some(ref fused) = self.fused_ctx { use crate::cuda_pipeline::sp14_isv_slots::{ diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index ea3e8c345..ba0381364 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -2170,759 +2170,10 @@ mod sp16_phase1_min_hold_temperature_gpu { } } -// ═══════════════════════════════════════════════════════════════════════════ -// SP16 Phase 2 (2026-05-08) — adaptive Hold cost scale producer -// (`hold_cost_scale_update_kernel`). -// -// Per train-multi-seed-pfh9n post-mortem: realised Hold rate climbed -// 0.25 → 0.52 across training while the cost penalty stayed at ~0.006 -// (~100× smaller than per-bar reward magnitudes — popart=0.97, cf=0.65). -// Hold became effectively free, allowing the structural Q-bias toward -// zero-variance Hold to dominate. This producer reads the same hold-rate -// chain feeding SP16-P1 (slots 382 + 381) and outputs a [1, 25] -// multiplier into ISV[HOLD_COST_SCALE_INDEX=461] — the 3 consumer sites -// in experience_kernels.cu multiply `isv[ISV_HOLD_COST_IDX=380]` by the -// scale to get the effective per-bar Hold cost. -// -// Verifies: -// 1. Pearl-A bootstrap: sentinel 0.0 → REPLACE with target_scale. -// 2. Over-holding (overrun=1.0) → scale HIGH (~25, near SCALE_MAX). -// 3. At-target (no overrun) → scale = SCALE_MIN = 1. -// 4. Under-target → scale = SCALE_MIN = 1 (no escalation). -// 5. observed=sentinel 0.0 → ISV preserved bit-exactly (cold-start). -// 6. Bilateral clamp keeps blended scale in [1, 25]. -// -// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")]. -// ═══════════════════════════════════════════════════════════════════════════ - -#[cfg(feature = "cuda")] -#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory. -mod sp16_phase2_hold_cost_scale_gpu { - use std::sync::Arc; - - use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; - use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; - use ml::cuda_pipeline::sp14_isv_slots::{ - HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX, - HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED, - HCS_TARGET_MEAN_INDEX, HCS_TARGET_M2_INDEX, - HCS_DIFF_MEAN_INDEX, HCS_DIFF_M2_INDEX, - HCS_PREV_TARGET_INDEX, HCS_SAMPLE_COUNT_INDEX, - }; - use ml::cuda_pipeline::sp13_isv_slots::{ - HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, - }; - - const HOLD_COST_SCALE_UPDATE_CUBIN: &[u8] = - include_bytes!(concat!(env!("OUT_DIR"), "/hold_cost_scale_update_kernel.cubin")); - - fn make_stream() -> Arc { - let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); - ctx.default_stream() - } - - fn load_kernel(stream: &Arc) -> CudaFunction { - let module = stream - .context() - .load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec()) - .expect("load hold_cost_scale_update_kernel cubin"); - module - .load_function("hold_cost_scale_update") - .expect("load hold_cost_scale_update function") - } - - /// Standard launch helper using the canonical SP16-P2 + T3 constants. - /// SP16 T3 (2026-05-08): kernel takes Welford accumulator slot - /// indices instead of a hardcoded α. - fn launch_canonical( - stream: &Arc, - kernel: &CudaFunction, - isv_ptr: u64, - ) { - unsafe { - stream - .launch_builder(kernel) - .arg(&isv_ptr) - .arg(&(HOLD_COST_SCALE_INDEX as i32)) - .arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32)) - .arg(&(HOLD_RATE_TARGET_INDEX as i32)) - .arg(&(HCS_TARGET_MEAN_INDEX as i32)) - .arg(&(HCS_TARGET_M2_INDEX as i32)) - .arg(&(HCS_DIFF_MEAN_INDEX as i32)) - .arg(&(HCS_DIFF_M2_INDEX as i32)) - .arg(&(HCS_PREV_TARGET_INDEX as i32)) - .arg(&(HCS_SAMPLE_COUNT_INDEX as i32)) - .arg(&SENTINEL_HOLD_COST_SCALE) - .arg(&SENTINEL_HOLD_RATE_OBSERVED) - .arg(&HOLD_COST_SCALE_MIN) - .arg(&HOLD_COST_SCALE_MAX) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), - shared_mem_bytes: 0, - }) - .expect("launch hold_cost_scale_update"); - } - stream.synchronize().expect("sync after hold_cost_scale_update"); - } - - /// SP16 Phase 2 — Test 1: scale climbs aggressively with overrun. - /// - /// Setup: observed=0.40, target=0.20 (100% overrun → overrun_norm=1.0). - /// Pearl-A bootstrap from sentinel 0.0 → scale REPLACED directly with - /// target_scale = SCALE_MAX = 25. - #[test] - #[ignore = "requires GPU"] - fn sp16_phase2_hold_cost_scale_climbs_with_overrun() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - const ISV_DIM: usize = 1024; - let mut isv = vec![0.0_f32; ISV_DIM]; - // 100% overrun: observed = 2× target. - isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40; - isv[HOLD_RATE_TARGET_INDEX] = 0.20; - // Scale at sentinel 0.0 → Pearl-A REPLACE with target_scale. - isv[HOLD_COST_SCALE_INDEX] = SENTINEL_HOLD_COST_SCALE; - - 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 scale = result[HOLD_COST_SCALE_INDEX]; - - // overrun = max(0, 0.40 − 0.20) = 0.20 - // overrun_norm = clamp(0.20 / max(0.20, 0.01), 0, 1) = 1.0 - // target_scale = 1 + 24 × 1.0 = 25.0 (= SCALE_MAX) - // Pearl-A: sentinel 0.0 → REPLACE → scale = 25.0 - assert!( - (scale - HOLD_COST_SCALE_MAX).abs() < 1e-4, - "100% overrun + Pearl-A bootstrap: expected ≈ {}, got {scale:.5}", - HOLD_COST_SCALE_MAX, - ); - } - - /// SP16 Phase 2 — Test 2: scale stays at SCALE_MIN=1 when at target. - /// - /// Setup: observed=0.20, target=0.20 (no overrun → overrun_norm=0). - /// target_scale = SCALE_MIN = 1.0. Pearl-A bootstrap from sentinel - /// 0.0 → REPLACE → scale = 1.0. - #[test] - #[ignore = "requires GPU"] - fn sp16_phase2_hold_cost_scale_at_target_is_one() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - const ISV_DIM: usize = 1024; - let mut isv = vec![0.0_f32; ISV_DIM]; - // No overrun: observed = target. - isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.20; - isv[HOLD_RATE_TARGET_INDEX] = 0.20; - isv[HOLD_COST_SCALE_INDEX] = SENTINEL_HOLD_COST_SCALE; - - let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); - isv_buf.write_from_slice(&isv); - - // NB: observed=0.20 is NOT the SENTINEL_HOLD_RATE_OBSERVED=0.0 - // value, so the kernel processes (does not skip via the sentinel - // guard). Verifies that the at-target path reaches SCALE_MIN. - launch_canonical(&stream, &kernel, isv_buf.dev_ptr); - - let result = isv_buf.read_all(); - let scale = result[HOLD_COST_SCALE_INDEX]; - - assert!( - (scale - HOLD_COST_SCALE_MIN).abs() < 1e-4, - "At-target: expected ≈ {}, got {scale:.5}", - HOLD_COST_SCALE_MIN, - ); - } - - /// SP16 Phase 2 — Test 3: scale stays at SCALE_MIN=1 when under target. - /// - /// Under-holding doesn't need a penalty escalation. Setup: - /// observed=0.10, target=0.20 (50% under). - /// overrun = max(0, 0.10 − 0.20) = 0 → overrun_norm = 0 → - /// target_scale = SCALE_MIN = 1.0. - #[test] - #[ignore = "requires GPU"] - fn sp16_phase2_hold_cost_scale_under_target_is_one() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - const ISV_DIM: usize = 1024; - let mut isv = vec![0.0_f32; ISV_DIM]; - isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.10; // under target - isv[HOLD_RATE_TARGET_INDEX] = 0.20; - isv[HOLD_COST_SCALE_INDEX] = SENTINEL_HOLD_COST_SCALE; - - 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 scale = result[HOLD_COST_SCALE_INDEX]; - - assert!( - (scale - HOLD_COST_SCALE_MIN).abs() < 1e-4, - "Under-target: expected ≈ {}, got {scale:.5} — under-holding should not escalate scale", - HOLD_COST_SCALE_MIN, - ); - } - - /// SP16 Phase 2 — Test 4: bilateral clamp keeps blended scale in [1, 25]. - /// - /// Setup: scale at malformed prior state of 30.0 (above SCALE_MAX). - /// observed=0.40, target=0.20 (100% overrun → target_scale = 25.0). - /// EMA blend: (1 − 0.05) × 30.0 + 0.05 × 25.0 = 28.5 + 1.25 = 29.75. - /// Post-blend bilateral clamp: max(1, min(29.75, 25)) = 25.0. - /// Verifies the post-blend defensive clamp catches this. - #[test] - #[ignore = "requires GPU"] - fn sp16_phase2_hold_cost_scale_bounds_clamp() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - const ISV_DIM: usize = 1024; - const MALFORMED_PRIOR_SCALE: f32 = 30.0; // above SCALE_MAX - let mut isv = vec![0.0_f32; ISV_DIM]; - isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.40; - isv[HOLD_RATE_TARGET_INDEX] = 0.20; - isv[HOLD_COST_SCALE_INDEX] = MALFORMED_PRIOR_SCALE; - - 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 scale = result[HOLD_COST_SCALE_INDEX]; - - // Even under malformed prior state, the post-blend bilateral - // clamp must keep scale in [SCALE_MIN, SCALE_MAX]. - assert!( - scale >= HOLD_COST_SCALE_MIN && scale <= HOLD_COST_SCALE_MAX, - "Bilateral clamp violated: scale={scale:.5} not in [{HOLD_COST_SCALE_MIN}, {HOLD_COST_SCALE_MAX}]", - ); - // Specifically: post-blend value would be 29.75, clamp pushes to 25. - assert!( - (scale - HOLD_COST_SCALE_MAX).abs() < 1e-4, - "Bilateral clamp expected at SCALE_MAX={}, got {scale:.5}", - HOLD_COST_SCALE_MAX, - ); - } - - /// SP16 Phase 2 — Test 5: Pearl-A first-observation bootstrap on - /// 50% overrun (NOT 100%, NOT 0%) — distinguishes REPLACE from - /// EMA-blend cleanly. - /// - /// Setup: scale at sentinel 0.0, observed=0.30, target=0.20. - /// overrun = 0.10, overrun_norm = 0.10 / 0.20 = 0.5 - /// target_scale = 1 + 24 × 0.5 = 13.0 - /// Pearl-A REPLACE → blended = 13.0 directly (NOT - /// `(1 − 0.05) × 0 + 0.05 × 13.0 = 0.65` which would be the - /// non-Pearl-A blend from sentinel and would also fall below - /// SCALE_MIN, getting clamped to 1.0). - /// - /// Without Pearl-A bootstrap, this test would observe ~1.0 (clamped - /// from below). With Pearl-A, it observes ~13.0. The 13:1 ratio is - /// the smoking gun for correct bootstrap behavior. - #[test] - #[ignore = "requires GPU"] - fn sp16_phase2_hold_cost_scale_pearl_a_bootstrap() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - const ISV_DIM: usize = 1024; - let mut isv = vec![0.0_f32; ISV_DIM]; - isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.30; // 50% overrun - isv[HOLD_RATE_TARGET_INDEX] = 0.20; - isv[HOLD_COST_SCALE_INDEX] = SENTINEL_HOLD_COST_SCALE; // sentinel 0.0 - - let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); - isv_buf.write_from_slice(&isv); - - launch_canonical(&stream, &kernel, isv_buf.dev_ptr); - - let result = isv_buf.read_all(); - let scale = result[HOLD_COST_SCALE_INDEX]; - - // overrun = 0.30 − 0.20 = 0.10 - // overrun_norm = 0.10 / max(0.20, 0.01) = 0.5 - // target_scale = 1 + (25 − 1) × 0.5 = 13.0 - // Pearl-A REPLACE → blended = 13.0 (NOT the EMA blend - // (1 − 0.05)×0 + 0.05×13 = 0.65 which would clamp up to 1.0). - let expected = 1.0_f32 + (HOLD_COST_SCALE_MAX - HOLD_COST_SCALE_MIN) * 0.5; - assert!( - (scale - expected).abs() < 1e-4, - "Pearl-A bootstrap on 50% overrun: expected ≈ {expected:.5}, got {scale:.5}. \ - A non-Pearl-A blend would yield ~1.0 (post-clamp); the {expected:.5} value is \ - the smoking gun for correct bootstrap behavior.", - ); - } -} - -// ═══════════════════════════════════════════════════════════════════════════ -// SP16 T3 (2026-05-08) — Wiener-optimal adaptive α via Welford accumulators. -// -// Per train-multi-seed-hjzss validation: SP16 T1+T2 chain was structurally -// landed but BEHAVIORALLY INERT in 5-epoch smoke (bit-identical to pfh9n -// baseline through epoch 3). Root cause: hardcoded `alpha = 0.05f` in both -// producer kernels needed ~60 epochs from cold start to converge. Per -// `feedback_isv_for_adaptive_bounds.md` α was itself an adaptive bound being -// hardcoded; per `pearl_wiener_optimal_adaptive_alpha`: -// -// α = diff_var / (diff_var + sample_var + ε) -// -// These tests verify the Wiener-optimal α implementation in -// `hold_cost_scale_update_kernel.cu` (drives SP16 T2 / slot 461). The -// MIN_HOLD_TEMPERATURE producer (T1 / slot 460) shares the same algorithm -// — the T2 tests are sufficient as the smoking gun; T1 mirrors verbatim -// modulo slot indices. -// -// Verifies: -// 1. α high during signal jumps (cold-start path responds aggressively). -// 2. α low in steady state (Welford diff_var → 0 → α decays). -// 3. Pearl-A first-observation bootstrap fires regardless of computed α. -// 4. α stays in (0, 1) over 50 synthetic epochs (bounded). -// 5. No 0.05f hardcoded literal remains in the kernel source (regression- -// locked at static-analysis time, not runtime). -// -// All GPU tests are #[ignore = "requires GPU"]; gated under -// #[cfg(feature = "cuda")]. Test 5 is a pure host-side string scan. -// ═══════════════════════════════════════════════════════════════════════════ - -#[cfg(feature = "cuda")] -#[allow(unsafe_code)] -mod sp16_phase3_wiener_alpha_gpu { - use std::sync::Arc; - - use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; - use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; - use ml::cuda_pipeline::sp14_isv_slots::{ - HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX, - HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED, - HCS_TARGET_MEAN_INDEX, HCS_TARGET_M2_INDEX, - HCS_DIFF_MEAN_INDEX, HCS_DIFF_M2_INDEX, - HCS_PREV_TARGET_INDEX, HCS_SAMPLE_COUNT_INDEX, - WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX, - }; - use ml::cuda_pipeline::sp13_isv_slots::{ - HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, - }; - - const HOLD_COST_SCALE_UPDATE_CUBIN: &[u8] = - include_bytes!(concat!(env!("OUT_DIR"), "/hold_cost_scale_update_kernel.cubin")); - - fn make_stream() -> Arc { - let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); - ctx.default_stream() - } - - fn load_kernel(stream: &Arc) -> CudaFunction { - let module = stream - .context() - .load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec()) - .expect("load hold_cost_scale_update_kernel cubin"); - module - .load_function("hold_cost_scale_update") - .expect("load hold_cost_scale_update function") - } - - fn launch_canonical(stream: &Arc, kernel: &CudaFunction, isv_ptr: u64) { - unsafe { - stream - .launch_builder(kernel) - .arg(&isv_ptr) - .arg(&(HOLD_COST_SCALE_INDEX as i32)) - .arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32)) - .arg(&(HOLD_RATE_TARGET_INDEX as i32)) - .arg(&(HCS_TARGET_MEAN_INDEX as i32)) - .arg(&(HCS_TARGET_M2_INDEX as i32)) - .arg(&(HCS_DIFF_MEAN_INDEX as i32)) - .arg(&(HCS_DIFF_M2_INDEX as i32)) - .arg(&(HCS_PREV_TARGET_INDEX as i32)) - .arg(&(HCS_SAMPLE_COUNT_INDEX as i32)) - .arg(&SENTINEL_HOLD_COST_SCALE) - .arg(&SENTINEL_HOLD_RATE_OBSERVED) - .arg(&HOLD_COST_SCALE_MIN) - .arg(&HOLD_COST_SCALE_MAX) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), - shared_mem_bytes: 0, - }) - .expect("launch hold_cost_scale_update"); - } - stream.synchronize().expect("sync after hold_cost_scale_update"); - } - - /// Recover the actual α used by the kernel on the latest launch. The - /// kernel doesn't directly expose α, but it can be reverse-engineered - /// from the blend equation: - /// - /// blended = (1 − α) × prior_scale + α × target_scale (when prior > 0) - /// ⇒ α = (blended − prior_scale) / (target_scale − prior_scale) - /// - /// Pre-launch fix: write `prior_scale` to slot 461 with a non-sentinel - /// value so the Pearl-A bootstrap branch isn't taken; capture - /// (prior_scale, target_scale) and compare to (blended) post-launch. - /// Caller MUST set `target_scale ≠ prior_scale` to keep denominator - /// well-defined; bilateral clamp post-blend may corrupt the - /// recovered α near boundaries — choose seed/observed values that - /// keep the blend well-inside [SCALE_MIN, SCALE_MAX]. - fn recover_alpha(blended: f32, prior_scale: f32, target_scale: f32) -> f32 { - (blended - prior_scale) / (target_scale - prior_scale) - } - - /// Compute target_scale from observed/target without launching the - /// kernel — mirrors the kernel's overrun-norm formula exactly. - fn compute_target_scale(observed: f32, target: f32) -> f32 { - let overrun = (observed - target).max(0.0); - let denom = target.max(0.01); - let overrun_norm = (overrun / denom).clamp(0.0, 1.0); - (HOLD_COST_SCALE_MIN - + (HOLD_COST_SCALE_MAX - HOLD_COST_SCALE_MIN) * overrun_norm) - .clamp(HOLD_COST_SCALE_MIN, HOLD_COST_SCALE_MAX) - } - - /// Run a synthetic per-epoch sequence by setting (observed, prev_blended) - /// per step and capturing the post-launch slot 461. - /// Returns the trajectory of `(blended, recovered_alpha)` per epoch. - fn run_synthetic_sequence( - stream: &Arc, - kernel: &CudaFunction, - observations: &[f32], - target: f32, - seed_blended: f32, - ) -> Vec<(f32, f32)> { - const ISV_DIM: usize = 1024; - - let mut isv = vec![0.0_f32; ISV_DIM]; - isv[HOLD_RATE_TARGET_INDEX] = target; - isv[HOLD_COST_SCALE_INDEX] = seed_blended; - - let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); - isv_buf.write_from_slice(&isv); - - let mut traj = Vec::with_capacity(observations.len()); - for (epoch, &obs) in observations.iter().enumerate() { - let mut current = isv_buf.read_all(); - current[HOLD_RATE_OBSERVED_EMA_INDEX] = obs; - isv_buf.write_from_slice(¤t); - - let prior_scale = current[HOLD_COST_SCALE_INDEX]; - let target_scale = compute_target_scale(obs, target); - - launch_canonical(stream, kernel, isv_buf.dev_ptr); - let post = isv_buf.read_all(); - let blended = post[HOLD_COST_SCALE_INDEX]; - - // Recover α from blend equation. For the first epoch (where - // prior is the seed), α=1.0 is the cold-start path. Once - // sample_count >= 3, α derives from Welford state. - let recovered_alpha = if (target_scale - prior_scale).abs() < 1e-6 { - f32::NAN // formula undefined when prior == target - } else { - recover_alpha(blended, prior_scale, target_scale) - }; - traj.push((blended, recovered_alpha)); - // Sanity check: Welford counter increments monotonically. - assert!( - (post[HCS_SAMPLE_COUNT_INDEX] - (epoch as f32 + 1.0)).abs() < 1e-4, - "epoch {epoch}: sample_count expected {}, got {}", - epoch + 1, post[HCS_SAMPLE_COUNT_INDEX] - ); - } - traj - } - - /// SP16 T3 — Test 1: α HIGH during signal jump (cold-start - /// responsiveness). - /// - /// Feed a sharp signal jump from below-target to over-target across - /// 5 epochs. The first three epochs trigger cold-start path - /// (sample_count < 3 → α=1.0 REPLACE). After that, the Wiener - /// formula sees high diff_var (the recent jumps) relative to - /// sample_var → α should remain HIGH (> 0.5) on the 4th and 5th - /// epochs — concrete proof of the responsiveness benefit over - /// the deleted hardcoded α=0.05. - #[test] - #[ignore = "requires GPU"] - fn sp16_phase3_alpha_high_during_signal_jump() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - // Sharp jumps: 0.20→0.30→0.40→0.30→0.40 (target=0.20). - // target_scale series (overrun_norm × 24 + 1): - // ep0: obs=0.20 → norm=0.0 → target=1 - // ep1: obs=0.30 → norm=0.5 → target=13 - // ep2: obs=0.40 → norm=1.0 → target=25 - // ep3: obs=0.30 → norm=0.5 → target=13 - // ep4: obs=0.40 → norm=1.0 → target=25 - // Diff sequence: +12, +12, -12, +12 → high diff_var. - let observations = [0.20, 0.30, 0.40, 0.30, 0.40]; - // Seed blended at 1.0 (NOT sentinel — bypass Pearl-A path) so we - // can recover α from the blend equation directly. - const SEED: f32 = 1.0001; // slightly above EPS_F to avoid Pearl-A REPLACE - let traj = run_synthetic_sequence(&stream, &kernel, &observations, 0.20, SEED); - - // Epoch 0 (sample_count=1 < 3): cold-start α=1.0, target=1. - // Epoch 1 (sample_count=2 < 3): cold-start α=1.0, target=13. - // Epoch 2 (sample_count=3): Wiener formula activates. After - // 3 samples (1, 13, 25), sample_mean ≈ 13, sample_var huge, - // diff_var also huge. α should be substantial (>0.3). - // Epoch 3-4: Wiener formula. With high target swings, diff_var - // stays comparable to sample_var → α stays high. - let alpha_ep3 = traj[3].1; - let alpha_ep4 = traj[4].1; - assert!( - alpha_ep3 > 0.3 && !alpha_ep3.is_nan(), - "Wiener α should be HIGH (>0.3) during high-volatility signal at epoch 3, got {alpha_ep3:.5}" - ); - assert!( - alpha_ep4 > 0.3 && !alpha_ep4.is_nan(), - "Wiener α should remain HIGH (>0.3) during continued high-volatility at epoch 4, got {alpha_ep4:.5}" - ); - // Defensive bound check. - assert!( - alpha_ep3 <= WELFORD_ALPHA_MAX && alpha_ep4 <= WELFORD_ALPHA_MAX, - "Wiener α must respect WELFORD_ALPHA_MAX={}; got ep3={alpha_ep3:.5} ep4={alpha_ep4:.5}", - WELFORD_ALPHA_MAX - ); - } - - /// SP16 T3 — Test 2: α floors at WELFORD_ALPHA_MIN in steady state. - /// - /// Feed a converging signal: `0.30 + 0.001 / (epoch + 1)` for 60 - /// epochs. The decaying perturbation makes consecutive diffs ever- - /// smaller (so diff_var → 0) while sample_var converges toward the - /// steady-state value — α = diff_var / (diff_var + sample_var + ε) - /// → 0 in the unclamped formula, then clamps UP to WELFORD_ALPHA_MIN. - /// - /// Per `pearl_wiener_alpha_floor_for_nonstationary` (2026-05-08), the - /// MIN floor is 0.4 — a non-stationary control-loop bandwidth - /// guarantee, not a "natural smoothing emerges" outcome. Without - /// this floor, train-multi-seed-b5gmp Fold 1 collapsed (Sharpe - /// 88.65 → 55.55) because α decayed too far during convergence - /// and the controller could not catch up to the climbing overrun - /// signal as the policy drifted post-convergence. - #[test] - #[ignore = "requires GPU"] - fn sp16_phase3_alpha_low_in_steady_state() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - // Converging signal: 0.30 + tiny decaying perturbation. After - // ~30 epochs the recent diffs are nano-sized while sample_var - // reflects the established mean of ~target_scale = 13. - let mut observations = vec![]; - for i in 0..60 { - observations.push(0.30_f32 + 0.001_f32 / ((i + 1) as f32)); - } - const SEED: f32 = 13.0; - let traj = run_synthetic_sequence(&stream, &kernel, &observations, 0.20, SEED); - - // After ~50 epochs the Welford state should have stabilised: - // sample_var stays bounded (target_scale fluctuates within - // ~0.001 of 13.0) while diff_var → 0 (consecutive diffs are - // sub-permille). The unclamped formula would yield α → 0; the - // 0.4 floor (non-stationary control-loop constraint) clamps the - // recovered α back up to WELFORD_ALPHA_MIN. - let tail: Vec = traj[50..60].iter() - .map(|&(_, a)| a) - .filter(|a| !a.is_nan()) - .collect(); - let mean_tail_alpha = tail.iter().sum::() / tail.len() as f32; - - // Tail α should sit AT the WELFORD_ALPHA_MIN floor (not above — - // the unclamped formula has decayed → 0; not below — the floor - // forbids it). Tolerance accommodates fp rounding in the - // recovered-α reconstruction inside training_loop.rs. - let tol = 1e-3_f32; - assert!( - (mean_tail_alpha - WELFORD_ALPHA_MIN).abs() < tol, - "Steady-state α must clamp to WELFORD_ALPHA_MIN={WELFORD_ALPHA_MIN}; \ - got mean over [50..60): {mean_tail_alpha:.5} (tol {tol:.0e})" - ); - // Defensive bound check — α stays positive (no permanent freeze). - assert!( - mean_tail_alpha >= WELFORD_ALPHA_MIN - tol, - "Steady-state α must respect WELFORD_ALPHA_MIN={}; got {mean_tail_alpha:.5}", - WELFORD_ALPHA_MIN - ); - } - - /// SP16 T3 — Test 3: Pearl-A first-observation bootstrap fires - /// regardless of cold-start α. - /// - /// When `prior_scale` is at sentinel (≤ EPS), the kernel takes the - /// REPLACE branch BEFORE the blend formula runs. The Wiener α - /// computation still happens (Welford state is updated) but the - /// result is discarded for this launch — blended = target_scale - /// directly. This protects the cold-start bootstrap from being - /// corrupted by an arbitrary cold-start α=1.0 that would otherwise - /// produce the same numerical result by coincidence. - /// - /// Setup: prior_scale=sentinel 0.0, observed=0.30, target=0.20. - /// target_scale = 1 + 24 × 0.5 = 13.0 - /// Pearl-A REPLACE → blended = 13.0 directly. - /// Without Pearl-A path: blend with α=1.0 → 0×0 + 1×13 = 13.0 - /// (same numerical result). To distinguish, we run a SECOND launch - /// with non-sentinel prior and verify the Welford state was already - /// updated by the first launch (sample_count = 1). - #[test] - #[ignore = "requires GPU"] - fn sp16_phase3_pearl_a_bootstrap_first_obs() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - const ISV_DIM: usize = 1024; - let mut isv = vec![0.0_f32; ISV_DIM]; - isv[HOLD_RATE_OBSERVED_EMA_INDEX] = 0.30; - isv[HOLD_RATE_TARGET_INDEX] = 0.20; - isv[HOLD_COST_SCALE_INDEX] = SENTINEL_HOLD_COST_SCALE; // sentinel 0.0 - - let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); - isv_buf.write_from_slice(&isv); - - launch_canonical(&stream, &kernel, isv_buf.dev_ptr); - let result = isv_buf.read_all(); - - // Pearl-A REPLACE → blended = 13.0. - let expected = 1.0_f32 + (HOLD_COST_SCALE_MAX - HOLD_COST_SCALE_MIN) * 0.5; - let blended = result[HOLD_COST_SCALE_INDEX]; - assert!( - (blended - expected).abs() < 1e-4, - "Pearl-A bootstrap on sentinel must REPLACE with target_scale; expected {expected:.5}, got {blended:.5}" - ); - // Welford accumulators must still have been advanced (the kernel - // updates them BEFORE checking the Pearl-A branch). - assert!( - (result[HCS_SAMPLE_COUNT_INDEX] - 1.0).abs() < 1e-4, - "Welford sample_count must be 1 after first launch (state advances regardless of Pearl-A), got {}", - result[HCS_SAMPLE_COUNT_INDEX] - ); - // prev_target tracks target_scale for next launch's diff calc. - assert!( - (result[HCS_PREV_TARGET_INDEX] - expected).abs() < 1e-4, - "prev_target must equal target_scale after first launch, expected {expected:.5}, got {}", - result[HCS_PREV_TARGET_INDEX] - ); - } - - /// SP16 T3 — Test 4: α stays in [WELFORD_ALPHA_MIN, WELFORD_ALPHA_MAX] - /// over 50 synthetic epochs (post-cold-start). - /// - /// Run a complex signal (ramp + noise) for 50 epochs and verify - /// every recovered α at sample_count ≥ 3 lies within the Wiener - /// formula's defensive bounds. Cold-start epochs (sample_count < 3) - /// take the α=1.0 REPLACE branch by design — those are NOT subject - /// to the [0.01, 0.95] clamp because the formula isn't applied - /// (Welford state is too immature). The test skips epochs 0-1 - /// (sample_count = 1, 2) and asserts on epochs 2+ where Wiener α - /// is the active code path. Catches numerical pathologies (α hitting - /// denormal, α exceeding 0.95 from rounding, α clamped to 0 from - /// variance underflow). - #[test] - #[ignore = "requires GPU"] - fn sp16_phase3_alpha_naturally_bounded() { - let stream = make_stream(); - let kernel = load_kernel(&stream); - - // 50-epoch sequence: linear ramp 0.20 → 0.40 with sinusoidal noise. - let mut observations = vec![]; - for i in 0..50 { - let base = 0.20 + 0.20 * (i as f32 / 49.0); - let noise = 0.02 * ((i as f32 * 0.7).sin()); - observations.push((base + noise).clamp(0.0, 1.0)); - } - const SEED: f32 = 1.5; - let traj = run_synthetic_sequence(&stream, &kernel, &observations, 0.20, SEED); - - // Skip epochs 0-1 (sample_count < 3 → cold-start REPLACE → α=1.0 - // by design, NOT subject to Wiener bounds). Wiener formula - // activates from epoch 2 onward (sample_count = 3). - for (epoch, &(_, alpha)) in traj.iter().enumerate().skip(2) { - if alpha.is_nan() { - // Skip epochs where the recovered-α formula is undefined - // (target == prior). Common in steady-state intervals. - continue; - } - assert!( - alpha > 0.0, - "epoch {epoch}: α must be > 0, got {alpha:.5} (variance underflow → permanent freeze)" - ); - assert!( - alpha <= WELFORD_ALPHA_MAX + 1e-3, // small tolerance for fp rounding via blend - "epoch {epoch} (post-cold-start): α must respect WELFORD_ALPHA_MAX={}, got {alpha:.5}", - WELFORD_ALPHA_MAX - ); - } - } -} - -// ═══════════════════════════════════════════════════════════════════════════ -// SP16 T3 — host-only regression test (Test 5). -// Verifies the kernel sources do NOT contain the deleted hardcoded -// `0.05f` literal in the EMA blend math. Pure string scan — no GPU -// required. Per `feedback_no_legacy_aliases.md` and -// `feedback_isv_for_adaptive_bounds.md`. -// ═══════════════════════════════════════════════════════════════════════════ - -/// SP16 T3 — Test 5: regression-locked check that no `alpha = 0.05f` -/// or comparable hardcoded blend constant remains in the producer kernel -/// sources. Runs at compile time via `include_str!`; failure points -/// directly at the introduced regression. -#[test] -fn sp16_phase3_no_hardcoded_alpha() { - const T1_SOURCE: &str = include_str!( - "../src/cuda_pipeline/min_hold_temperature_update_kernel.cu" - ); - const T2_SOURCE: &str = include_str!( - "../src/cuda_pipeline/hold_cost_scale_update_kernel.cu" - ); - - // The blend equation pre-T3 was `(1.0f - alpha) * current + alpha * target_*` - // with `alpha = 0.05f`. Search for the canonical form. - for (name, src) in [("T1 min_hold_temperature", T1_SOURCE), - ("T2 hold_cost_scale", T2_SOURCE)] { - // The hardcoded literal must not appear in any blend assignment. - // Allow `0.05f` to appear in commentary (e.g. "the deleted α=0.05 - // constant"), but NOT as a kernel-local variable initializer or - // an arithmetic operand. We approximate by requiring the literal - // not to appear in any line that also contains an `alpha` token - // outside of comments. - for (lineno, line) in src.lines().enumerate() { - // Skip C-style comment lines (the kernel uses leading `*` for - // doc comments inside the algorithm explanation block). - let trimmed = line.trim_start(); - if trimmed.starts_with("//") || trimmed.starts_with("*") - || trimmed.starts_with("/*") { - continue; - } - // Hard fail on any executable line that contains both `alpha` - // and `0.05`. The pre-T3 implementation had exactly one such - // site per kernel (`const float alpha = 0.05f;`). - if line.contains("alpha") && line.contains("0.05") { - panic!( - "{name} kernel line {} reintroduces hardcoded alpha=0.05: {}", - lineno + 1, line - ); - } - } - // Defence in depth: the new kernel ABI passes Welford slot indices, - // not an `alpha` scalar. Verify the new param signature appears - // (catches accidental revert to the pre-T3 ABI). - assert!( - src.contains("target_mean_idx") && src.contains("sample_count_idx"), - "{name} kernel does not reference Welford accumulator slots — Wiener α implementation missing" - ); - } -} +// SP16 Phase 2 / Phase 3 oracle tests (11 GPU tests + helpers + 1 +// hardcoded-alpha regression test) DELETED by SP18 Phase 1 +// (2026-05-09). The hold_cost_scale_update_kernel cubin they +// include_bytes! is gone. See docs/sp18-wireup-audit.md (A2). // ═══════════════════════════════════════════════════════════════════════════ // SP16 Phase 0 (2026-05-08) — Q-by-action HEALTH_DIAG diagnostic. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index ab08e0542..3ba3014e9 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,110 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP18 v2 Phase 1 Tasks 1.3-1.6: atomic deletion of SP13/SP16 hold_cost_scale chain + +Atomic deletion of the entire reactive Hold-cost controller chain (SP13 P0a host +controller + SP16 P2 producer + SP16 T3 Wiener-α machinery) per DD7=c. Replaced +by the SP18 D-leg structural Hold opportunity-cost reward (Phase 2 lands the +`compute_sp18_hold_opportunity_cost` device fn at SP18-bound caps in slots +[483..493)). + +**Scope: 18 sites in one commit (8 plan + 10 audit-expanded A1-A10)** per +`feedback_no_partial_refactor`. The expansion was pre-validated by the consumer +audit script (`scripts/audit_sp18_consumers.sh`); see `docs/sp18-wireup-audit.md` +for the A1-A10 inventory. + +**INTERIM STATE — L40S DISPATCH FORBIDDEN:** The 3 reward subtraction sites in +`experience_kernels.cu` are now placeholder-commented but contain no Hold cost. +Phase 2 must land before any L40S smoke. The reward shape is currently +biased toward Hold (no per-bar penalty exists). + +**Files modified:** + +- `crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu` — DELETED. +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — 3 cost subtraction + sites (segment_complete + per-bar positioned-Hold + per-bar flat-Hold) + replaced with placeholder comments noting the Phase 2 device-fn replacement. +- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — DELETE + `HOLD_COST_SCALE_INDEX=461`, `SENTINEL_HOLD_COST_SCALE`, + `HOLD_COST_SCALE_MIN/MAX`, `SP16_P2_HOLD_COST_SCALE_SLOT_BASE/END`, + `HCS_TARGET_MEAN/M2/DIFF_MEAN/DIFF_M2/PREV_TARGET/SAMPLE_COUNT_INDEX` (slots + 462..468). Updated `SP16_T3_WIENER_SLOT_BASE` 462 → 468 (HCS block retired; + MHT block remains at [468..474)). Layout-locked tests + `sp16_p2_hold_cost_scale_slot_layout_locked` + `all_sp16_p2_slots_fit_within_isv_total_dim` + removed; `sp16_t3_wiener_welford_slot_layout_locked` updated to assert MHT + block only. +- `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs` — DELETE + `HOLD_COST_CONTROLLER_GAIN`, `HOLD_COST_FLOOR_RATIO`, `HOLD_COST_CEIL_RATIO`. + RETAIN `HOLD_COST_BASE` (constructor-init-only diagnostic at slot 380, never + updated) and `HOLD_RATE_TARGET_DEFAULT` (constructor-init at slot 381). +- `crates/ml/src/cuda_pipeline/state_layout.cuh` — DELETE + `ISV_HOLD_COST_SCALE_IDX` + 6 `ISV_HCS_*_IDX` mirror constants. Range + [461..468) becomes RESERVED gap (SP14-C.1 RESERVED-gap pattern). +- `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` — DELETE + `HoldCostScaleUpdateOps` struct + impl (~120 lines) + `HOLD_COST_SCALE_UPDATE_CUBIN` + import. +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — DELETE + `HOLD_COST_SCALE_UPDATE_CUBIN` static, `hold_cost_scale_update_ops` struct + field, constructor `let hold_cost_scale_update_ops = …` line + struct field-init + line, `launch_hold_cost_scale_update` method (~80 lines). Layout fingerprint + seed updated: `HOLD_COST_SCALE=461;HCS_TARGET_MEAN=462;…;HCS_SAMPLE_COUNT=467;` + collapsed to `RESERVED_GAP_461_TO_468=SP18_P1_RETIRED;`. ISV_TOTAL_DIM stays + at 507 — gap is documentation, not size shrink (per recommendation against + compaction). +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — DELETE 5 sites: + (a) `launch_hold_cost_scale_update` call site + warning log, + (b) hold_cost_scale_diag HEALTH_DIAG emit block (~50 lines reading slot 461 + + Welford accumulators), + (c) host SP13 P0a Hold-pricing controller block (~60 lines reading slots + 381+382, computing `cost = HOLD_COST_BASE × (1 + GAIN × excess)`, writing + slot 380), + (d) 7 dispatch arms in `reset_named_state`: `sp16_phase2_hold_cost_scale_adaptive` + + `sp16_t3_hcs_target_mean/m2/diff_mean/diff_m2/prev_target/sample_count`. +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — DELETE 7 + `RegistryEntry` instances (slot 461 + 6 HCS_*) and the `lock_sp18_v2_pp4_retired_chain` + test (`sp16_t2_t3_slots_marked_retired`). Per user call: locking deletion + is pointless ceremony, no inverse-contract replacement. +- `crates/ml/build.rs` — DELETE `hold_cost_scale_update_kernel.cu` from the + cubin manifest. +- `crates/ml/tests/sp14_oracle_tests.rs` — DELETE 11 GPU oracle tests + + helpers (`sp16_phase2_*` × 5, `sp16_phase3_*` × 5, plus `recover_alpha`, + `compute_target_scale`, `run_synthetic_sequence` helpers, plus the standalone + `sp16_phase3_no_hardcoded_alpha` source-string scan). 749 lines removed. + These `include_bytes!` the deleted cubin and cannot survive. + +**Observation chain preserved per DD7=c:** + +- Slot 380 (HOLD_COST_INDEX): constructor-init at `HOLD_COST_BASE=0.005`, + never updated post-startup. Read-only diagnostic anchor. +- Slot 381 (HOLD_RATE_TARGET_INDEX): constructor-init at + `HOLD_RATE_TARGET_DEFAULT=0.20`, never updated. +- Slot 382 (HOLD_RATE_OBSERVED_EMA_INDEX): per-step `hold_rate_observer_kernel` + + Pearls A+D EMA chain still wired (used by MIN_HOLD_TEMPERATURE producer at + slot 460 — different consumer chain). +- MHT_* slots [468..474): MIN_HOLD_TEMPERATURE Welford accumulators still + wired, dispatch arms preserved. + +**Validation:** + +- `cargo check --workspace` clean. +- `cargo test -p ml --lib state_reset_registry` — 5 passed (including + `every_fold_and_soft_reset_entry_has_dispatch_arm` contract test). +- `cargo test -p ml --lib sp14_isv_slots` — 23 passed (including + `sp18_combined_slot_layout_locked`, `sp18_shrink_perturb_slot_layout_locked`, + `sp16_t3_wiener_welford_slot_layout_locked`). +- `cargo test -p ml --lib sp18` — 4 passed. +- Audit fingerprint diff (pre vs post): `HOLD_COST_SCALE_INDEX` consumer count + 70 → 25 (only docstring archaeology + ISV_TOTAL_DIM giant comment + one + comment in sp14_isv_slots referencing the retirement); HCS_* consumer count + 108 → 17 (only registry archaeology + giant ISV_TOTAL_DIM comment + dqn- + wire-up-audit prose); per-bar Hold-cost subtraction sites count 13 → 0 + (zero hits in production code — the smoking gun for complete deletion). + +**ISV slot range [461..468) is RESERVED gap; layout fingerprint seed updated +in same commit so checkpoint compatibility is preserved (`ISV_TOTAL_DIM=507` +unchanged).** + ## 2026-05-09 — SP18 v2 ISV-adaptive shrink-and-perturb at fold transition Replaces the hardcoded `α=0.8 / σ=0.01` constants in `fused_training.rs::reset_for_fold`'s `shrink_and_perturb` call with ISV-driven adaptive bounds per `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`. The driving signal is the relative weight drift `||current − best||₂ / ||best||₂` already computed each epoch by the SP18 v2 weight-drift diagnostic kernel (commit `aab13a83f`). diff --git a/docs/sp18-wireup-audit.md b/docs/sp18-wireup-audit.md index fe3bbe2a4..ad7346248 100644 --- a/docs/sp18-wireup-audit.md +++ b/docs/sp18-wireup-audit.md @@ -67,7 +67,47 @@ Two options: ## Status -**HALTED** at end of Phase 1 Task 1.1 — pre-Task-1.2 (pre-commit hook), pre-Task-1.3-1.5 (atomic deletion), pre-Task-1.6 (close-out). The audit script `scripts/audit_sp18_consumers.sh` is committed; the snapshot below is locked for `--check` mode. +**Phase 1 Tasks 1.3-1.6 LANDED (2026-05-09)** — atomic deletion of all 18 sites +(8 plan + 10 audit-expanded A1-A10) committed in a single atomic commit per +`feedback_no_partial_refactor`. The branch is now in INTERIM STATE: 3 reward +sites (`experience_kernels.cu` segment_complete + per-bar positioned-Hold + +per-bar flat-Hold) are placeholder-commented but contain no Hold cost; +**L40S dispatch is FORBIDDEN** until Phase 2 lands the +`compute_sp18_hold_opportunity_cost` device fn replacing the deleted +subtraction. + +Per DD7c, observation chain preserved: slot 380 (HOLD_COST_INDEX, constructor- +init-only diagnostic at HOLD_COST_BASE=0.005), slot 381 +(HOLD_RATE_TARGET_INDEX, constructor-init at 0.20), slot 382 +(HOLD_RATE_OBSERVED_EMA_INDEX, FoldReset chain still wired). MHT_* slots +[468..474) preserved (MIN_HOLD_TEMPERATURE chain RETAINED). + +Sites deleted (atomic commit): +- 3 reward subtraction sites in experience_kernels.cu (placeholder comments) +- hold_cost_scale_update_kernel.cu (file deleted) +- HoldCostScaleUpdateOps in gpu_aux_trunk.rs (struct + impl, ~120 lines) +- launch_hold_cost_scale_update + cubin static + struct field + constructor + + field-init in gpu_dqn_trainer.rs (5 sites) +- HOLD_COST_SCALE_INDEX, HCS_* slots, sentinel/bounds, layout-locked tests in + sp14_isv_slots.rs +- HOLD_COST_CONTROLLER_GAIN/FLOOR/CEIL constants in sp13_isv_slots.rs +- ISV_HOLD_COST_SCALE_IDX + ISV_HCS_* mirror constants in state_layout.cuh +- build.rs cubin manifest entry +- 7 fold-reset registry entries (slot 461 + HCS_* 462..468) +- 7 dispatch arms in training_loop.rs +- Host controller block in training_loop.rs (lines ~4171-4226) +- launcher call site + hold_cost_scale_diag emit block in training_loop.rs +- 11 GPU oracle tests + helpers in sp14_oracle_tests.rs (~750 lines) +- lock_sp18_v2_pp4_retired_chain test (deleted, no inverse-contract + replacement — pointless ceremony per user call) +- Layout fingerprint seed updated (HOLD_COST_SCALE + HCS_* removed; range + [461..468) is RESERVED gap; ISV_TOTAL_DIM stays at 507 per checkpoint + compatibility — RESERVED-gap pattern from SP14-C.1) + +**Validation: cargo check --workspace clean; cargo test -p ml --lib passes +(state_reset_registry::every_fold_and_soft_reset_entry_has_dispatch_arm, +sp16_t3_wiener_welford_slot_layout_locked, sp18_combined_slot_layout_locked +all OK).** ## B-leg verification (Phase 1.1 sub-objective) @@ -107,106 +147,85 @@ re-capture is committed as part of Task 1.6 (close-out). ``` === D-leg: Slot 380 (HOLD_COST_INDEX) consumers === - total_hits=54 - crates/ml/tests/sp14_oracle_tests.rs 1 - crates/ml/build.rs 2 - crates/ml/src/cuda_pipeline/experience_kernels.cu 4 + total_hits=42 + crates/ml/build.rs 1 + crates/ml/src/cuda_pipeline/experience_kernels.cu 1 crates/ml/src/cuda_pipeline/state_layout.cuh 1 crates/ml/src/cuda_pipeline/gpu_experience_collector.rs 1 crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu 1 crates/ml/src/cuda_pipeline/sp5_isv_slots.rs 1 crates/ml/src/cuda_pipeline/sp13_isv_slots.rs 1 - crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 8 - crates/ml/src/cuda_pipeline/sp14_isv_slots.rs 3 - crates/ml/src/trainers/dqn/trainer/training_loop.rs 6 - crates/ml/src/trainers/dqn/state_reset_registry.rs 2 - docs/sp18-wireup-audit.md 4 - docs/dqn-wire-up-audit.md 6 + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 6 + crates/ml/src/trainers/dqn/trainer/training_loop.rs 3 + crates/ml/src/trainers/dqn/state_reset_registry.rs 1 + docs/sp18-wireup-audit.md 5 + docs/dqn-wire-up-audit.md 7 docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md 4 docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md 5 docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md 4 === D-leg: Slot 461 (HOLD_COST_SCALE_INDEX) consumers === - total_hits=70 - crates/ml/tests/sp14_oracle_tests.rs 20 - crates/ml/build.rs 2 - crates/ml/src/cuda_pipeline/experience_kernels.cu 6 - crates/ml/src/cuda_pipeline/state_layout.cuh 1 - crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu 1 - crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 7 - crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs 2 - crates/ml/src/cuda_pipeline/sp14_isv_slots.rs 3 - crates/ml/src/trainers/dqn/trainer/training_loop.rs 4 - crates/ml/src/trainers/dqn/state_reset_registry.rs 2 - docs/sp18-wireup-audit.md 4 - docs/dqn-wire-up-audit.md 5 + total_hits=30 + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 1 + crates/ml/src/cuda_pipeline/sp14_isv_slots.rs 2 + docs/sp18-wireup-audit.md 6 + docs/dqn-wire-up-audit.md 8 docs/superpowers/plans/2026-05-08-sp16-hold-collapse-structural-fix.md 10 docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md 3 === D-leg: Slots [462..468) (HCS_* Welford) consumers === - total_hits=108 - crates/ml/tests/sp14_oracle_tests.rs 24 - crates/ml/src/cuda_pipeline/state_layout.cuh 6 - crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu 12 - crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 12 - crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs 4 - crates/ml/src/cuda_pipeline/sp14_isv_slots.rs 12 - crates/ml/src/trainers/dqn/trainer/training_loop.rs 16 - crates/ml/src/trainers/dqn/state_reset_registry.rs 15 - docs/dqn-wire-up-audit.md 6 + total_hits=19 + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 1 + crates/ml/src/trainers/dqn/state_reset_registry.rs 9 + docs/dqn-wire-up-audit.md 8 docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md 1 === D-leg: hold_cost_scale_update_kernel references === - total_hits=103 - crates/ml/tests/sp14_oracle_tests.rs 16 - crates/ml/build.rs 2 - crates/ml/src/cuda_pipeline/experience_kernels.cu 13 - crates/ml/src/cuda_pipeline/state_layout.cuh 2 - crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu 1 - crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 10 - crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs 6 + total_hits=67 + crates/ml/tests/sp14_oracle_tests.rs 1 + crates/ml/build.rs 1 + crates/ml/src/cuda_pipeline/state_layout.cuh 1 + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 5 + crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs 1 crates/ml/src/cuda_pipeline/sp14_isv_slots.rs 3 - crates/ml/src/trainers/dqn/trainer/training_loop.rs 2 - crates/ml/src/trainers/dqn/state_reset_registry.rs 1 - docs/sp18-wireup-audit.md 12 - docs/dqn-wire-up-audit.md 10 + crates/ml/src/trainers/dqn/trainer/training_loop.rs 1 + docs/sp18-wireup-audit.md 11 + docs/dqn-wire-up-audit.md 18 docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md 4 docs/superpowers/plans/2026-05-08-sp16-hold-collapse-structural-fix.md 9 docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md 12 === D-leg: hold_rate_observer_kernel (RETAINED — diag chain stays per DD7=c) === - total_hits=208 + total_hits=162 crates/ml/tests/sp13_phase0_oracle_tests.rs 19 - crates/ml/tests/sp14_oracle_tests.rs 40 - crates/ml/build.rs 7 + crates/ml/tests/sp14_oracle_tests.rs 20 + crates/ml/build.rs 5 crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu 6 crates/ml/src/cuda_pipeline/state_layout.cuh 1 crates/ml/src/cuda_pipeline/gpu_experience_collector.rs 12 crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu 4 - crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu 2 crates/ml/src/cuda_pipeline/sp5_isv_slots.rs 3 crates/ml/src/cuda_pipeline/sp13_isv_slots.rs 2 - crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 18 - crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs 9 - crates/ml/src/cuda_pipeline/sp14_isv_slots.rs 7 - crates/ml/src/trainers/dqn/trainer/training_loop.rs 23 - crates/ml/src/trainers/dqn/state_reset_registry.rs 5 - docs/sp18-wireup-audit.md 5 + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs 10 + crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs 5 + crates/ml/src/cuda_pipeline/sp14_isv_slots.rs 5 + crates/ml/src/trainers/dqn/trainer/training_loop.rs 13 + crates/ml/src/trainers/dqn/state_reset_registry.rs 3 + docs/sp18-wireup-audit.md 7 docs/isv-slots.md 1 - docs/dqn-wire-up-audit.md 9 + docs/dqn-wire-up-audit.md 11 docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md 6 docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md 1 docs/superpowers/plans/2026-05-05-sp14-aux-q-wire-earned-gradient-flow.md 5 docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md 20 docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md 3 === D-leg: Per-bar Hold-cost subtraction sites (3 lines we change in Phase 1.5) === - total_hits=13 - crates/ml/src/cuda_pipeline/experience_kernels.cu 13 + total_hits= === D-leg: state_layout.cuh mirror constants (HOLD_COST + HCS) === - total_hits=9 - crates/ml/src/cuda_pipeline/state_layout.cuh 9 + total_hits=2 + crates/ml/src/cuda_pipeline/state_layout.cuh 2 === D-leg: state_reset_registry slot 461 entries === - total_hits=22 - crates/ml/src/trainers/dqn/state_reset_registry.rs 22 + total_hits=3 + crates/ml/src/trainers/dqn/state_reset_registry.rs 3 === D-leg: Cubin manifest references in build.rs === - total_hits=4 - crates/ml/build.rs 4 + total_hits=3 + crates/ml/build.rs 3 === B-leg: td_lambda_kernel launch sites === total_hits=9 crates/ml/src/cuda_pipeline/gpu_experience_collector.rs 8