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

feat(sp18 v2 P1.T3-T6): atomic delete SP13/SP16 hold_cost_scale chain

Browse Source 
Deletes the entire reactive Hold-cost controller chain from SP13 P0a / SP16
P2 / SP16 T3 atomically — replaced by the SP18 D-leg structural opportunity-
cost reward landing in Phase 2.

Per DD7(c): observation chain preserved (slots 380, 382). Slot 380 becomes
constructor-init-only diagnostic at HOLD_COST_BASE=0.005, never updated.

18 consumer sites deleted (8 from plan checklist + 10 from A1-A10 audit
expansion):

Production code:
  - 3 reward subtraction sites (replaced with Phase 2 placeholder comments)
  - hold_cost_scale_update_kernel.cu (file deleted)
  - HoldCostScaleUpdateOps struct in gpu_aux_trunk.rs (~120 lines)
  - launch_hold_cost_scale_update in gpu_dqn_trainer.rs (forwarder)
  - Host controller block in training_loop.rs (~60 lines)
  - Slot constants [461..468) in sp14_isv_slots.rs
  - HOLD_COST_CONTROLLER_GAIN/FLOOR/CEIL constants in sp13_isv_slots.rs
  - state_layout.cuh mirror constants for [461..468)
  - build.rs cubin manifest entry
  - 7 fold-reset registry entries + 7 dispatch arms

Tests:
  - sp14_oracle_tests.rs lines 2173-2926 (11 tests + helpers, ~750 lines)
    — these include_bytes! the deleted cubin and cannot survive
  - lock_sp18_v2_pp4_retired_chain test (deleted, no inverse-contract
    replacement — locking a deletion is pointless ceremony per user call)

Layout fingerprint: range [461..468) is RESERVED gap (SP14-C.1 pattern;
ISV_TOTAL_DIM stays at 507 for checkpoint compatibility). Layout fingerprint
seed updated atomically: HOLD_COST_SCALE + HCS_* slots replaced with
RESERVED_GAP_461_TO_468=SP18_P1_RETIRED marker.

INTERIM STATE: 3 reward sites are now missing the per-bar Hold cost. This
is forbidden to L40S-dispatch until Phase 2 lands the
compute_sp18_hold_opportunity_cost device fn that replaces the deleted
subtraction. The audit doc records this hold.

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,
sp18_shrink_perturb_slot_layout_locked, sp18_fold_reset_entries_present —
all OK). Audit fingerprint: per-bar Hold-cost subtraction sites count
13 → 0 (smoking gun for complete deletion).

Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
Audit: docs/sp18-wireup-audit.md (full A1-A10 expansion + post-deletion
       fingerprint)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 09:51:03 +02:00
parent b43413e4d3
commit 3c318953a7
13 changed files with 299 additions and 1814 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
crates/ml/build.rs
View File

@@ -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

49
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -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;
}

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

