From 1a3bcf97b89a2b4b68ba9bdde7ad7a9ad40a40d8 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 8 May 2026 15:56:46 +0200 Subject: [PATCH] feat(sp16-p2): adaptive Hold cost scale via ISV[461] MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Per train-multi-seed-pfh9n post-mortem: observed_hold_rate climbed 0.25 → 0.52 across training while cost penalty (~0.006) was 100× smaller than per-bar reward magnitudes (popart=0.97, cf=0.65). Hold action was effectively free, allowing Q(Hold) to dominate via structural low-variance bias. Fix: scale Hold cost adaptively. ISV[HOLD_COST_SCALE_INDEX=461] tracks: scale = clamp(1.0 + 24.0 × max(0, observed - target) / max(target, 0.01), 1.0, 25.0) Effective cost at 100% overrun (observed=2× target): 0.006 × 25 = 0.15, competitive with per-bar reward magnitudes (~0.01-0.1). At/below target: scale = 1.0 (no extra penalty). Pearl-A bootstrap + Welford slow EMA (α=0.05). Mirrors T1's MIN_HOLD_TEMPERATURE pattern (same input signals: ISV[382] observed, ISV[381] target). Producer: hold_cost_scale_update_kernel.cu — single-thread cold-path, per-epoch boundary, AFTER MIN_HOLD_TEMPERATURE in training_loop.rs. Consumer migration (atomic per feedback_no_partial_refactor): 3 sites in experience_kernels.cu — segment_complete branch (line ~3089), per-bar positioned-Hold branch (line ~3553), per-bar flat-Hold branch (line ~3617). Cold-start fallback: scale=1.0 when slot ≤ 0 or out-of-bounds (bit-identical pre-Phase-2 cost magnitude). ISV_TOTAL_DIM: 461 → 462. Behavioral tests (5/5 PASS on RTX 3050): - sp16_phase2_hold_cost_scale_climbs_with_overrun - sp16_phase2_hold_cost_scale_at_target_is_one - sp16_phase2_hold_cost_scale_under_target_is_one - sp16_phase2_hold_cost_scale_bounds_clamp - sp16_phase2_hold_cost_scale_pearl_a_bootstrap Regression: SP14 oracle suite 30/30 PASS, SP15 phase 1 oracle suite 36/36 PASS. Instrumentation: HEALTH_DIAG[N]: hold_cost_scale_diag obs/tgt/norm/scale. Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 12 + .../src/cuda_pipeline/experience_kernels.cu | 46 ++- crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs | 106 ++++++- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 120 +++++++- .../hold_cost_scale_update_kernel.cu | 189 ++++++++++++ crates/ml/src/cuda_pipeline/sp14_isv_slots.rs | 111 +++++++ crates/ml/src/cuda_pipeline/state_layout.cuh | 12 + .../src/trainers/dqn/state_reset_registry.rs | 5 + .../src/trainers/dqn/trainer/training_loop.rs | 72 +++++ crates/ml/tests/sp14_oracle_tests.rs | 285 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 77 +++++ 11 files changed, 1027 insertions(+), 8 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 82c777436..a756fdc96 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -867,6 +867,18 @@ 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", // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer. // Single-block 256-thread kernel computing RMS = sqrt(mean(x²)) // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 89f3294af..e0229a0ba 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -3084,9 +3084,22 @@ extern "C" __global__ void experience_env_step( * with segment_complete (Hold keeps current position, doesn't close * trades) but trail-fire can force-exit while the policy picks Hold; * this branch covers that edge. Per-bar (non-segment_complete) - * Hold-cost subtraction is in the per-bar branches below. */ + * 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) { - base_reward -= isv_signals_ptr[ISV_HOLD_COST_IDX]; + 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; } /* SP12 v3 fix (2026-05-04): asymmetric bounded cap (loss aversion). * SP11 (commit 35db31089) made the cap symmetric ±10 to fix slot-63 @@ -3548,9 +3561,21 @@ extern "C" __global__ void experience_env_step( * structurally (HOLD_COST_BASE × CEIL_RATIO = 0.005 max per bar), * well below the SP12 asymmetric cap [-10, +5] — no per-bar clamp * required. Spec §"Change 1: Price Hold (replaces v2's Eliminate - * Hold)". */ + * 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) { - r_micro -= isv_signals_ptr[ISV_HOLD_COST_IDX]; + 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; } reward = r_micro; micro_reward_per_sample[out_off] = r_micro; @@ -3612,9 +3637,18 @@ extern "C" __global__ void experience_env_step( * "stay flat") pays the same per-bar cost as a policy picking * Hold while positioned (i.e. "stay long/short"). The cost is * structurally bounded by HOLD_COST_BASE × CEIL_RATIO = 0.005 - * per bar, well below the SP12 asymmetric cap [-10, +5]. */ + * 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) { - r_opp_cost -= isv_signals_ptr[ISV_HOLD_COST_IDX]; + 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; } reward = r_opp_cost; reward_components_per_sample[out_off * 6 + 4] = r_opp_cost; diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index b8a8d85e4..8ca64ecaf 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -52,7 +52,8 @@ 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, H_S2_AUX_RMS_EMA_CUBIN, + AVG_WIN_HOLD_TIME_UPDATE_CUBIN, DD_SATURATION_FLOOR_UPDATE_CUBIN, + HOLD_COST_SCALE_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, MIN_HOLD_TEMPERATURE_UPDATE_CUBIN, REWARD_CAP_UPDATE_CUBIN, }; @@ -1182,3 +1183,106 @@ impl HS2AuxRmsEmaOps { Ok(()) } } + +/// SP16 Phase 2 (2026-05-08): adaptive Hold cost scale producer ops. +/// +/// Mirrors `MinHoldTemperatureUpdateOps` (same input chain — hold-rate +/// overrun — same per-epoch cadence — same Pearl-A bootstrap pattern). +/// The kernel is a single-block single-thread cold-path producer that +/// reads `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and +/// `ISV[HOLD_RATE_TARGET_INDEX=381]`, maps overrun to a [1, 25] +/// multiplier, and Pearl-A-bootstrap + α=0.05 EMA blends into +/// `ISV[HOLD_COST_SCALE_INDEX=461]`. +/// +/// # Why this struct exists separately from MinHoldTemperatureUpdateOps +/// +/// Both producers read the SAME inputs (slots 382 + 381) and run at the +/// SAME cadence (per-epoch boundary), but write to DIFFERENT outputs +/// (slot 460 — temperature in [5, 50] vs slot 461 — scale in [1, 25]). +/// Combining them into a single fused kernel was rejected because: +/// 1. The two outputs are downstream of two semantically-distinct +/// consumer paths (min-hold soft saturation vs Hold cost +/// magnitude); fusing them couples kernels that should be +/// independently auditable. +/// 2. Per-epoch launch overhead is negligible (single-thread cold +/// kernel), so the fusion buys no measurable perf. +/// 3. Mirrors the SP14-P0-A and SP14-P1-Producer split — ONE producer +/// kernel per ISV slot, ONE wiring point per consumer. +/// +/// # Pearls applied +/// - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. +/// - `pearl_first_observation_bootstrap.md` — sentinel 0.0 → REPLACE. +/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` +/// pre-loaded at construction; on-device guards. +/// - `pearl_symmetric_clamp_audit.md` — bilateral +/// `fmaxf(lo, fminf(x, hi))` clamp on target_scale + post-blend. +/// - `feedback_isv_for_adaptive_bounds.md` — bounds [1, 25] are +/// Category-1 dimensional safety floors (NOT tuning). +/// - `feedback_no_partial_refactor.md` — kernel + launcher + 3 consumer +/// sites + tests + audit doc updated atomically in the same commit. +#[allow(missing_debug_implementations)] +pub(crate) struct HoldCostScaleUpdateOps { + update_kernel: CudaFunction, +} + +impl HoldCostScaleUpdateOps { + pub(crate) fn new(stream: &Arc) -> Result { + let context = stream.context(); + let module = context + .load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("hold_cost_scale_update cubin load: {e}")))?; + let update_kernel = module + .load_function("hold_cost_scale_update") + .map_err(|e| MLError::ModelError(format!("hold_cost_scale_update load: {e}")))?; + Ok(Self { update_kernel }) + } + + /// Launch the adaptive Hold cost scale producer. + /// + /// 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). + /// - `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). + /// - `alpha`: HOLD_COST_SCALE_EMA_ALPHA (0.05). + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch( + &self, + stream: &Arc, + isv_ptr: u64, + scale_idx: i32, + hold_rate_observed_idx: i32, + hold_rate_target_idx: i32, + sentinel_scale: f32, + sentinel_observed_hold_rate: f32, + scale_min: f32, + scale_max: f32, + alpha: 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(&sentinel_scale) + .arg(&sentinel_observed_hold_rate) + .arg(&scale_min) + .arg(&scale_max) + .arg(&alpha) + .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(()) + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 84016b932..0dd30ed50 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2165,6 +2165,33 @@ 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")); + /// 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 /// `h_s2_aux [B, SH2]` (aux trunk final output) and EMA-blending into @@ -2493,7 +2520,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -pub(crate) const ISV_TOTAL_DIM: usize = 461; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon + Class A P0-A adaptive reward caps + Class A P1-Producer adaptive Bayesian Kelly priors + Class A audit-fix Batch 4-A adaptive DD saturation floor + Class A audit-fix Batch 4-B adaptive plan_threshold floor + adaptive MIN_HOLD_TEMPERATURE. Bumped 459 → 461 by Class A audit-fix Batch 4-B (2026-05-08): added 2 slots [459..461) — PLAN_THRESHOLD_FLOOR_ADAPTIVE (459) + MIN_HOLD_TEMPERATURE_ADAPTIVE (460). PLAN_THRESHOLD_FLOOR_ADAPTIVE replaces the hardcoded `0.1f` floor in `plan_threshold_update_kernel.cu:44` (`fmaxf(0.1f, 0.5f * ema)`) — Design Y inline producer (no new kernel; same kernel writes both slot 49 and the slow-EMA-shadow floor at slot 459 in the same launch, then uses the freshly-written floor as the lower bound). Pearl-A bootstrap (sentinel 0.1 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence — the floor is a long-term shadow of `0.5 × readiness_ema` and must move much slower than the readiness EMA itself). Bounds [0.05, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (probability units). MIN_HOLD_TEMPERATURE_ADAPTIVE replaces the deleted epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch). Producer kernel `min_hold_temperature_update_kernel` reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary directional accuracy), maps to a [5, 50] temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × confusion` where `confusion = 1 − clamp((short_ema − 0.5)/0.5, 0, 1)` (so dir_acc ≈ 0.5 plateau → temp HIGH/permissive exit ramp; dir_acc → 1.0 → temp LOW/strict informed-commitment), and slow-EMA-blends into slot 460. HIGH T = forgiving (`deficit/(deficit+T)` saturation factor → 0); LOW T = sharp (→ 1) per `compute_min_hold_penalty` at trade_physics.cuh:575. Pearl-A bootstrap (sentinel 50.0 matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start under both regimes) + α=0.05 mid-cadence EMA. Bounds [5, 50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (matches the original schedule range). Decouples temperature from epoch number — when WR plateaus, the old schedule pinned LOW temp (sharp) at end of training (max punishment exactly when most needed forgiveness); the signal-driven version reacts to actual realized skill. Per `feedback_no_legacy_aliases.md` the legacy schedule constants + helper function are deleted atomically. Bumped 458 → 459 by Class A audit-fix Batch 4-A (2026-05-08): added 1 adaptive-DD-saturation-floor slot [458..459) — DD_SATURATION_FLOOR_ADAPTIVE (458). Replaces the hardcoded `0.25f` saturation floor in `trade_physics.cuh::apply_margin_cap` (the upper end of the linear position-size scaling ramp dd_scale = max(0.05, 1.0 − dd_frac / saturation_floor)). Distinct semantic role from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound). Producer kernel `dd_saturation_floor_update_kernel` aggregates per-env DD_MAX from the existing `dd_state_per_env[env*6 + 1]` tile (same source as the SP15 dd_state_kernel writes to), computes p75 estimator via Welford rule (mean + Z_75 × sigma) × safety_factor=1.5, and slow-EMA blends. Pearl-A bootstrap (sentinel 0.25 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.01 slow EMA (DD distribution is a slow-moving fold-volatility property, shouldn't track per-epoch noise). Bounds [0.10, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Atomically deleted alongside this slot addition: legacy `compute_drawdown_penalty` device function in trade_physics.cuh + its single caller at experience_kernels.cu:3822 + the `w_dd` and `dd_threshold` kernel args (Item 2 Case A — SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting per `feedback_no_partial_refactor`). Bumped 454 → 458 by Class A P1-Producer (2026-05-08): added 4 adaptive-Bayesian-prior slots [454..458) — KELLY_PRIOR_WINS (454) + KELLY_PRIOR_LOSSES (455) + KELLY_PRIOR_SUM_WINS (456) + KELLY_PRIOR_SUM_LOSSES (457). Per Class A audit, hardcoded `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f / prior_sum_losses=0.01f` (kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::kelly_position_cap:304-307) was a static "I don't know" Bayesian prior frozen at sprint-1 defaults — the cold-start of every new fold ignored what prior folds had learned about the realized Kelly distribution. Producer kernel `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats from the existing `portfolio_state[n_envs, PS_STRIDE]` buffer (same source as kelly_cap_update consumes) and slow-EMA blends into the four slots. Pearl-A bootstrap (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded values for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds). Bounds counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Bumped 452 → 454 by Class A P0-A (2026-05-08): added 2 adaptive-reward-cap slots [452..454) — REWARD_POS_CAP_ADAPTIVE (452) + REWARD_NEG_CAP_ADAPTIVE (453). Per Class A audit, hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha — controller signals (sharpe EMA, var_q, q_gap) all derive from the CAPPED buffer, so the controller cannot select for trades it cannot see. The state_layout comment explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need"; 11 superprojects of WR-stuck-at-46-48% revealed adaptive need. Producer kernel `reward_cap_update_kernel` computes p99(winning_returns) × safety_factor=1.5 → POS cap and NEG = -2 × POS (preserves Kahneman 2:1 loss-aversion asymmetry per `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from hardcoded scalar to producer-time multiplier, single source of truth). Pearl-A bootstrap (sentinels 5.0/-10.0) + Wiener-α=0.01 slow EMA. Bounds POS∈[1, 50], NEG∈[-100, -2] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) +pub(crate) const ISV_TOTAL_DIM: usize = 462; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon + Class A P0-A adaptive reward caps + Class A P1-Producer adaptive Bayesian Kelly priors + Class A audit-fix Batch 4-A adaptive DD saturation floor + Class A audit-fix Batch 4-B adaptive plan_threshold floor + adaptive MIN_HOLD_TEMPERATURE + SP16 Phase 2 adaptive Hold cost scale. Bumped 461 → 462 by SP16 Phase 2 (2026-05-08): added 1 slot [461..462) — HOLD_COST_SCALE_INDEX (461). Replaces the static `ISV[HOLD_COST_INDEX=380]` per-bar Hold-cost magnitude (~0.006) when over-holding is observed. Producer kernel `hold_cost_scale_update_kernel` reads `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and `ISV[HOLD_RATE_TARGET_INDEX=381]` (same SP13 hold-pricing chain feeding SP16-P1's MIN_HOLD_TEMPERATURE — same cadence, same input pattern), maps overrun into a [1, 25] multiplier via `target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × clamp(max(0, observed - target) / max(target, 0.01), 0, 1)`, and Pearl-A-bootstrap + α=0.05 mid-cadence EMA blends into slot 461. Consumer: `experience_kernels.cu` per-bar Hold cost subtractions (3 sites: segment_complete branch line 3089, per-bar positioned-Hold branch line 3553, per-bar flat-Hold branch line 3617) multiply `isv[ISV_HOLD_COST_IDX=380]` by `isv[HOLD_COST_SCALE_INDEX=461]`. Cold-start fallback: when slot 461 is at sentinel 0.0 OR outside [1, 25], the consumer uses scale=1.0 (bit-identical to pre-Phase-2 v3-P0a Hold-cost magnitude). At observed=2× target → scale=25× → effective cost ~0.15/bar competitive with per-bar reward magnitudes ~0.01-0.1 but well below SP12 asymmetric cap [-10, +5]. Bounds [1, 25] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md`. Bumped 459 → 461 by Class A audit-fix Batch 4-B (2026-05-08): added 2 slots [459..461) — PLAN_THRESHOLD_FLOOR_ADAPTIVE (459) + MIN_HOLD_TEMPERATURE_ADAPTIVE (460). PLAN_THRESHOLD_FLOOR_ADAPTIVE replaces the hardcoded `0.1f` floor in `plan_threshold_update_kernel.cu:44` (`fmaxf(0.1f, 0.5f * ema)`) — Design Y inline producer (no new kernel; same kernel writes both slot 49 and the slow-EMA-shadow floor at slot 459 in the same launch, then uses the freshly-written floor as the lower bound). Pearl-A bootstrap (sentinel 0.1 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence — the floor is a long-term shadow of `0.5 × readiness_ema` and must move much slower than the readiness EMA itself). Bounds [0.05, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (probability units). MIN_HOLD_TEMPERATURE_ADAPTIVE replaces the deleted epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch). Producer kernel `min_hold_temperature_update_kernel` reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary directional accuracy), maps to a [5, 50] temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × confusion` where `confusion = 1 − clamp((short_ema − 0.5)/0.5, 0, 1)` (so dir_acc ≈ 0.5 plateau → temp HIGH/permissive exit ramp; dir_acc → 1.0 → temp LOW/strict informed-commitment), and slow-EMA-blends into slot 460. HIGH T = forgiving (`deficit/(deficit+T)` saturation factor → 0); LOW T = sharp (→ 1) per `compute_min_hold_penalty` at trade_physics.cuh:575. Pearl-A bootstrap (sentinel 50.0 matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start under both regimes) + α=0.05 mid-cadence EMA. Bounds [5, 50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (matches the original schedule range). Decouples temperature from epoch number — when WR plateaus, the old schedule pinned LOW temp (sharp) at end of training (max punishment exactly when most needed forgiveness); the signal-driven version reacts to actual realized skill. Per `feedback_no_legacy_aliases.md` the legacy schedule constants + helper function are deleted atomically. Bumped 458 → 459 by Class A audit-fix Batch 4-A (2026-05-08): added 1 adaptive-DD-saturation-floor slot [458..459) — DD_SATURATION_FLOOR_ADAPTIVE (458). Replaces the hardcoded `0.25f` saturation floor in `trade_physics.cuh::apply_margin_cap` (the upper end of the linear position-size scaling ramp dd_scale = max(0.05, 1.0 − dd_frac / saturation_floor)). Distinct semantic role from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound). Producer kernel `dd_saturation_floor_update_kernel` aggregates per-env DD_MAX from the existing `dd_state_per_env[env*6 + 1]` tile (same source as the SP15 dd_state_kernel writes to), computes p75 estimator via Welford rule (mean + Z_75 × sigma) × safety_factor=1.5, and slow-EMA blends. Pearl-A bootstrap (sentinel 0.25 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.01 slow EMA (DD distribution is a slow-moving fold-volatility property, shouldn't track per-epoch noise). Bounds [0.10, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Atomically deleted alongside this slot addition: legacy `compute_drawdown_penalty` device function in trade_physics.cuh + its single caller at experience_kernels.cu:3822 + the `w_dd` and `dd_threshold` kernel args (Item 2 Case A — SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting per `feedback_no_partial_refactor`). Bumped 454 → 458 by Class A P1-Producer (2026-05-08): added 4 adaptive-Bayesian-prior slots [454..458) — KELLY_PRIOR_WINS (454) + KELLY_PRIOR_LOSSES (455) + KELLY_PRIOR_SUM_WINS (456) + KELLY_PRIOR_SUM_LOSSES (457). Per Class A audit, hardcoded `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f / prior_sum_losses=0.01f` (kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::kelly_position_cap:304-307) was a static "I don't know" Bayesian prior frozen at sprint-1 defaults — the cold-start of every new fold ignored what prior folds had learned about the realized Kelly distribution. Producer kernel `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats from the existing `portfolio_state[n_envs, PS_STRIDE]` buffer (same source as kelly_cap_update consumes) and slow-EMA blends into the four slots. Pearl-A bootstrap (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded values for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds). Bounds counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Bumped 452 → 454 by Class A P0-A (2026-05-08): added 2 adaptive-reward-cap slots [452..454) — REWARD_POS_CAP_ADAPTIVE (452) + REWARD_NEG_CAP_ADAPTIVE (453). Per Class A audit, hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha — controller signals (sharpe EMA, var_q, q_gap) all derive from the CAPPED buffer, so the controller cannot select for trades it cannot see. The state_layout comment explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need"; 11 superprojects of WR-stuck-at-46-48% revealed adaptive need. Producer kernel `reward_cap_update_kernel` computes p99(winning_returns) × safety_factor=1.5 → POS cap and NEG = -2 × POS (preserves Kahneman 2:1 loss-aversion asymmetry per `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from hardcoded scalar to producer-time multiplier, single source of truth). Pearl-A bootstrap (sentinels 5.0/-10.0) + Wiener-α=0.01 slow EMA. Bounds POS∈[1, 50], NEG∈[-100, -2] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -7635,6 +7662,13 @@ 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, + // ── Speculative inference cache ── /// Cached trunk output [B, SH2] from between-bar speculative forward. speculative_h_s2: CudaSlice, @@ -15061,6 +15095,79 @@ 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_EMA_ALPHA, HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX, + HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED, + }; + 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" + ); + + 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, + SENTINEL_HOLD_COST_SCALE, + SENTINEL_HOLD_RATE_OBSERVED, + HOLD_COST_SCALE_MIN, + HOLD_COST_SCALE_MAX, + HOLD_COST_SCALE_EMA_ALPHA, + )?; + + Ok(()) + } + /// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary /// producer. /// @@ -23502,6 +23609,15 @@ 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)?; + // ── Speculative inference cache ── let speculative_h_s2 = stream.alloc_zeros::(b * sh2) .map_err(|e| MLError::ModelError(format!("speculative_h_s2 alloc: {e}")))?; @@ -24342,6 +24458,8 @@ 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, speculative_h_s2, speculative_features, speculative_valid: false, diff --git a/crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu b/crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu new file mode 100644 index 000000000..3c20797bf --- /dev/null +++ b/crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu @@ -0,0 +1,189 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP16 Phase 2 (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 + Welford + * EMA, but maps the overrun into the Hold-cost-scale dimension instead + * of the min-hold-temperature dimension. + * + * Why this kernel exists (post-mortem of train-multi-seed-pfh9n): + * + * HEALTH_DIAG[2]: hold_pricing observed_rate=0.2501 target=0.2000 + * cost=0.006251 + * + * The realised Hold rate climbed 0.25 → 0.52 across training while the + * cost penalty stayed at ~0.006 — ~100× smaller than per-bar reward + * magnitudes (popart=0.97, cf=0.65 in `reward_split`). Hold became + * effectively free, allowing the structural Q-bias toward zero-variance + * Hold to dominate. This kernel addresses the cost-magnitude axis: when + * the model is over-holding (overrun > 0), the per-bar Hold cost is + * scaled up multiplicatively at the 3 consumer sites in + * `experience_kernels.cu` (segment_complete branch line ~3089, per-bar + * positioned-Hold branch line ~3553, per-bar flat-Hold branch line ~3617). + * + * Mapping: + * + * overrun = max(0, observed − target) + * overrun_normalized = clamp(overrun / max(target, 0.01), 0, 1) + * target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_normalized + * + * - observed = 2× target → overrun_norm = 1.0 → scale = 25 + * (effective cost ~0.006 × 25 = 0.15/bar — competitive with per-bar + * reward magnitudes ~0.01-0.1 but well below SP12 cap [-10, +5]) + * - observed = target → overrun_norm = 0 → scale = 1 + * (no extra penalty — bit-identical to pre-Phase-2 cost magnitude) + * - observed < target → overrun_norm = 0 → scale = 1 + * (under-holding — no penalty escalation needed) + * + * Why bounds [1, 25] are dimensional safety, not tuning: + * - SCALE_MIN=1: the floor is "no extra penalty" — lower bound enforced + * by the formula (overrun is clamped at 0 from below). Below 1 would + * mean "discount the cost below baseline" which is the opposite of + * what the controller is supposed to do under overrun. + * - SCALE_MAX=25: at this ceiling, the effective Hold cost is ~0.15/bar. + * Above 25× the static base, single Hold bars start dominating per- + * bar reward magnitudes by an order of magnitude — overcorrection + * territory. The audit-doc 25× ceiling matches the ratio found in + * the pfh9n post-mortem (cost ~0.006, reward ~0.97 → ratio ~150 → + * tightening to 6:1 via 25× scale is the conservative point). + * + * Algorithm (single-block single-thread — cold path, per-epoch boundary): + * + * Phase 1 (single thread): read observed_hold_rate (slot 382), + * target_hold_rate (slot 381), and current scale (slot 461) from ISV. + * + * Phase 2 (sentinel guard): if observed_hold_rate is at sentinel 0.0 + * (no per-step Hold observations have landed yet, e.g. epoch 0 of + * any fold before the first hold_rate_observer launch chains in), + * keep ISV slot 461 unchanged — the consumer falls back to scale=1.0 + * when slot 461 is at sentinel 0.0 (bit-identical pre-Phase-2 cost + * magnitude). This is the bootstrap-protection equivalent of the + * SP16-P1 dir_acc sentinel guard. + * + * Phase 3 (compute target scale): + * + * overrun = fmaxf(0, observed − target) + * overrun_norm = fminf(1, overrun / fmaxf(target, 0.01)) + * target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × overrun_norm + * bilateral clamp to [SCALE_MIN, SCALE_MAX] + * + * Phase 4 (Pearl-A bootstrap + Welford-style EMA): + * + * if current is at sentinel 0.0 (within EPS) → REPLACE with target_scale. + * else → blended = (1 − α) × current + α × target_scale + * re-clamp post-blend. + * + * Pearls + invariants: + * + * - `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 for sentinel + * detection. + * + * - `pearl_first_observation_bootstrap.md` — first valid hold-rate + * observation replaces sentinel 0.0 directly; no blend. + * + * - `feedback_no_stubs.md` — full body, no return-zero placeholders. + * + * - `feedback_isv_for_adaptive_bounds.md` — bounds [1, 25] are + * dimensional safety floors (not tuning constants). + * + * - `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 in the + * same commit. + * + * Args: + * isv — `[ISV_TOTAL_DIM]` device f32, modified + * in place at index `scale_idx`. + * 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). + * sentinel_scale — SENTINEL_HOLD_COST_SCALE (0.0). + * sentinel_observed_hold_rate — SENTINEL_HOLD_RATE_OBSERVED (0.0; matches + * the SP13 hold-rate observer cold-start — + * shared sentinel by design with SP16-P1). + * scale_min — HOLD_COST_SCALE_MIN (1.0). + * scale_max — HOLD_COST_SCALE_MAX (25.0). + * alpha — HOLD_COST_SCALE_EMA_ALPHA (0.05). + * + * Launch: grid=(1, 1, 1), block=(1, 1, 1). + * No shared memory — single-thread cold-path kernel. + * ══════════════════════════════════════════════════════════════════════════ */ + +#include + +#define EPS_F 1e-6f + +extern "C" __global__ +void hold_cost_scale_update( + float* __restrict__ isv, + int scale_idx, + int hold_rate_observed_idx, + int hold_rate_target_idx, + float sentinel_scale, + float sentinel_observed_hold_rate, + float scale_min, + float scale_max, + float alpha) +{ + /* 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 (matches the SP13 hold_rate_observer fold-reset sentinel — see + * `state_reset_registry.rs::sp13_hold_rate_observed_ema`). Keep the + * scale slot unchanged; the consumer falls back to scale=1.0 (bit- + * identical pre-Phase-2 cost magnitude). This mirrors the SP16-P1 + * MIN_HOLD_TEMPERATURE_FALLBACK=50.0 sentinel-detect path. */ + if (fabsf(observed_hold_rate - sentinel_observed_hold_rate) < EPS_F) { + return; + } + + /* Hold-rate overrun mapping. The denominator floor of 0.01 keeps + * the formula numerically stable when the target is configured to + * a tiny value (and protects against accidental sentinel-0 pollution + * on the target slot). target=0.20 (default) → effective denominator + * is the target itself; the floor only matters in pathological + * misconfiguration. Mirrors the SP16-P1 MIN_HOLD_TEMPERATURE producer + * for consistency (same input chain). */ + const float overrun = fmaxf(0.0f, observed_hold_rate - target_hold_rate); + const float 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; + + /* Bilateral clamp per `pearl_symmetric_clamp_audit.md`. Defends + * against arithmetic round-off pushing target outside the + * dimensional-safety window. */ + target_scale = fmaxf(scale_min, fminf(target_scale, scale_max)); + + /* 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 with target_scale. The consumer's + * cold-start fallback (scale=1.0 when isv[461] ≤ 0 or out-of-bounds) + * therefore kicks in for exactly one launch — the producer's first + * valid observation lands a real scale and the consumer reads it + * thereafter. Mirrors the SP14-P1-Producer Pearl-A pattern. */ + const float current = isv[scale_idx]; + float blended; + if (current <= EPS_F) { + blended = target_scale; + } else { + blended = (1.0f - alpha) * current + alpha * target_scale; + } + /* Re-clamp post-blend (defensive — handles malformed prior state + * outside [scale_min, scale_max]). */ + blended = fmaxf(scale_min, fminf(blended, scale_max)); + + isv[scale_idx] = blended; +} diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index 6e7c1cd5c..830901e11 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -445,6 +445,80 @@ 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. +// +// 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; + +// EMA blend rate. Mid-cadence (matches SP16-P1 MIN_HOLD_TEMPERATURE +// α=0.05): the scale should track realised Hold-rate overrun on the +// per-epoch cadence (faster than per-fold beliefs, slower than +// per-step signals — same input chain as MIN_HOLD_TEMPERATURE). +// Pearl-A bootstrap replaces sentinel directly on first valid +// observation; α=0.05 thereafter. +pub const HOLD_COST_SCALE_EMA_ALPHA: f32 = 0.05; + +pub const SP16_P2_HOLD_COST_SCALE_SLOT_BASE: usize = 461; +pub const SP16_P2_HOLD_COST_SCALE_SLOT_END: usize = 462; + #[cfg(test)] mod tests { use super::*; @@ -701,4 +775,41 @@ mod tests { SP14_AUDIT_4B_TEMP_SLOT_END, ISV_TOTAL_DIM, ); } + + /// 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); + // EMA α matches SP16-P1 MIN_HOLD_TEMPERATURE mid-cadence (same + // input chain — hold-rate overrun — same target cadence). + assert_eq!(HOLD_COST_SCALE_EMA_ALPHA, 0.05); + } + + #[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, + ); + } } diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index 23f894a9b..fb1751694 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -201,6 +201,18 @@ #define ISV_HOLD_RATE_TARGET_IDX 381 // SP13 v3 target Hold-pick rate (static, default 0.20) #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). +// ──────────────────────────────────────────────────────────────────────────── +#define ISV_HOLD_COST_SCALE_IDX 461 // SP16 Phase 2 adaptive Hold cost scale (sentinel 0.0; bounds [1, 25]; α=0.05) + // NOTE: this macro is kept at 383 for the historic SP13-era boundary marker // and is referenced only in commentary inside .cu/.cuh files. The actual // ISV bus dimension is the Rust `ISV_TOTAL_DIM` constant in diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 008bed078..f4768473b 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1173,6 +1173,11 @@ 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] — SP16 Phase 2 (2026-05-08) adaptive Hold cost scale (multiplicative scalar applied to the static per-bar Hold cost ISV[HOLD_COST_INDEX=380] at the 3 consumer sites in experience_kernels.cu — segment_complete branch line ~3089, per-bar positioned-Hold branch line ~3553, per-bar flat-Hold branch line ~3617). Producer kernel `hold_cost_scale_update_kernel` reads ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382] and ISV[HOLD_RATE_TARGET_INDEX=381] (same input chain as SP16-P1 MIN_HOLD_TEMPERATURE — same per-epoch cadence, same Pearl-A bootstrap pattern) and maps the overrun into a [1, 25] multiplier via `target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × clamp(max(0, observed - target) / max(target, 0.01), 0, 1)`, slow-EMA-blends into slot 461. 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. The adaptive scale pushes effective cost to ~0.15/bar at 100% overrun (observed=2× target), competitive with per-bar reward magnitudes but well below the SP12 asymmetric cap [-10, +5] so no clamp interaction. FoldReset sentinel SENTINEL_HOLD_COST_SCALE=0.0 — Pearl-A first-observation bootstrap: the consumer's cold-start fallback (scale=1.0 when slot ≤ 0 or out-of-bounds) fires for exactly one launch, then the producer's first valid observation lands a real scale and the consumer reads it thereafter. Cold-start sentinel guard at the producer: when observed_hold_rate is at SENTINEL_HOLD_RATE_OBSERVED=0.0 (matches sp13_hold_rate_observed_ema fold-reset sentinel), the kernel keeps slot 461 unchanged — the consumer falls back to scale=1.0. α=0.05 mid-cadence EMA. Bounds [1, 25] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md`. SCALE_MIN=1: the floor is `no extra penalty` (lower bound enforced by overrun being clamped at 0). SCALE_MAX=25: at this ceiling, a single Hold bar's effective cost ~0.15 — competitive with per-bar reward magnitudes ~0.01-0.1 but well below the SP12 cap, so no clamp interaction. Same launch-ordering constraint as SP16-P1: producer fires per-epoch boundary AFTER the per-step hold_rate_observer chain populates slot 382 (per `pearl_canary_input_freshness_launch_order.md`).", + }, RegistryEntry { name: "sp14_audit_4b_min_hold_temperature_adaptive", category: ResetCategory::FoldReset, diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index daf097938..25d1a5999 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -538,6 +538,32 @@ impl DQNTrainer { if let Err(e) = fused.trainer().launch_min_hold_temperature_update() { tracing::warn!(epoch, "launch_min_hold_temperature_update failed (non-fatal): {e}"); } + // SP16 Phase 2 (2026-05-08): adaptive Hold cost scale. + // + // Same input chain as SP16-P1 MIN_HOLD_TEMPERATURE (slots + // 382 + 381) but writes to slot 461 — a [1, 25] + // multiplier on the static per-bar Hold cost + // `ISV[ISV_HOLD_COST_IDX=380]` (~0.006). Per train- + // multi-seed-pfh9n post-mortem: realised Hold rate + // climbed 0.25 → 0.52 while cost stayed ~100× below + // per-bar reward magnitudes — Hold was effectively free + // and structural Q-bias toward zero-variance Hold + // dominated. The scale climbs to 25× at 100% overrun + // (observed=2× target), pushing effective cost to + // ~0.15/bar — competitive with per-bar reward + // magnitudes (~0.01-0.1) but well below the SP12 + // asymmetric cap [-10, +5] so no clamp interaction. + // + // Launch ordering: AFTER MIN_HOLD_TEMPERATURE (same + // input chain — both read slot 382). The per-step + // hold_rate_observer chain populates slot 382 every + // step, so by the time we arrive at the per-epoch + // 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}"); + } // 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. @@ -566,6 +592,36 @@ impl DQNTrainer { epoch, obs, tgt, overrun_norm, target_temp, temp, ); } + // 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; + 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; + 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}", + epoch, obs, tgt, overrun_norm, target_scale, scale, + ); + } } // D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary. @@ -8495,6 +8551,22 @@ 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, + ); + } + } // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // a stateful EMA) — rewrite the constructor's value at fold diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index d342ba751..10e436b0f 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -2135,6 +2135,291 @@ 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_EMA_ALPHA, HOLD_COST_SCALE_INDEX, HOLD_COST_SCALE_MAX, + HOLD_COST_SCALE_MIN, SENTINEL_HOLD_COST_SCALE, SENTINEL_HOLD_RATE_OBSERVED, + }; + use ml::cuda_pipeline::sp13_isv_slots::{ + HOLD_RATE_OBSERVED_EMA_INDEX, HOLD_RATE_TARGET_INDEX, + }; + + const HOLD_COST_SCALE_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/hold_cost_scale_update_kernel.cubin")); + + fn make_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(HOLD_COST_SCALE_UPDATE_CUBIN.to_vec()) + .expect("load hold_cost_scale_update_kernel cubin"); + module + .load_function("hold_cost_scale_update") + .expect("load hold_cost_scale_update function") + } + + /// Standard launch helper using the canonical SP16-P2 constants. + fn launch_canonical( + stream: &Arc, + kernel: &CudaFunction, + isv_ptr: u64, + ) { + unsafe { + stream + .launch_builder(kernel) + .arg(&isv_ptr) + .arg(&(HOLD_COST_SCALE_INDEX as i32)) + .arg(&(HOLD_RATE_OBSERVED_EMA_INDEX as i32)) + .arg(&(HOLD_RATE_TARGET_INDEX as i32)) + .arg(&SENTINEL_HOLD_COST_SCALE) + .arg(&SENTINEL_HOLD_RATE_OBSERVED) + .arg(&HOLD_COST_SCALE_MIN) + .arg(&HOLD_COST_SCALE_MAX) + .arg(&HOLD_COST_SCALE_EMA_ALPHA) + .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 Phase 0 (2026-05-08) — Q-by-action HEALTH_DIAG diagnostic. // diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 20a2ad294..a941bbc32 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -8646,3 +8646,80 @@ The original SP16 Phase 1 plan presumed the Fold-1 zero meant "producer didn't f 2. **Per-epoch cadence vs hold-rate observation cadence.** The MIN_HOLD_TEMPERATURE producer fires per-epoch boundary; the hold-rate EMA updates per-step. Fold-1 epoch-1 sees the EMA after only one epoch's worth of per-step updates — typically enough to push the EMA off sentinel 0.0, but with little time-history. The instrumentation will reveal the exact magnitude of the inputs at every transition, so we can falsify "the kernel responds correctly" by reading the smoke log. 3. **Phase 2 dependency check.** The original SP16 plan's Phase 2 increases the magnitude of the per-bar hold cost. That's orthogonal to this signal swap — the temperature controls when the min-hold *exit* penalty bites; the cost controls how much sitting in Hold itself costs. Both are independently active in production. + +## SP16 Phase 2: adaptive Hold cost scale via ISV[461] (2026-05-08) + +### Problem + +`train-multi-seed-pfh9n` HEALTH_DIAG epoch 2 emit: +``` +HEALTH_DIAG[2]: hold_pricing observed_rate=0.2501 target=0.2000 cost=0.006251 +``` + +By Fold-1 final, observed Hold rate had climbed to 0.516 (51.6%). The cost penalty stayed at ~0.006 across the entire training because the SP13-v3 P0a controller (gain=5, ceil=5×) saturated quickly at small magnitudes. Per-bar reward magnitudes at the same time were ~100× larger (popart=0.97, cf=0.65 in `reward_split`). Hold was effectively free; the structural Q-bias toward zero-variance Hold (Adam's `m/sqrt(v)` favours stable targets — Hold has zero per-bar PnL variance) had nothing to push against. + +The Phase 1 (revised) MIN_HOLD_TEMPERATURE addresses *when* the min-hold penalty bites at trade exit; this Phase 2 addresses *how much* the per-bar Hold cost itself costs. Both are independently active in production: the temperature is the sharpness of the soft-saturation curve at exit, the scale is the magnitude of the per-bar penalty itself. + +### Design + +ISV[HOLD_COST_SCALE_INDEX=461] — multiplicative scalar on the static per-bar Hold cost (`ISV[ISV_HOLD_COST_IDX=380]` ~= 0.006). Producer kernel `hold_cost_scale_update_kernel.cu` reads the SAME inputs as SP16-P1 (`ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and `ISV[HOLD_RATE_TARGET_INDEX=381]` — slots feeding the SP13 v3 P0a Hold-pricing controller chain) and maps overrun into a 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 + = 1 + 24 × overrun_norm +``` + +Pearl-A first-observation bootstrap (sentinel 0.0 → REPLACE) + α=0.05 mid-cadence EMA blend thereafter. Bilateral clamp pre- and post-blend per `pearl_symmetric_clamp_audit.md`. + +At observed=2× target (overrun_norm=1) → scale=25× → effective cost ~0.006 × 25 = 0.15/bar — competitive with per-bar reward magnitudes (~0.01-0.1) but well below the SP12 asymmetric cap [-10, +5] so no clamp interaction. At/below target → scale=1.0 (no extra penalty; bit-identical to pre-Phase-2 cost magnitude). + +### Slot allocation + +ISV[461..462) — `HOLD_COST_SCALE_INDEX=461` (sentinel 0.0, bounds [1.0, 25.0], α=0.05). Locked by `sp16_p2_hold_cost_scale_slot_layout_locked` test in `sp14_isv_slots.rs`. Bumped `ISV_TOTAL_DIM` 461 → 462. + +### Consumer migration (atomic per `feedback_no_partial_refactor.md`) + +The 3 per-bar Hold cost subtraction sites in `experience_kernels.cu` are migrated atomically in the same commit: + +1. **segment_complete branch** (line ~3089): `base_reward -= isv[ISV_HOLD_COST_IDX] * hold_cost_scale` +2. **per-bar positioned-Hold branch** (line ~3553): `r_micro -= isv[ISV_HOLD_COST_IDX] * hold_cost_scale` +3. **per-bar flat-Hold branch** (line ~3617): `r_opp_cost -= isv[ISV_HOLD_COST_IDX] * hold_cost_scale` + +Cold-start fallback at the consumer: when `isv[461]` is at sentinel 0.0 OR outside [1.0, 25.0], use scale=1.0 (bit-identical pre-Phase-2 cost magnitude — the consumer's first launch fires before the producer has populated the slot). The producer's first valid observation lands a real scale and the consumer reads it thereafter. + +### Producer launch ordering + +Per-epoch boundary, immediately AFTER the SP16-P1 MIN_HOLD_TEMPERATURE producer launch (training_loop.rs line ~538). Both share the same input slots (382 + 381) and the per-step `hold_rate_observer` chain populates slot 382 every step, so by the time we arrive at the per-epoch boundary the input has been refreshed N×B times. + +Per `pearl_canary_input_freshness_launch_order.md`: this launch order is the smoking-gun-prone constraint — if the producer fired BEFORE the per-step observer, slot 382 would still be at sentinel 0.0 and the producer would skip the update via the cold-start sentinel guard. + +### Files touched (atomic commit) + +- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — added `HOLD_COST_SCALE_INDEX=461` + sentinel + bounds + α + slot-layout-locked test (~70 lines). +- `crates/ml/src/cuda_pipeline/state_layout.cuh` — C #define mirror `ISV_HOLD_COST_SCALE_IDX=461`. +- `crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu` — NEW producer kernel (~165 lines, mirrors `min_hold_temperature_update_kernel.cu`). +- `crates/ml/build.rs` — cubin manifest entry. +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — bumped `ISV_TOTAL_DIM` 461 → 462; cubin static; trainer field; constructor init; `launch_hold_cost_scale_update` method. +- `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` — `HoldCostScaleUpdateOps` struct + import. +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — 3 consumer sites migrated atomically. +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — registry entry `sp16_phase2_hold_cost_scale_adaptive`. +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — launch call site after MIN_HOLD_TEMPERATURE; HEALTH_DIAG `hold_cost_scale_diag` instrumentation; reset dispatch arm. +- `crates/ml/tests/sp14_oracle_tests.rs` — 5 behavioral tests in module `sp16_phase2_hold_cost_scale_gpu`. + +### Validation + +- `cargo check -p ml --tests --all-targets` — clean (warnings only, no errors). +- `cargo test -p ml --lib --release sp14_isv_slots` — 18/18 PASS (slot layout locks). +- `cargo test -p ml --test sp14_oracle_tests --release sp16_phase2 -- --ignored` — 5/5 PASS. +- `cargo test -p ml --test sp14_oracle_tests --release -- --ignored` — 30/30 PASS (no SP14/SP16-P1 regressions). +- `cargo test -p ml --test sp15_phase1_oracle_tests --release -- --ignored` — 36/36 PASS (no SP15 regressions). + +### Concerns + +1. **25× ceiling justification.** SCALE_MAX=25 was chosen as the conservative point between "competitive with per-bar reward magnitudes" (~0.01-0.1) and "still below the SP12 asymmetric cap [-10, +5]". At 25×, single Hold bars cost ~0.15 — meaningful pressure but not catastrophic. The ratio could need tuning if Fold-2+ smoke shows persistent over-holding even at scale=25; the next phase candidate would lift SCALE_MAX to 50 or 100 (and re-evaluate the SP12 cap interaction). + +2. **Hold-cost consumer site found at exactly the canonical site.** `experience_kernels.cu:3089` (segment_complete branch) was the source of the `cost=0.006251` diagnostic value emitted at HEALTH_DIAG[2] in pfh9n. The 2 sibling sites (3553, 3617) cover the per-bar positioned-Hold and per-bar flat-Hold branches. All three are migrated atomically with the same fallback semantics so the per-bar Hold cost subtraction is consistent across all 3 reward composition paths. + +3. **Producer launch ordering.** The SP16-P1 producer fires AFTER the per-step `hold_rate_observer` chain via the `if (collector.executed) {…}` block in training_loop.rs. The new SP16-P2 producer fires immediately after SP16-P1 inside the same block, so the launch ordering is correct by construction. The `pearl_canary_input_freshness_launch_order.md` invariant is preserved.