From 8a25b330fcc2a0b197eccec242e1cfcffb8ccfdf Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 3 May 2026 21:36:11 +0200 Subject: [PATCH] fix(sp9): targets as Invariant-1 anchors + divergence sentinel handling MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two targeted fixes to the SP9 Kelly cold-start warmup floor (Fix 37, commit `48a8b9ee7`) for quirks observed in smoke-test-wrwkz on a 5-epoch L40S smoke. Atomic per `feedback_no_partial_refactor`. **Quirk 1 — EMA target saturation:** ep1 HEALTH_DIAG showed `stat_count_tgt=419 div_tgt=105002 temp_tgt=1.0 conf [stat=1.000 bhv=0.333 tmp=1.000] floor=0.0`. Pearl A sentinel-bootstrap on the 3 EMA target updaters (`kelly_*_target_ema_kernel.cu`) caused the first observation to replace the sentinel directly, so `target = first_obs`, `current/target = 1.0`, `confidence = 1.0`, `floor = base × (1 − 1.0) = 0` — defeating the cold-start mechanism. The principled fix per `feedback_isv_for_adaptive_bounds.md`: "sufficient" is defined in absolute terms via Invariant-1 numerical anchors. The 3 target ISV slots become constructor-written constants (written ONCE in trainer constructor, identical pattern to existing `config.cql_alpha → ISV[CQL_ALPHA_INDEX=48]` from Plan 1 Task 12). - ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] = 100.0 (trades) - ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] = 2.0 (ratio) - ISV[KELLY_TEMPORAL_TARGET_INDEX=335] = 5.0 (epochs) **Quirk 2 — divergence ratio explosion:** ep1 showed `divergence=70005`. Algebraically `intent_f / max(eval_f, 1e-6) = 0.07 / 1e-6 = 70 000` because `ISV[EVAL_DIST_F_INDEX=338]` was still at Pearl A sentinel 0 before the first val window populated it. Algebraically the warmup-floor kernel still derived `behavioral_conf = 0` (correct semantic), but the synthetic 70 000 in HEALTH_DIAG masked the real signal flow. Fix: explicit sentinel detection in `intent_eval_divergence_compute_kernel.cu`: if (eval_f < SENTINEL_THRESHOLD=1e-5) divergence = SENTINEL_DIVERGENCE=1e6 // forces behavioral_conf=0 else divergence = intent_f / eval_f Same `behavioral_conf = 0` outcome but with explicit provenance — HEALTH_DIAG now reads `divergence=1e6` until eval_f matures. **Atomic structure:** - DELETE 3 EMA target updater kernels (`.cu` files) - DELETE 3 entries from `crates/ml/build.rs::kernels_with_common` - DELETE 3 cubin static byte arrays + 3 struct fields + 3 cubin loaders + 3 fields in trainer construction tuple in `gpu_dqn_trainer.rs` - DELETE 3 launches in `launch_sp9_kelly_warmup_floor` (chain shrinks 5 → 2: q_var_mag_ema + main warmup-floor) - DELETE 3 scratch slots; SP5_SCRATCH_TOTAL 268 → 265; SCRATCH_SP9_EVAL_DIST_BASE slides 265 → 262 - ADD constructor-write of 3 Invariant-1 anchors (100.0, 2.0, 5.0) immediately before layout-fingerprint write - UPDATE 3 dispatch arms in `reset_named_state` to rewrite the Invariant-1 anchors at fold boundary (NOT sentinel 0) — these are constants, not stateful EMAs - UPDATE 3 registry descriptions to reflect Invariant-1 anchor semantic + smoke-test-wrwkz evidence - UPDATE `intent_eval_divergence_compute_kernel.cu` with sentinel-detect branch - ISV slot layout UNCHANGED (3 target slots @ ISV[333..336) remain); Wiener-buffer linear span (`SP5_PRODUCER_COUNT=165`) UNCHANGED — 3 wiener triples for deleted producers become reserved-unused like Pearl 6's [525..543) carve-out - audit doc Fix 37.1 entry with full provenance + smoke evidence **Verification:** - `SQLX_OFFLINE=true cargo check -p ml` — clean (only pre-existing 18 warnings; no new errors or warnings). - `SQLX_OFFLINE=true cargo test -p ml --lib state_reset` — 4/4 pass including `every_fold_and_soft_reset_entry_has_dispatch_arm`. - `SQLX_OFFLINE=true cargo test -p ml --lib sp5_isv_slots` — 9/9 pass including `sp9_slots_contiguous_above_sp8_block`. **Expected smoke signature post-fix:** - floor: starts ≈ base_floor × 1.0, decays as confidence accumulates - divergence: 1e6 until first val window, small (≤ 5) afterward - conf: stat=count/100, bhv=0 until eval_dist matures, tmp=ep/5 - combined_conf reaches 1.0 only when at least one axis genuinely matures (typically temporal at epoch 5 in fold 0) Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 23 +- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 205 ++++++++---------- .../intent_eval_divergence_compute_kernel.cu | 48 +++- .../kelly_divergence_target_ema_kernel.cu | 40 ---- .../kelly_sample_count_target_ema_kernel.cu | 42 ---- .../kelly_temporal_target_ema_kernel.cu | 41 ---- .../kelly_warmup_floor_compute_kernel.cu | 6 +- .../src/trainers/dqn/state_reset_registry.rs | 6 +- .../src/trainers/dqn/trainer/training_loop.rs | 23 +- docs/dqn-wire-up-audit.md | 167 ++++++++++++++ 10 files changed, 334 insertions(+), 267 deletions(-) delete mode 100644 crates/ml/src/cuda_pipeline/kelly_divergence_target_ema_kernel.cu delete mode 100644 crates/ml/src/cuda_pipeline/kelly_sample_count_target_ema_kernel.cu delete mode 100644 crates/ml/src/cuda_pipeline/kelly_temporal_target_ema_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 8c8e866c8..fb26c76c2 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -525,25 +525,28 @@ fn main() { // `pearl_controller_anchors_isv_driven.md`. "loss_balance_max_budget_compute_kernel.cu", // SP9 Fix 37 (2026-05-03): Kelly cold-start warmup floor — fully - // ISV-driven. Six new producer kernels per + // ISV-driven. Four producer kernels per // `pearl_cold_start_exit_signal_or.md` and // `pearl_controller_anchors_isv_driven.md`: // 1. eval_intent_dist_compute — replaces DtoH read_eval_intent_magnitude_distribution // 2. intent_eval_divergence_compute — behavioral-axis numerator // 3. q_var_mag_ema_compute — base_floor self-relative baseline - // 4. kelly_sample_count_target_ema_compute — statistical-axis denominator - // 5. kelly_divergence_target_ema_compute — behavioral-axis denominator - // 6. kelly_temporal_target_ema_compute — temporal-axis denominator - // 7. kelly_warmup_floor_compute — main producer (combines all axes) - // All 7 are single-block kernels writing to producer_step_scratch_buf, + // 4. kelly_warmup_floor_compute — main producer (combines all axes) + // All 4 are single-block kernels writing to producer_step_scratch_buf, // chained through apply_pearls_ad_kernel for Pearl A sentinel-bootstrap - // + Pearl D Wiener-α steady-state smoothing. + // + Pearl D Wiener-α steady-state smoothing. Per Fix 37.1 + // (2026-05-03): the 3 EMA target updaters that previously sat between + // q_var_mag_ema and kelly_warmup_floor are deleted — the targets at + // ISV[333..336) are Invariant-1 numerical anchors (sample-sufficiency + // count, healthy divergence ratio, fold-temporal threshold) written + // ONCE at constructor time per `feedback_isv_for_adaptive_bounds.md`, + // not EMA-tracked. Pearl A sentinel-bootstrap immediately saturated + // the EMAs to the first observation, which made conf=1.0 in epoch 1 + // (`floor=0`) — defeating the cold-start mechanism. Constructor-write + // gives a fixed denominator so the ratio yields meaningful confidence. "eval_intent_dist_compute_kernel.cu", "intent_eval_divergence_compute_kernel.cu", "q_var_mag_ema_compute_kernel.cu", - "kelly_sample_count_target_ema_kernel.cu", - "kelly_divergence_target_ema_kernel.cu", - "kelly_temporal_target_ema_kernel.cu", "kelly_warmup_floor_compute_kernel.cu", // SP7 (2026-05-03): GPU dispatch for cql/c51 budget consumer. // Eliminates capture-time host-branch freeze in diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 6f338b357..1688e3c0b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -510,27 +510,6 @@ static SP9_INTENT_EVAL_DIVERGENCE_COMPUTE_CUBIN: &[u8] = static SP9_Q_VAR_MAG_EMA_COMPUTE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/q_var_mag_ema_compute_kernel.cubin")); -/// SP9 (Fix 37, 2026-05-03): EMA target for the statistical-confidence -/// axis. Single block, 1 thread; reads `ISV[KELLY_SAMPLE_COUNT_INDEX=283]` -/// and writes to scratch. Downstream apply_pearls_ad_kernel smooths into -/// ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333]. -static SP9_KELLY_SAMPLE_COUNT_TARGET_EMA_CUBIN: &[u8] = - include_bytes!(concat!(env!("OUT_DIR"), "/kelly_sample_count_target_ema_kernel.cubin")); - -/// SP9 (Fix 37, 2026-05-03): EMA target for the behavioral-confidence -/// axis. Single block, 1 thread; reads `ISV[INTENT_EVAL_DIVERGENCE_INDEX=332]` -/// and writes to scratch. Downstream apply_pearls_ad_kernel smooths into -/// ISV[KELLY_DIVERGENCE_TARGET_INDEX=334]. -static SP9_KELLY_DIVERGENCE_TARGET_EMA_CUBIN: &[u8] = - include_bytes!(concat!(env!("OUT_DIR"), "/kelly_divergence_target_ema_kernel.cubin")); - -/// SP9 (Fix 37, 2026-05-03): EMA target for the temporal-confidence axis. -/// Single block, 1 thread; reads `ISV[EPOCH_IDX_INDEX=39]` and writes to -/// scratch. Downstream apply_pearls_ad_kernel smooths into -/// ISV[KELLY_TEMPORAL_TARGET_INDEX=335]. -static SP9_KELLY_TEMPORAL_TARGET_EMA_CUBIN: &[u8] = - include_bytes!(concat!(env!("OUT_DIR"), "/kelly_temporal_target_ema_kernel.cubin")); - /// SP9 (Fix 37, 2026-05-03): main Kelly cold-start warmup-floor producer. /// Single block, 1 thread; reads 8 ISV inputs (q_var_mag, q_var_mag_ema, /// kelly_sample_count, sample_count_target, intent_eval_divergence, @@ -1210,12 +1189,14 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// 1 float for kelly_warmup_floor (SCRATCH_SP9_KELLY_WARMUP_FLOOR=259, at [259..260)) /// 1 float for q_var_mag_ema (SCRATCH_SP9_Q_VAR_MAG_EMA=260, at [260..261)) /// 1 float for intent_eval_divergence (SCRATCH_SP9_INTENT_EVAL_DIVERGENCE=261, at [261..262)) -/// 1 float for sample_count_target (SCRATCH_SP9_SAMPLE_COUNT_TARGET=262, at [262..263)) -/// 1 float for divergence_target (SCRATCH_SP9_DIVERGENCE_TARGET=263, at [263..264)) -/// 1 float for temporal_target (SCRATCH_SP9_TEMPORAL_TARGET=264, at [264..265)) -/// 3 floats for eval_dist[Q,H,F] (SCRATCH_SP9_EVAL_DIST_BASE=265, at [265..268)) -/// Combined: 259 + 9 = 268 scratch slots [0..268). -pub const SP5_SCRATCH_TOTAL: usize = 268; +/// 3 floats for eval_dist[Q,H,F] (SCRATCH_SP9_EVAL_DIST_BASE=262, at [262..265)) +/// Fix 37.1 (2026-05-03): the 3 EMA target updaters that previously sat +/// at [262..265) (sample_count_target / divergence_target / temporal_target) +/// are deleted — their target ISV slots are constructor-written +/// Invariant-1 anchors per `feedback_isv_for_adaptive_bounds.md` and +/// require no scratch storage. The eval_dist base slides 265 → 262. +/// Combined: 259 + 6 = 265 scratch slots [0..265). +pub const SP5_SCRATCH_TOTAL: usize = 265; /// SP5 Layer D Task D1 (rewrite, 2026-05-02): scratch index base for /// `pnl_aggregation_update`. @@ -1307,17 +1288,20 @@ pub const SCRATCH_LB_MAX_BUDGET_CQL: usize = 251; // 251..255 pub const SCRATCH_LB_MAX_BUDGET_C51: usize = 255; // 255..259 /// SP9 (Fix 37, 2026-05-03): scratch slots for Kelly cold-start warmup -/// floor + EMA targets + intent/eval divergence + eval_dist GPU -/// migration. Each producer kernel writes to its scratch slot; downstream +/// floor + intent/eval divergence + eval_dist GPU migration. Each +/// producer kernel writes to its scratch slot; downstream /// `apply_pearls_ad_kernel` smooths into the corresponding ISV slot via /// Pearl A sentinel-bootstrap + Pearl D Wiener-α steady-state smoothing. +/// +/// Fix 37.1 (2026-05-03): the 3 EMA target updater scratch slots +/// (sample_count_target / divergence_target / temporal_target) are +/// deleted — their target ISV slots are constructor-written +/// Invariant-1 anchors per `feedback_isv_for_adaptive_bounds.md` and +/// require no scratch storage. The eval_dist base slides 265 → 262. pub const SCRATCH_SP9_KELLY_WARMUP_FLOOR: usize = 259; // → ISV[KELLY_WARMUP_FLOOR_INDEX=330] pub const SCRATCH_SP9_Q_VAR_MAG_EMA: usize = 260; // → ISV[Q_VAR_MAG_EMA_INDEX=331] pub const SCRATCH_SP9_INTENT_EVAL_DIVERGENCE: usize = 261; // → ISV[INTENT_EVAL_DIVERGENCE_INDEX=332] -pub const SCRATCH_SP9_SAMPLE_COUNT_TARGET: usize = 262; // → ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] -pub const SCRATCH_SP9_DIVERGENCE_TARGET: usize = 263; // → ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] -pub const SCRATCH_SP9_TEMPORAL_TARGET: usize = 264; // → ISV[KELLY_TEMPORAL_TARGET_INDEX=335] -pub const SCRATCH_SP9_EVAL_DIST_BASE: usize = 265; // → ISV[EVAL_DIST_Q/H/F_INDEX=336..339) +pub const SCRATCH_SP9_EVAL_DIST_BASE: usize = 262; // → ISV[EVAL_DIST_Q/H/F_INDEX=336..339) /// SP5 Task A7: scratch index base for pearl_8_trail_update trail_dist[4] output block. /// Slots [199..203): per-direction trail-stop distance (Short=0, Hold=1, Long=2, Flat=3). @@ -4605,7 +4589,7 @@ pub struct GpuDqnTrainer { /// SP9 (Fix 37, 2026-05-03): GPU producer for eval-side magnitude /// intent distribution. Replaces DtoH `read_eval_intent_magnitude_distribution()`. /// Single-block, 3-thread kernel (one per Q/H/F bin); writes 3 floats - /// to `scratch[SCRATCH_SP9_EVAL_DIST_BASE=265..268)`. Followed by 3 + /// to `scratch[SCRATCH_SP9_EVAL_DIST_BASE=262..265)`. Followed by 3 /// apply_pearls_ad_kernel launches → ISV[EVAL_DIST_Q/H/F_INDEX=336..339). /// Loaded from `eval_intent_dist_compute_kernel.cubin`. sp9_eval_intent_dist_compute_kernel: CudaFunction, @@ -4623,26 +4607,6 @@ pub struct GpuDqnTrainer { /// apply_pearls_ad_kernel launch → ISV[Q_VAR_MAG_EMA_INDEX=331]. /// Loaded from `q_var_mag_ema_compute_kernel.cubin`. sp9_q_var_mag_ema_compute_kernel: CudaFunction, - /// SP9 (Fix 37, 2026-05-03): EMA target for the statistical-confidence - /// axis. Single-block, single-thread; reads `ISV[KELLY_SAMPLE_COUNT_INDEX=283]`; - /// writes to `scratch[SCRATCH_SP9_SAMPLE_COUNT_TARGET=262]`. Followed - /// by 1 apply_pearls_ad_kernel launch → - /// ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333]. Loaded from - /// `kelly_sample_count_target_ema_kernel.cubin`. - sp9_kelly_sample_count_target_ema_kernel: CudaFunction, - /// SP9 (Fix 37, 2026-05-03): EMA target for the behavioral-confidence - /// axis. Single-block, single-thread; reads - /// `ISV[INTENT_EVAL_DIVERGENCE_INDEX=332]`; writes to - /// `scratch[SCRATCH_SP9_DIVERGENCE_TARGET=263]`. Followed by 1 - /// apply_pearls_ad_kernel launch → ISV[KELLY_DIVERGENCE_TARGET_INDEX=334]. - /// Loaded from `kelly_divergence_target_ema_kernel.cubin`. - sp9_kelly_divergence_target_ema_kernel: CudaFunction, - /// SP9 (Fix 37, 2026-05-03): EMA target for the temporal-confidence - /// axis. Single-block, single-thread; reads `ISV[EPOCH_IDX_INDEX=39]`; - /// writes to `scratch[SCRATCH_SP9_TEMPORAL_TARGET=264]`. Followed by 1 - /// apply_pearls_ad_kernel launch → ISV[KELLY_TEMPORAL_TARGET_INDEX=335]. - /// Loaded from `kelly_temporal_target_ema_kernel.cubin`. - sp9_kelly_temporal_target_ema_kernel: CudaFunction, /// SP9 (Fix 37, 2026-05-03): main Kelly cold-start warmup-floor producer. /// Single-block, single-thread; reads 8 ISV inputs and computes /// base_floor × (1 − max(statistical_conf, behavioral_conf, temporal_conf)) @@ -12096,20 +12060,40 @@ impl GpuDqnTrainer { } /// SP9 (Fix 37, 2026-05-03): launch the four cold-path EMA producers - /// (q_var_mag_ema + 3 EMA targets) and the main warmup-floor producer. + /// (q_var_mag_ema) and the main warmup-floor producer. /// - /// All five kernels are single-block, single-thread cold-path producers + /// Both kernels are single-block, single-thread cold-path producers /// chained through `apply_pearls_ad_kernel`. The launch order matches /// the SP9 dependency DAG: /// 1. q_var_mag_ema_compute — reads ISV[FLATNESS_BASE+1] - /// 2. kelly_sample_count_target_ema_compute — reads ISV[KELLY_SAMPLE_COUNT_INDEX] - /// 3. kelly_divergence_target_ema_compute — reads ISV[INTENT_EVAL_DIVERGENCE_INDEX] - /// 4. kelly_temporal_target_ema_compute — reads ISV[EPOCH_IDX_INDEX] - /// 5. kelly_warmup_floor_compute — reads all 8 inputs above + ISV - /// targets just produced + /// 2. kelly_warmup_floor_compute — reads 8 inputs (q_var_mag, + /// q_var_mag_ema, kelly_sample_count, + /// sample_count_target, intent_eval_divergence, + /// divergence_target, epoch_idx, + /// temporal_target). The 3 target + /// slots are constructor-written + /// Invariant-1 anchors per Fix 37.1 + /// (`feedback_isv_for_adaptive_bounds.md`). + /// + /// Fix 37.1 (2026-05-03): the 3 EMA target updaters that previously sat + /// between #1 and #2 (kelly_sample_count_target_ema / + /// kelly_divergence_target_ema / kelly_temporal_target_ema) are deleted. + /// Their target slots ISV[333..336) are now numerical-stability anchors + /// — `KELLY_SAMPLE_COUNT_TARGET_INDEX = 100.0`, + /// `KELLY_DIVERGENCE_TARGET_INDEX = 2.0`, + /// `KELLY_TEMPORAL_TARGET_INDEX = 5.0` — written ONCE at trainer + /// constructor time. Pearl A sentinel-bootstrap on the original EMA + /// path saturated all three to the very first observation in epoch 1 + /// (`stat_count_tgt=419, div_tgt=70005, temp_tgt=1`), which made + /// `current/target = 1.0`, `confidence = 1.0` and + /// `floor = base × (1 − 1.0) = 0` — defeating the cold-start floor. + /// A fixed denominator yields a meaningful target-relative ratio that + /// rises monotonically as samples / divergence-stability / fold-time + /// accumulate. /// /// Must run AFTER: - /// * `launch_intent_eval_divergence_compute` (for input #3) + /// * `launch_intent_eval_divergence_compute` (so ISV[332] is fresh + /// before the warmup-floor kernel reads it) /// Must run BEFORE: /// * any `unified_env_step_core` (which reads ISV[KELLY_WARMUP_FLOOR_INDEX]) pub(crate) fn launch_sp9_kelly_warmup_floor(&self) -> Result<(), MLError> { @@ -12185,34 +12169,9 @@ impl GpuDqnTrainer { "q_var_mag_ema_compute", )?; - // 2. kelly_sample_count_target_ema (reads ISV[283]) - launch_simple_ema( - &self.sp9_kelly_sample_count_target_ema_kernel, - KELLY_SAMPLE_COUNT_INDEX, - SCRATCH_SP9_SAMPLE_COUNT_TARGET, - KELLY_SAMPLE_COUNT_TARGET_INDEX, - "kelly_sample_count_target_ema_compute", - )?; - - // 3. kelly_divergence_target_ema (reads ISV[332]) - launch_simple_ema( - &self.sp9_kelly_divergence_target_ema_kernel, - INTENT_EVAL_DIVERGENCE_INDEX, - SCRATCH_SP9_DIVERGENCE_TARGET, - KELLY_DIVERGENCE_TARGET_INDEX, - "kelly_divergence_target_ema_compute", - )?; - - // 4. kelly_temporal_target_ema (reads ISV[39] EPOCH_IDX_INDEX) - launch_simple_ema( - &self.sp9_kelly_temporal_target_ema_kernel, - EPOCH_IDX_INDEX, - SCRATCH_SP9_TEMPORAL_TARGET, - KELLY_TEMPORAL_TARGET_INDEX, - "kelly_temporal_target_ema_compute", - )?; - - // 5. kelly_warmup_floor_compute — combines all 8 inputs. + // 2. kelly_warmup_floor_compute — combines all 8 inputs (target + // slots ISV[333..336) read directly from constructor-written + // Invariant-1 anchors per Fix 37.1; no producer chain). let q_var_mag_idx_i32: i32 = (FLATNESS_BASE + 1) as i32; let q_var_mag_ema_idx_i32: i32 = Q_VAR_MAG_EMA_INDEX as i32; let kelly_sample_count_idx_i32: i32 = KELLY_SAMPLE_COUNT_INDEX as i32; @@ -15552,27 +15511,6 @@ impl GpuDqnTrainer { module.load_function("q_var_mag_ema_compute") .map_err(|e| MLError::ModelError(format!("sp9 q_var_mag_ema_compute load: {e}")))? }; - let sp9_kelly_sample_count_target_ema_kernel = { - let module = stream.context() - .load_cubin(SP9_KELLY_SAMPLE_COUNT_TARGET_EMA_CUBIN.to_vec()) - .map_err(|e| MLError::ModelError(format!("sp9 kelly_sample_count_target_ema cubin load: {e}")))?; - module.load_function("kelly_sample_count_target_ema_compute") - .map_err(|e| MLError::ModelError(format!("sp9 kelly_sample_count_target_ema_compute load: {e}")))? - }; - let sp9_kelly_divergence_target_ema_kernel = { - let module = stream.context() - .load_cubin(SP9_KELLY_DIVERGENCE_TARGET_EMA_CUBIN.to_vec()) - .map_err(|e| MLError::ModelError(format!("sp9 kelly_divergence_target_ema cubin load: {e}")))?; - module.load_function("kelly_divergence_target_ema_compute") - .map_err(|e| MLError::ModelError(format!("sp9 kelly_divergence_target_ema_compute load: {e}")))? - }; - let sp9_kelly_temporal_target_ema_kernel = { - let module = stream.context() - .load_cubin(SP9_KELLY_TEMPORAL_TARGET_EMA_CUBIN.to_vec()) - .map_err(|e| MLError::ModelError(format!("sp9 kelly_temporal_target_ema cubin load: {e}")))?; - module.load_function("kelly_temporal_target_ema_compute") - .map_err(|e| MLError::ModelError(format!("sp9 kelly_temporal_target_ema_compute load: {e}")))? - }; let sp9_kelly_warmup_floor_compute_kernel = { let module = stream.context() .load_cubin(SP9_KELLY_WARMUP_FLOOR_COMPUTE_CUBIN.to_vec()) @@ -17674,6 +17612,48 @@ impl GpuDqnTrainer { * `training_loop.rs` prevents the producer launch until * a meaningful prev value exists). */ *sig_ptr.add(Q_DRIFT_RATE_INDEX) = 0.0_f32; + // SP9 Fix 37.1 (2026-05-03): the 3 Kelly cold-start exit-axis + // targets are Invariant-1 numerical anchors per + // `feedback_isv_for_adaptive_bounds.md`, not EMA-tracked. + // The original Fix 37 routed each through a single-input + // EMA producer + Pearl A sentinel-bootstrap + Pearl D + // Wiener-α; smoke-test-wrwkz revealed this caused the + // first observation to replace the sentinel directly, + // making target == current_obs in epoch 1 (`stat_count_tgt + // = 419`, `temp_tgt = 1`). Consequently `current/target = + // 1.0`, `confidence = 1.0`, `floor = base × (1 − 1.0) = 0` + // — defeating the cold-start mechanism entirely. + // + // The principled fix per the pearl: define what + // "sufficient" means in absolute terms via Invariant-1 + // anchors. The target-relative ratio + // `current_observation / fixed_anchor` then yields a + // self-normalised confidence in [0, 1] that rises + // monotonically as samples / divergence-stability / + // fold-time accumulate. + // + // ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX = 333] = 100.0 + // — minimum trade samples before statistical confidence + // releases the cap. + // ISV[KELLY_DIVERGENCE_TARGET_INDEX = 334] = 2.0 + // — divergence ratio above which intent-vs-eval are + // considered diverged enough to keep the cap closed. + // ISV[KELLY_TEMPORAL_TARGET_INDEX = 335] = 5.0 + // — minimum epochs in fold before temporal confidence + // releases the cap (safety net per Risk 3 in spec). + // + // FoldReset rewrites the same anchors (in `reset_named_state`) + // since these are constants, not stateful EMAs — they must + // never reach 0 sentinel between folds. + use crate::cuda_pipeline::sp5_isv_slots::{ + KELLY_SAMPLE_COUNT_TARGET_INDEX, + KELLY_DIVERGENCE_TARGET_INDEX, + KELLY_TEMPORAL_TARGET_INDEX, + }; + *sig_ptr.add(KELLY_SAMPLE_COUNT_TARGET_INDEX) = 100.0_f32; + *sig_ptr.add(KELLY_DIVERGENCE_TARGET_INDEX) = 2.0_f32; + *sig_ptr.add(KELLY_TEMPORAL_TARGET_INDEX) = 5.0_f32; + // Layout fingerprint (ISV[58..60)). Compile-time structural hash of // the slot layout; checkpoint load fails-fast on mismatch. // Stored as a u64 split across two f32 lanes using raw bit-cast so @@ -18799,9 +18779,6 @@ impl GpuDqnTrainer { sp9_eval_intent_dist_compute_kernel, sp9_intent_eval_divergence_compute_kernel, sp9_q_var_mag_ema_compute_kernel, - sp9_kelly_sample_count_target_ema_kernel, - sp9_kelly_divergence_target_ema_kernel, - sp9_kelly_temporal_target_ema_kernel, sp9_kelly_warmup_floor_compute_kernel, lb_budget_dispatch_kernel, saxpy_f32_dev_ptr_aux, diff --git a/crates/ml/src/cuda_pipeline/intent_eval_divergence_compute_kernel.cu b/crates/ml/src/cuda_pipeline/intent_eval_divergence_compute_kernel.cu index dd4f300c9..8d4267fd6 100644 --- a/crates/ml/src/cuda_pipeline/intent_eval_divergence_compute_kernel.cu +++ b/crates/ml/src/cuda_pipeline/intent_eval_divergence_compute_kernel.cu @@ -26,8 +26,11 @@ // ratio rises sharply, signalling the controller. The behavioral_conf axis // of the warmup-floor formula maps `clamp(1 − divergence/divergence_target, 0, 1)` // — divergence above target → confidence = 0 (cap stays closed); divergence -// at or below target → confidence rises (cap relaxes). EMA target itself is -// adaptive (Wiener-α via Pearl D) — see `kelly_divergence_target_ema_kernel.cu`. +// at or below target → confidence rises (cap relaxes). The divergence +// target ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] is an Invariant-1 numerical +// anchor (2.0) per Fix 37.1 (`feedback_isv_for_adaptive_bounds.md`), +// constructor-written once and rewritten on FoldReset — see audit Fix 37.1 +// for why the original Fix 37 EMA-tracked target was abandoned. // // Single block, single thread; cheap arithmetic. No atomicAdd // (`feedback_no_atomicadd.md`); `__threadfence_system()` after the write. @@ -47,10 +50,14 @@ extern "C" __global__ void intent_eval_divergence_compute( ) { if (blockIdx.x != 0 || threadIdx.x != 0) return; - /* Numerical-stability anchor (Invariant 1) — prevents divide-by-zero - * when intent_f or eval_f is exactly zero. Same EPS scale used by - * `train_active_frac_compute_kernel` and the SP4 Pearl D recipe. */ - const float EPS_DIV = 1e-6f; + /* Numerical-stability anchor (Invariant 1) — distinguishes a freshly- + * fold-reset / never-populated `eval_f` (Pearl A sentinel = 0) from a + * tiny-but-real `eval_f` value. Anything strictly below this threshold + * is treated as "uninitialised" because `eval_dist` writes are + * naturally bounded by [0, 1] from `eval_intent_dist_compute_kernel`'s + * normalised output, and 1e-5 sits well below any realisable value + * after even a single observation enters the Pearl A bootstrap path. */ + const float SENTINEL_THRESHOLD = 1e-5f; /* Sum action_counts over Full magnitude bins: exp_idx = dir*3 + 2 for * dir ∈ {0..4}. Bins 2 (Short Full), 5 (Hold Full), 8 (Long Full), @@ -74,11 +81,30 @@ extern "C" __global__ void intent_eval_divergence_compute( if (eval_f < 0.0f) eval_f = 0.0f; if (eval_f > 1.0f) eval_f = 1.0f; - /* Divergence = intent_f / eval_f. Both numerator and denominator are - * floored at EPS_DIV to prevent NaN/Inf when either is 0. Range under - * healthy training: ~1; under val-Flat-collapse: large (intent up, - * eval down). */ - float divergence = fmaxf(intent_f, EPS_DIV) / fmaxf(eval_f, EPS_DIV); + /* Fix 37.1 (2026-05-03): handle Pearl A sentinel for `eval_f`. Before + * the first val-window has populated `ISV[EVAL_DIST_F_INDEX=338]`, + * the slot still carries its sentinel 0 (per + * `pearl_first_observation_bootstrap.md`). Reading it as a real + * denominator and flooring at any `EPS_DIV` produced a synthetic + * `divergence ≈ intent_f / 1e-6 ≈ 70 000`, which the warmup-floor + * kernel then mapped to `behavioral_conf = clamp(1 − 70000/2, 0, 1) = 0` + * — algebraically the right answer "no behavioral signal yet", but + * coupled to a magic-floor `EPS_DIV` and easily missed in + * smoke logs (the synthetic 70 000 looked like a runaway). + * Instead: explicitly emit `SENTINEL_DIVERGENCE` (large enough that + * `divergence / divergence_target ≫ 1` regardless of target) so the + * downstream warmup-floor kernel's `behavioral_conf` clamp lands + * cleanly at 0 with provenance traceable in HEALTH_DIAG. Once + * `eval_dist_compute` fires its first val-window observation, + * Pearl A's first-observation-replacement raises `eval_f` above the + * sentinel threshold and the real ratio takes over. */ + const float SENTINEL_DIVERGENCE = 1e6f; + float divergence; + if (eval_f < SENTINEL_THRESHOLD) { + divergence = SENTINEL_DIVERGENCE; + } else { + divergence = intent_f / eval_f; + } scratch_buf[scratch_idx] = divergence; __threadfence_system(); diff --git a/crates/ml/src/cuda_pipeline/kelly_divergence_target_ema_kernel.cu b/crates/ml/src/cuda_pipeline/kelly_divergence_target_ema_kernel.cu deleted file mode 100644 index 28e10f5aa..000000000 --- a/crates/ml/src/cuda_pipeline/kelly_divergence_target_ema_kernel.cu +++ /dev/null @@ -1,40 +0,0 @@ -// crates/ml/src/cuda_pipeline/kelly_divergence_target_ema_kernel.cu -// -// SP9 (Fix 37, 2026-05-03): GPU producer for EMA-tracked intent vs eval -// divergence target — the behavioral-confidence axis denominator of the -// OR'd cold-start exit per `pearl_cold_start_exit_signal_or.md`. -// -// Reads ISV[INTENT_EVAL_DIVERGENCE_INDEX=332] (the SP9 ratio -// `intent_f / eval_f` produced by `intent_eval_divergence_compute_kernel.cu`). -// Writes that observation to `scratch_buf[scratch_idx]`. -// -// Downstream `apply_pearls_ad_kernel` chains: -// - Pearl A (sentinel 0 at fold boundary → first observation replaces slot) -// - Pearl D (Wiener-optimal α adapts to divergence-stability) -// → ISV[KELLY_DIVERGENCE_TARGET_INDEX=334]. -// -// Per `pearl_controller_anchors_isv_driven.md`: the "healthy divergence" -// threshold is itself signal-driven via Pearl D smoothing — when divergence -// is stable around any value, that value becomes the target; when divergence -// shifts (val-Flat-collapse onset), the EMA lags and the consumer's ratio -// `divergence / target` rises sharply, dropping behavioral_conf to 0. -// -// Single-block, single-thread cold-path kernel. No atomicAdd; no shared -// memory; `__threadfence_system()` after the write. - -#include - -extern "C" __global__ void kelly_divergence_target_ema_compute( - const float* __restrict__ isv_signals, - int intent_eval_divergence_isv_index, // INTENT_EVAL_DIVERGENCE_INDEX = 332 - float* __restrict__ scratch_buf, - int scratch_idx -) { - if (blockIdx.x != 0 || threadIdx.x != 0) return; - - float observed = isv_signals[intent_eval_divergence_isv_index]; - if (observed < 0.0f) observed = 0.0f; - - scratch_buf[scratch_idx] = observed; - __threadfence_system(); -} diff --git a/crates/ml/src/cuda_pipeline/kelly_sample_count_target_ema_kernel.cu b/crates/ml/src/cuda_pipeline/kelly_sample_count_target_ema_kernel.cu deleted file mode 100644 index dcd6eaa77..000000000 --- a/crates/ml/src/cuda_pipeline/kelly_sample_count_target_ema_kernel.cu +++ /dev/null @@ -1,42 +0,0 @@ -// crates/ml/src/cuda_pipeline/kelly_sample_count_target_ema_kernel.cu -// -// SP9 (Fix 37, 2026-05-03): GPU producer for EMA-tracked Kelly -// sample-sufficiency target — the statistical-confidence axis denominator -// of the OR'd cold-start exit per `pearl_cold_start_exit_signal_or.md`. -// -// Reads ISV[KELLY_SAMPLE_COUNT_INDEX=283] (Pearl 6's cumulative-via-max -// trade count). Writes that observation to `scratch_buf[scratch_idx]`. -// -// Downstream `apply_pearls_ad_kernel` chains: -// - Pearl A (sentinel 0 at fold boundary → first observation replaces slot) -// - Pearl D (Wiener-optimal α — variance ratio adapts the smoothing rate) -// → ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333]. -// -// Why an EMA-tracked target instead of a hardcoded constant: per -// `pearl_controller_anchors_isv_driven.md`, the threshold for "enough -// statistical confidence to release the cap" should track the running -// value, not be tuned. Pearl D's variance-ratio-derived α slows the EMA -// when the signal is stable and speeds it up when it shifts — Wiener- -// optimal blending. The target-relative ratio `kelly_sample_count / -// sample_count_target` becomes a self-normalised confidence in [0, 1] -// (clamped in the consumer). -// -// Single-block, single-thread cold-path kernel. No atomicAdd; no shared -// memory; `__threadfence_system()` after the write. - -#include - -extern "C" __global__ void kelly_sample_count_target_ema_compute( - const float* __restrict__ isv_signals, - int kelly_sample_count_isv_index, // KELLY_SAMPLE_COUNT_INDEX = 283 - float* __restrict__ scratch_buf, - int scratch_idx -) { - if (blockIdx.x != 0 || threadIdx.x != 0) return; - - float observed = isv_signals[kelly_sample_count_isv_index]; - if (observed < 0.0f) observed = 0.0f; - - scratch_buf[scratch_idx] = observed; - __threadfence_system(); -} diff --git a/crates/ml/src/cuda_pipeline/kelly_temporal_target_ema_kernel.cu b/crates/ml/src/cuda_pipeline/kelly_temporal_target_ema_kernel.cu deleted file mode 100644 index 6f57cadb9..000000000 --- a/crates/ml/src/cuda_pipeline/kelly_temporal_target_ema_kernel.cu +++ /dev/null @@ -1,41 +0,0 @@ -// crates/ml/src/cuda_pipeline/kelly_temporal_target_ema_kernel.cu -// -// SP9 (Fix 37, 2026-05-03): GPU producer for EMA-tracked temporal target -// — the temporal-confidence axis denominator of the OR'd cold-start exit -// per `pearl_cold_start_exit_signal_or.md`. -// -// Reads ISV[EPOCH_IDX_INDEX=39] (the existing per-fold epoch counter). -// Writes that observation to `scratch_buf[scratch_idx]`. -// -// Downstream `apply_pearls_ad_kernel` chains: -// - Pearl A (sentinel 0 at fold boundary → first observation replaces slot) -// - Pearl D (Wiener-optimal α adapts to the steady-state epoch growth rate) -// → ISV[KELLY_TEMPORAL_TARGET_INDEX=335]. -// -// Per `pearl_controller_anchors_isv_driven.md`: even the temporal target -// is signal-driven — the Wiener-EMA tracks the running epoch_idx so the -// consumer's `epoch_idx / temporal_target` ratio stays in a useful range -// regardless of fold length or maximum epoch count. Per Risk 3 in the SP9 -// spec: this axis is the safety net — if statistical and behavioral both -// stay dormant (e.g. fold 0 has no validation samples yet), temporal alone -// eventually fires once epoch_idx_in_fold reaches its EMA-target. -// -// Single-block, single-thread cold-path kernel. No atomicAdd; no shared -// memory; `__threadfence_system()` after the write. - -#include - -extern "C" __global__ void kelly_temporal_target_ema_compute( - const float* __restrict__ isv_signals, - int epoch_idx_isv_index, // EPOCH_IDX_INDEX = 39 - float* __restrict__ scratch_buf, - int scratch_idx -) { - if (blockIdx.x != 0 || threadIdx.x != 0) return; - - float observed = isv_signals[epoch_idx_isv_index]; - if (observed < 0.0f) observed = 0.0f; - - scratch_buf[scratch_idx] = observed; - __threadfence_system(); -} diff --git a/crates/ml/src/cuda_pipeline/kelly_warmup_floor_compute_kernel.cu b/crates/ml/src/cuda_pipeline/kelly_warmup_floor_compute_kernel.cu index 13d35089f..3f6861414 100644 --- a/crates/ml/src/cuda_pipeline/kelly_warmup_floor_compute_kernel.cu +++ b/crates/ml/src/cuda_pipeline/kelly_warmup_floor_compute_kernel.cu @@ -24,11 +24,11 @@ // - q_var_mag ← ISV[FLATNESS_BASE+1=207] (Pearl 2 magnitude flatness) // - q_var_mag_ema ← ISV[Q_VAR_MAG_EMA_INDEX=331] (SP9 EMA-tracked baseline) // - kelly_sample_count ← ISV[KELLY_SAMPLE_COUNT_INDEX=283] (Pearl 6 cumulative) -// - sample_count_target ← ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] (SP9 EMA target) +// - sample_count_target ← ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] (Fix 37.1 Invariant-1 anchor = 100.0) // - intent_eval_divergence← ISV[INTENT_EVAL_DIVERGENCE_INDEX=332] (SP9 producer) -// - divergence_target ← ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] (SP9 EMA target) +// - divergence_target ← ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] (Fix 37.1 Invariant-1 anchor = 2.0) // - epoch_idx ← ISV[EPOCH_IDX_INDEX=39] (existing per-fold counter) -// - temporal_target ← ISV[KELLY_TEMPORAL_TARGET_INDEX=335] (SP9 EMA target) +// - temporal_target ← ISV[KELLY_TEMPORAL_TARGET_INDEX=335] (Fix 37.1 Invariant-1 anchor = 5.0) // // Writes: 1 float to `scratch_buf[scratch_idx]`. Downstream // `apply_pearls_ad_kernel` smooths into ISV[KELLY_WARMUP_FLOOR_INDEX=330]. diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 2c07628b0..6db1d18cd 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -720,17 +720,17 @@ impl StateResetRegistry { RegistryEntry { name: "sp9_kelly_sample_count_target", category: ResetCategory::FoldReset, - description: "ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] — SP9 Fix 37 EMA-tracked sample-sufficiency target. Statistical-axis denominator of the OR'd cold-start exit. Produced by `kelly_sample_count_target_ema_kernel` from `ISV[KELLY_SAMPLE_COUNT_INDEX=283]` (Pearl 6's cumulative-via-max trade count); smoothed by Pearl D Wiener-α. The consumer's `kelly_sample_count / sample_count_target` ratio yields a self-normalised `statistical_conf` in [0, 1]. FoldReset sentinel 0 — Pearl A bootstraps from the new fold's first observation.", + description: "ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] — SP9 Fix 37.1 (2026-05-03) Invariant-1 numerical anchor: minimum trade samples (100.0) for statistical-confidence release of the OR'd cold-start exit. Originally Fix 37 routed this through `kelly_sample_count_target_ema_kernel` + Pearl A/D smoothing, but smoke-test-wrwkz revealed sentinel-bootstrap saturated the EMA to the first observation in epoch 1, making `current/target = 1.0` and defeating the cold-start floor. Now constructor-written ONCE per `feedback_isv_for_adaptive_bounds.md`. FoldReset rewrites the same Invariant-1 anchor (NOT sentinel 0) since this is a constant, not a stateful EMA. The consumer's `kelly_sample_count / sample_count_target` ratio yields a self-normalised `statistical_conf` in [0, 1] that rises monotonically as samples accumulate.", }, RegistryEntry { name: "sp9_kelly_divergence_target", category: ResetCategory::FoldReset, - description: "ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] — SP9 Fix 37 EMA-tracked behavioral-gap target. Behavioral-axis denominator of the OR'd cold-start exit. Produced by `kelly_divergence_target_ema_kernel` from `ISV[INTENT_EVAL_DIVERGENCE_INDEX=332]`; smoothed by Pearl D Wiener-α. The consumer's `1 − divergence / divergence_target` ratio yields a self-normalised `behavioral_conf` in [0, 1]. FoldReset sentinel 0 — Pearl A bootstraps from the new fold's first observation.", + description: "ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] — SP9 Fix 37.1 (2026-05-03) Invariant-1 numerical anchor: divergence ratio (2.0) above which intent-vs-eval is considered diverged enough to keep the cap closed (behavioral-axis denominator of the OR'd cold-start exit). Originally Fix 37 routed this through `kelly_divergence_target_ema_kernel` + Pearl A/D smoothing, but smoke-test-wrwkz revealed sentinel-bootstrap saturated the EMA to the first observation in epoch 1 (`div_tgt=70 005`), making the consumer's ratio collapse to 1.0 and defeating the cold-start floor. Now constructor-written ONCE per `feedback_isv_for_adaptive_bounds.md`. FoldReset rewrites the same Invariant-1 anchor (NOT sentinel 0). The consumer's `1 − divergence / divergence_target` ratio yields a self-normalised `behavioral_conf` in [0, 1] that drops sharply when divergence exceeds the threshold.", }, RegistryEntry { name: "sp9_kelly_temporal_target", category: ResetCategory::FoldReset, - description: "ISV[KELLY_TEMPORAL_TARGET_INDEX=335] — SP9 Fix 37 EMA-tracked temporal target. Temporal-axis denominator of the OR'd cold-start exit; the safety net per Risk 3 in the SP9 spec (always fires once epoch_idx ≥ target if statistical and behavioral both stay dormant). Produced by `kelly_temporal_target_ema_kernel` from `ISV[EPOCH_IDX_INDEX=39]`; smoothed by Pearl D Wiener-α. FoldReset sentinel 0 — Pearl A bootstraps from the new fold's first observation.", + description: "ISV[KELLY_TEMPORAL_TARGET_INDEX=335] — SP9 Fix 37.1 (2026-05-03) Invariant-1 numerical anchor: minimum epochs in fold (5.0) before temporal-confidence release of the OR'd cold-start exit (the safety net per Risk 3 in the SP9 spec — always fires once epoch_idx ≥ target if statistical and behavioral both stay dormant). Originally Fix 37 routed this through `kelly_temporal_target_ema_kernel` + Pearl A/D smoothing, but smoke-test-wrwkz revealed sentinel-bootstrap saturated the EMA to the first observation in epoch 1 (`temp_tgt=1.0`), making temporal_conf trivially `epoch_idx/1.0 ≥ 1` from epoch 1 onwards. Now constructor-written ONCE per `feedback_isv_for_adaptive_bounds.md`. FoldReset rewrites the same Invariant-1 anchor (NOT sentinel 0). The consumer's `epoch_idx_in_fold / temporal_target` ratio yields a self-normalised `temporal_conf` in [0, 1] that rises linearly with fold-time.", }, RegistryEntry { name: "sp9_eval_dist_q", diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index ac98315cf..58819a831 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -6896,21 +6896,38 @@ impl DQNTrainer { } } "sp9_kelly_sample_count_target" => { + // SP9 Fix 37.1 (2026-05-03): Invariant-1 numerical anchor — + // not a stateful EMA. Rewrite the constructor's value at + // fold boundary so the new fold sees the correct + // sample-sufficiency threshold (100 trades) for the + // statistical-confidence axis. Per + // `feedback_isv_for_adaptive_bounds.md`. if let Some(ref fused) = self.fused_ctx { use crate::cuda_pipeline::sp5_isv_slots::KELLY_SAMPLE_COUNT_TARGET_INDEX; - fused.trainer().write_isv_signal_at(KELLY_SAMPLE_COUNT_TARGET_INDEX, 0.0); + fused.trainer().write_isv_signal_at(KELLY_SAMPLE_COUNT_TARGET_INDEX, 100.0); } } "sp9_kelly_divergence_target" => { + // SP9 Fix 37.1 (2026-05-03): Invariant-1 numerical anchor — + // not a stateful EMA. Rewrite the constructor's value at + // fold boundary so the new fold sees the correct + // divergence threshold (2.0) for the behavioral-confidence + // axis. Per `feedback_isv_for_adaptive_bounds.md`. if let Some(ref fused) = self.fused_ctx { use crate::cuda_pipeline::sp5_isv_slots::KELLY_DIVERGENCE_TARGET_INDEX; - fused.trainer().write_isv_signal_at(KELLY_DIVERGENCE_TARGET_INDEX, 0.0); + fused.trainer().write_isv_signal_at(KELLY_DIVERGENCE_TARGET_INDEX, 2.0); } } "sp9_kelly_temporal_target" => { + // SP9 Fix 37.1 (2026-05-03): Invariant-1 numerical anchor — + // not a stateful EMA. Rewrite the constructor's value at + // fold boundary so the new fold sees the correct + // fold-temporal threshold (5 epochs) for the + // temporal-confidence axis. Per + // `feedback_isv_for_adaptive_bounds.md`. if let Some(ref fused) = self.fused_ctx { use crate::cuda_pipeline::sp5_isv_slots::KELLY_TEMPORAL_TARGET_INDEX; - fused.trainer().write_isv_signal_at(KELLY_TEMPORAL_TARGET_INDEX, 0.0); + fused.trainer().write_isv_signal_at(KELLY_TEMPORAL_TARGET_INDEX, 5.0); } } "sp9_eval_dist_q" => { diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 751da628f..7b25f094f 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -4679,3 +4679,170 @@ fmaxf chain. Per-fold persistence vs cross-fold persistence is the existing carve-out — Pearl 6's KELLY_F_SMOOTH stays cross-fold; SP9's warmup_floor is per-fold so each fold's cold-start protection re-engages cleanly. + +## Fix 37.1 — SP9 Kelly cold-start follow-up: targets as Invariant-1 anchors + divergence sentinel handling (2026-05-03) + +**Symptom:** smoke-test-wrwkz (5-epoch L40S smoke on Fix 37 commit +`48a8b9ee7`) ep1 HEALTH_DIAG line: + +``` +sp9_kelly_warmup [floor=0.0000 divergence=70005.93 q_var_mag_ema=6.85e-2 + conf [stat=1.000 bhv=0.333 tmp=1.000] + targets [stat_count_tgt=419 div_tgt=105002 temp_tgt=1.0]] +``` + +`floor=0` and `conf=1` from epoch 1 onward — the cold-start mechanism is +effectively never engaged. + +**Root cause (two distinct quirks):** + +*Quirk 1 — EMA target saturation:* The 3 target ISV slots +(`KELLY_SAMPLE_COUNT_TARGET_INDEX=333`, `KELLY_DIVERGENCE_TARGET_INDEX=334`, +`KELLY_TEMPORAL_TARGET_INDEX=335`) were originally Fix 37 routed through +`kelly_*_target_ema_kernel.cu` + `apply_pearls_ad_kernel`. Pearl A +sentinel-bootstrap on the EMA path causes the FIRST observation to +replace the sentinel directly (per +`pearl_first_observation_bootstrap.md`). So `target = first_obs`, ratio +`current/target = 1.0`, confidence `clamp(ratio, 0, 1) = 1.0`, +`floor = base × (1 − 1.0) = 0`. The cold-start mechanism is defeated +because the EMA target tracks the running observation rather than a +fixed threshold. + +The principled fix per `feedback_isv_for_adaptive_bounds.md`: +"sufficient" is defined in absolute terms via Invariant-1 numerical +anchors. The 3 target slots become constructor-written constants +(written ONCE in the trainer constructor, similar to existing +`config.cql_alpha → ISV[CQL_ALPHA_INDEX=48]`), and the target-relative +ratio `current_observation / fixed_anchor` then yields a self-normalised +confidence in [0, 1] that rises monotonically as samples / +divergence-stability / fold-time accumulate. + + - `ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] = 100.0f` — minimum trade + samples for statistical confidence to release the cap + - `ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] = 2.0f` — divergence ratio + above which intent-vs-eval is considered diverged enough to keep the + cap closed + - `ISV[KELLY_TEMPORAL_TARGET_INDEX=335] = 5.0f` — minimum epochs in + fold for temporal confidence to release the cap (safety net per + Risk 3 in spec) + +*Quirk 2 — divergence ratio explosion:* The +`intent_eval_divergence_compute_kernel.cu` previously computed +`divergence = max(intent_f, EPS_DIV) / max(eval_f, EPS_DIV)`. Before the +first val-window populates `ISV[EVAL_DIST_F_INDEX=338]`, the slot is at +Pearl A sentinel 0. With `intent_f ≈ 0.07` and `eval_f` at sentinel, +flooring at `EPS_DIV=1e-6` produced a synthetic +`divergence ≈ 0.07 / 1e-6 ≈ 70 000` (matched in smoke-test-wrwkz +exactly). Algebraically the warmup-floor kernel still derived +`behavioral_conf = clamp(1 − 70 000/2, 0, 1) = 0` (correct semantic +"no behavioral signal yet"), but the synthetic 70 000 in HEALTH_DIAG +masked the actual signal-flow and was easily misread as a runaway. + +The principled fix: explicitly detect the sentinel via threshold +`SENTINEL_THRESHOLD = 1e-5f` (well below any realisable `eval_f` after +the Pearl A first-observation replacement) and emit +`SENTINEL_DIVERGENCE = 1e6f`. Same `behavioral_conf = 0` outcome but +with explicit provenance — HEALTH_DIAG now reads `divergence=1e6` until +the first val window populates `eval_f`, at which point the real ratio +takes over. + +**Fix structure (atomic commit per `feedback_no_partial_refactor`):** + +- DELETE 3 EMA target updater kernels (`.cu` files): + - `kelly_sample_count_target_ema_kernel.cu` + - `kelly_divergence_target_ema_kernel.cu` + - `kelly_temporal_target_ema_kernel.cu` +- DELETE 3 entries from `crates/ml/build.rs::kernels_with_common` +- DELETE 3 cubin static byte arrays + 3 struct fields + 3 cubin loaders + + 3 fields in trainer construction tuple in `gpu_dqn_trainer.rs` +- DELETE 3 launches in `launch_sp9_kelly_warmup_floor` (chain shrinks + from 5 → 2 launches: q_var_mag_ema + main warmup-floor kernel) +- DELETE 3 scratch slots (`SCRATCH_SP9_SAMPLE_COUNT_TARGET=262`, + `SCRATCH_SP9_DIVERGENCE_TARGET=263`, `SCRATCH_SP9_TEMPORAL_TARGET=264`); + `SP5_SCRATCH_TOTAL` 268 → 265; `SCRATCH_SP9_EVAL_DIST_BASE` slides + 265 → 262 (3 floats reclaimed) +- ADD constructor-write of 3 Invariant-1 anchors (100.0, 2.0, 5.0) at + `gpu_dqn_trainer.rs` constructor-time, immediately before the layout- + fingerprint write — same pattern as existing + `*sig_ptr.add(CQL_ALPHA_INDEX) = config.cql_alpha` +- UPDATE 3 dispatch arms in `reset_named_state` to rewrite the same + Invariant-1 anchors at fold boundary (NOT sentinel 0) — these are + constants, not stateful EMAs; they must never reach 0 between folds +- UPDATE 3 registry descriptions in `state_reset_registry.rs` to reflect + the Invariant-1 anchor semantic + the smoke-test-wrwkz evidence +- UPDATE `intent_eval_divergence_compute_kernel.cu` with + `SENTINEL_THRESHOLD=1e-5f` / `SENTINEL_DIVERGENCE=1e6f` sentinel-detect + branch +- ISV slot layout UNCHANGED — the 3 target slots @ ISV[333..336) + remain. Wiener-buffer linear span (`SP5_PRODUCER_COUNT=165`) UNCHANGED + — the 3 wiener triples for the deleted producers become "reserved- + unused" similar to Pearl 6's [525..543) carve-out (no + `launch_apply_pearls` consumer fires for those slots, so the wiener + state stays zero-initialised) + +**Files touched:** + +- `crates/ml/src/cuda_pipeline/kelly_sample_count_target_ema_kernel.cu` (DELETED) +- `crates/ml/src/cuda_pipeline/kelly_divergence_target_ema_kernel.cu` (DELETED) +- `crates/ml/src/cuda_pipeline/kelly_temporal_target_ema_kernel.cu` (DELETED) +- `crates/ml/build.rs` — 3 entries removed; comment updated to "Four + producer kernels" +- `crates/ml/src/cuda_pipeline/intent_eval_divergence_compute_kernel.cu` — + sentinel-detect branch for uninitialised `eval_f` +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 3 cubin statics + + 3 struct fields + 3 cubin loaders + 3 tuple entries + 3 launches + removed; constructor adds 3 Invariant-1 anchor writes; + `SP5_SCRATCH_TOTAL` 268 → 265; `SCRATCH_SP9_EVAL_DIST_BASE` 265 → 262; + 3 obsolete `SCRATCH_SP9_*_TARGET` constants removed +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — 3 descriptions + updated to reflect Invariant-1 anchor semantic +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — 3 dispatch + arms updated to write 100.0 / 2.0 / 5.0 (NOT 0.0) +- `docs/dqn-wire-up-audit.md` (this entry) + +**Verification:** + +- `SQLX_OFFLINE=true cargo check -p ml` — clean (only pre-existing 18 + warnings; no new errors or warnings). +- `SQLX_OFFLINE=true cargo test -p ml --lib state_reset` — all 4 tests + pass (contract test `every_fold_and_soft_reset_entry_has_dispatch_arm` + still covers all 9 SP9 entries; their dispatch arms now write + Invariant-1 anchors instead of 0). +- `SQLX_OFFLINE=true cargo test -p ml --lib sp5_isv_slots` — all 9 + tests pass (slot layout unchanged; producer count unchanged). + +**Expected smoke-test signature post-fix (5-epoch L40S):** + +- `floor`: starts at base_floor × 1.0 in epoch 1 (Pearl D Wiener-α + smoothing of base_floor × (1 − combined_confidence) where confidence + is genuinely low), decays as confidence accumulates over epochs +- `divergence`: `1e6` (SENTINEL_DIVERGENCE) until the first val window + populates `ISV[EVAL_DIST_F_INDEX=338]`; small (≤ 5) afterward +- `conf [stat=X bhv=Y tmp=Z]`: `stat = sample_count/100` rises with + trade count; `bhv` stays 0 until eval_dist matures; `tmp = epoch_idx/5` + rises linearly with fold-time +- combined_conf reaches 1.0 only when at least one axis genuinely + matures — typically temporal axis at epoch 5 in fold 0 + +**Per-pearl provenance:** + +- `feedback_isv_for_adaptive_bounds.md` — Invariant-1 numerical anchors + are the survivable case for hardcoded constants (100.0, 2.0, 5.0 + define what "sufficient" means in absolute terms). +- `pearl_controller_anchors_isv_driven.md` — caveat: not every + anchor / target needs to be EMA-tracked. Cold-start exit thresholds + are the canonical case where a fixed threshold is the principled + design (matching test for "regime change") — the EMA anti-pattern + (target = current_obs) was a misapplication. +- `pearl_first_observation_bootstrap.md` — sentinel-detect branch in + divergence kernel makes "before first observation" explicit; same + semantic as Pearl A but with cleaner provenance in HEALTH_DIAG. +- `feedback_no_partial_refactor.md` — atomic commit (kernel deletions + + Rust deletions + constructor-writes + dispatch-arm updates + audit + entry, all in one commit). +- `feedback_no_stubs.md` — Invariant-1 anchors are real values defining + "sufficient" in absolute terms (100 trades / 2.0 ratio / 5 epochs), + not stubs. +- `feedback_no_cpu_compute_strict.md` — constructor-writes are GPU-side + ISV writes via mapped-pinned `sig_ptr`; same path as all existing + constructor-time ISV initialization (CQL_ALPHA, GAMMA_DIR, etc.).