@@ -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<CudaStream>) -> Result<Self, MLError> {
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<CudaStream>,
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.

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

@@ -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::<f32>(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,

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

@@ -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 <cuda_runtime.h>
#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;
}

29
crates/ml/src/cuda_pipeline/sp13_isv_slots.rs
View File

@@ -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)`

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

@@ -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);

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

@@ -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)

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

@@ -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.
}

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

@@ -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::{

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

@@ -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<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec())
.expect("load hold_cost_scale_update_kernel cubin");
module
.load_function("hold_cost_scale_update")
.expect("load hold_cost_scale_update function")
}
/// 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<CudaStream>,
kernel: &CudaFunction,
isv_ptr: u64,
) {
unsafe {
stream
.launch_builder(kernel)
.arg(&isv_ptr)
.arg(&(HOLD_COST_SCALE_INDEX as i32))
.arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32))
.arg(&(HOLD_RATE_TARGET_INDEX as i32))
.arg(&(HCS_TARGET_MEAN_INDEX as i32))
.arg(&(HCS_TARGET_M2_INDEX as i32))
.arg(&(HCS_DIFF_MEAN_INDEX as i32))
.arg(&(HCS_DIFF_M2_INDEX as i32))
.arg(&(HCS_PREV_TARGET_INDEX as i32))
.arg(&(HCS_SAMPLE_COUNT_INDEX as i32))
.arg(&SENTINEL_HOLD_COST_SCALE)
.arg(&SENTINEL_HOLD_RATE_OBSERVED)
.arg(&HOLD_COST_SCALE_MIN)
.arg(&HOLD_COST_SCALE_MAX)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch hold_cost_scale_update");
}
stream.synchronize().expect("sync after hold_cost_scale_update");
}
/// 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<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec())
.expect("load hold_cost_scale_update_kernel cubin");
module
.load_function("hold_cost_scale_update")
.expect("load hold_cost_scale_update function")
}
fn launch_canonical(stream: &Arc<CudaStream>, kernel: &CudaFunction, isv_ptr: u64) {
unsafe {
stream
.launch_builder(kernel)
.arg(&isv_ptr)
.arg(&(HOLD_COST_SCALE_INDEX as i32))
.arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32))
.arg(&(HOLD_RATE_TARGET_INDEX as i32))
.arg(&(HCS_TARGET_MEAN_INDEX as i32))
.arg(&(HCS_TARGET_M2_INDEX as i32))
.arg(&(HCS_DIFF_MEAN_INDEX as i32))
.arg(&(HCS_DIFF_M2_INDEX as i32))
.arg(&(HCS_PREV_TARGET_INDEX as i32))
.arg(&(HCS_SAMPLE_COUNT_INDEX as i32))
.arg(&SENTINEL_HOLD_COST_SCALE)
.arg(&SENTINEL_HOLD_RATE_OBSERVED)
.arg(&HOLD_COST_SCALE_MIN)
.arg(&HOLD_COST_SCALE_MAX)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch hold_cost_scale_update");
}
stream.synchronize().expect("sync after hold_cost_scale_update");
}
/// Recover the actual α used by the kernel on the latest launch. The
/// kernel doesn't directly expose α, but it can be reverse-engineered
/// from the blend equation:
///
/// blended = (1 − α) × prior_scale + α × target_scale (when prior > 0)
/// ⇒ α = (blended − prior_scale) / (target_scale − prior_scale)
///
/// Pre-launch fix: write `prior_scale` to slot 461 with a non-sentinel
/// value so the Pearl-A bootstrap branch isn't taken; capture
/// (prior_scale, target_scale) and compare to (blended) post-launch.
/// Caller MUST set `target_scale ≠ prior_scale` to keep denominator
/// well-defined; bilateral clamp post-blend may corrupt the
/// recovered α near boundaries — choose seed/observed values that
/// keep the blend well-inside [SCALE_MIN, SCALE_MAX].
fn recover_alpha(blended: f32, prior_scale: f32, target_scale: f32) -> f32 {
(blended - prior_scale) / (target_scale - prior_scale)
}
/// Compute target_scale from observed/target without launching the
/// kernel — mirrors the kernel's overrun-norm formula exactly.
fn compute_target_scale(observed: f32, target: f32) -> f32 {
let overrun = (observed - target).max(0.0);
let denom = target.max(0.01);
let overrun_norm = (overrun / denom).clamp(0.0, 1.0);
(HOLD_COST_SCALE_MIN
+ (HOLD_COST_SCALE_MAX - HOLD_COST_SCALE_MIN) * overrun_norm)
.clamp(HOLD_COST_SCALE_MIN, HOLD_COST_SCALE_MAX)
}
/// Run a synthetic per-epoch sequence by setting (observed, prev_blended)
/// per step and capturing the post-launch slot 461.
/// Returns the trajectory of `(blended, recovered_alpha)` per epoch.
fn run_synthetic_sequence(
stream: &Arc<CudaStream>,
kernel: &CudaFunction,
observations: &[f32],
target: f32,
seed_blended: f32,
) -> Vec<(f32, f32)> {
const ISV_DIM: usize = 1024;
let mut isv = vec![0.0_f32; ISV_DIM];
isv[HOLD_RATE_TARGET_INDEX] = target;
isv[HOLD_COST_SCALE_INDEX] = seed_blended;
let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv");
isv_buf.write_from_slice(&isv);
let mut traj = Vec::with_capacity(observations.len());
for (epoch, &obs) in observations.iter().enumerate() {
let mut current = isv_buf.read_all();
current[HOLD_RATE_OBSERVED_EMA_INDEX] = obs;
isv_buf.write_from_slice(&current);
let prior_scale = current[HOLD_COST_SCALE_INDEX];
let target_scale = compute_target_scale(obs, target);
launch_canonical(stream, kernel, isv_buf.dev_ptr);
let post = isv_buf.read_all();
let blended = post[HOLD_COST_SCALE_INDEX];
// Recover α from blend equation. For the first epoch (where
// prior is the seed), α=1.0 is the cold-start path. Once
// sample_count >= 3, α derives from Welford state.
let recovered_alpha = if (target_scale - prior_scale).abs() < 1e-6 {
f32::NAN // formula undefined when prior == target
} else {
recover_alpha(blended, prior_scale, target_scale)
};
traj.push((blended, recovered_alpha));
// Sanity check: Welford counter increments monotonically.
assert!(
(post[HCS_SAMPLE_COUNT_INDEX] - (epoch as f32 + 1.0)).abs() < 1e-4,
"epoch {epoch}: sample_count expected {}, got {}",
epoch + 1, post[HCS_SAMPLE_COUNT_INDEX]
);
}
traj
}
/// SP16 T3 — Test 1: α HIGH during signal jump (cold-start
/// responsiveness).
///
/// Feed a sharp signal jump from below-target to over-target across
/// 5 epochs. The first three epochs trigger cold-start path
/// (sample_count < 3 → α=1.0 REPLACE). After that, the Wiener
/// formula sees high diff_var (the recent jumps) relative to
/// sample_var → α should remain HIGH (> 0.5) on the 4th and 5th
/// epochs — concrete proof of the responsiveness benefit over
/// the deleted hardcoded α=0.05.
#[test]
#[ignore = "requires GPU"]
fn sp16_phase3_alpha_high_during_signal_jump() {
let stream = make_stream();
let kernel = load_kernel(&stream);
// Sharp jumps: 0.20→0.30→0.40→0.30→0.40 (target=0.20).
// target_scale series (overrun_norm × 24 + 1):
// ep0: obs=0.20 → norm=0.0 → target=1
// ep1: obs=0.30 → norm=0.5 → target=13
// ep2: obs=0.40 → norm=1.0 → target=25
// ep3: obs=0.30 → norm=0.5 → target=13
// ep4: obs=0.40 → norm=1.0 → target=25
// Diff sequence: +12, +12, -12, +12 → high diff_var.
let observations = [0.20, 0.30, 0.40, 0.30, 0.40];
// Seed blended at 1.0 (NOT sentinel — bypass Pearl-A path) so we
// can recover α from the blend equation directly.
const SEED: f32 = 1.0001; // slightly above EPS_F to avoid Pearl-A REPLACE
let traj = run_synthetic_sequence(&stream, &kernel, &observations, 0.20, SEED);
// Epoch 0 (sample_count=1 < 3): cold-start α=1.0, target=1.
// Epoch 1 (sample_count=2 < 3): cold-start α=1.0, target=13.
// Epoch 2 (sample_count=3): Wiener formula activates. After
// 3 samples (1, 13, 25), sample_mean ≈ 13, sample_var huge,
// diff_var also huge. α should be substantial (>0.3).
// Epoch 3-4: Wiener formula. With high target swings, diff_var
// stays comparable to sample_var → α stays high.
let alpha_ep3 = traj[3].1;
let alpha_ep4 = traj[4].1;
assert!(
alpha_ep3 > 0.3 && !alpha_ep3.is_nan(),
"Wiener α should be HIGH (>0.3) during high-volatility signal at epoch 3, got {alpha_ep3:.5}"
);
assert!(
alpha_ep4 > 0.3 && !alpha_ep4.is_nan(),
"Wiener α should remain HIGH (>0.3) during continued high-volatility at epoch 4, got {alpha_ep4:.5}"
);
// Defensive bound check.
assert!(
alpha_ep3 <= WELFORD_ALPHA_MAX && alpha_ep4 <= WELFORD_ALPHA_MAX,
"Wiener α must respect WELFORD_ALPHA_MAX={}; got ep3={alpha_ep3:.5} ep4={alpha_ep4:.5}",
WELFORD_ALPHA_MAX
);
}
/// SP16 T3 — Test 2: α 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<f32> = traj[50..60].iter()
.map(|&(_, a)| a)
.filter(|a| !a.is_nan())
.collect();
let mean_tail_alpha = tail.iter().sum::<f32>() / tail.len() as f32;
// Tail α should 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.

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

@@ -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`).

149
docs/sp18-wireup-audit.md
View File

@@ -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).
<!-- BEGIN audit_sp18_consumers.sh FINGERPRINT -->
```
=== 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