diff --git a/crates/ml/build.rs b/crates/ml/build.rs index c7097b07b..55dd3f753 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -939,30 +939,17 @@ fn main() { // `feedback_no_partial_refactor.md` (matching Phase 1.1-1.3 + // 3.1-3.3 precedent). "regret_signal_kernel.cu", - // SP15 Phase 3.5 (2026-05-06): confidence-aware Hold floor — - // bounded sigmoid `hold_floor = α × σ(k × (entropy − ε₀))`. - // Single source file with one `extern "C" __global__` symbol - // (`hold_floor_kernel`) → one cubin → one launcher - // (`launch_sp15_hold_floor`). Mirrors the Phase 3.3 - // `dd_penalty_kernel.cu` / 3.4 `regret_signal_kernel.cu` - // single-kernel-per-cubin pattern. Action-selection level - // (NOT reward shaping) — the follow-up consumer adds - // `hold_floor` to Q_hold pre-argmax/Thompson selection so - // Hold becomes uncertainty expression rather than the - // distributional default. α / k / ε₀ are ISV-tracked - // (slots 426 / 427 / 428); the follow-up signal-driven - // producers updating these from rolling 95th percentile of - // |Q_dir| (NOT running max — outlier-ratchet vulnerable per - // spec second-review #6), running variance of entropy, and - // 75th percentile of running entropy distribution are - // documented Phase 3.5 follow-up sub-tasks. Phase 3.5 lands - // kernel + launcher + 4 ISV anchor seeds + 4 registry - // entries + 4 dispatch arms only; the action-selection - // wiring (add `hold_floor` to Q_hold pre-argmax/Thompson) - // is deferred to a follow-up atomic commit per - // `feedback_no_partial_refactor.md` (matching Phase 1.1-1.3 - // + 3.1-3.4 precedent). - "hold_floor_kernel.cu", + // SP15 Phase 3.5 — `hold_floor_kernel.cu` REMOVED in Phase 3.5.b + // (2026-05-06): the bounded-sigmoid `hold_floor = α × σ(k × + // (entropy − ε₀))` is now an inline `__device__` computation + // inside `experience_action_select` reading the same ISV slots + // 426/427/428/429 directly. The standalone kernel was scaffolding + // for the deferred consumer per `feedback_no_partial_refactor`; + // when the consumer landed inline, the kernel + launcher + + // cubin manifest entry became dead code per + // `feedback_wire_everything_up.md` + `feedback_no_legacy_aliases.md`. + // ISV slots + state_reset_registry entries remain (still + // consumed by the inline computation). // SP15 Phase 3.5.2 (2026-05-06): asymmetric reward under DD — // for gains `r_adjusted = r × (1 + λ × dd_pct)`, for losses // unchanged. Applied BEFORE the SP12 v3 NEG/POS reward cap diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 71d98bcce..26e4b2832 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1254,6 +1254,88 @@ extern "C" __global__ void experience_action_select( e_dir[d] = compute_e_c51_inline(probs_d, atom_vals_d, n_atoms); } + /* SP15 Phase 3.5.b + 3.5.3.b (2026-05-06) — confidence-aware Hold + * floor + cooldown mask, both inline. + * + * 3.5.b — confidence-aware Hold floor: + * `hold_floor = α × σ(k × (entropy − ε₀))` per spec §8.2 (3.5) + * post-amendment-2 fix. Action-selection level (NOT reward shaping). + * When the policy is uncertain (high entropy across the direction + * posterior) the Hold component receives a positive floor that + * makes Hold the principled "uncertainty expression" rather than + * the distributional default. α / k / ε₀ are ISV-tracked + * (slots 426 / 427 / 428). + * + * `entropy` here is the Shannon entropy of the per-step direction + * policy `p[d] = softmax(e_dir[d])` — the same quantity the + * downstream "75th-percentile entropy distribution" producer (Phase + * 3.5 follow-up) tracks via `ENTROPY_DIST_REF=429`. Computed once + * inline from the b0_size = 4 e_dir floats already in registers; + * no extra kernel launch + scratch buffer. The standalone + * `hold_floor_kernel.cu` + `launch_sp15_hold_floor` was scaffolding + * for the deferred consumer wiring; with the consumer landing inline + * the kernel + launcher + cubin are deleted in the same commit per + * `feedback_no_legacy_aliases` + `feedback_wire_everything_up`. + * + * 3.5.3.b — cooldown mask: + * When `ISV[COOLDOWN_BARS_REMAINING=435] > 0`, force the picked + * direction to Hold by setting the Thompson-effective Q for every + * non-Hold direction to -INFINITY before argmax. The cooldown + * counter producer (`cooldown_kernel`) remains intact — only the + * action-selection consumer is wired here. Cooldown supersedes + * hold_floor — when forcing Hold the floor is moot. + * + * Implementation: build a local q_eff_dir[4] = e_dir[] copy with + * hold_floor added at DIR_HOLD, then optionally mask non-Hold to + * -INFINITY when cooldown > 0. Pass 2 below consumes q_eff_dir + * instead of e_dir for the temperature blend. e_dir itself is + * preserved for the conviction + q_gap consumers downstream — those + * read the raw E[Q] posterior independent of the gating signals + * (per spec the conviction signal must reflect the policy's + * confidence in WHERE to trade, not in whether to gate). */ + float q_eff_dir[4]; + float hold_floor = 0.0f; + int cooldown_active = 0; + if (isv_signals_ptr != NULL && b0_size <= 4) { + /* Direction-policy entropy. Stable softmax over e_dir[0..b0_size). */ + float ed_max = e_dir[0]; + for (int d = 1; d < b0_size; d++) { + if (e_dir[d] > ed_max) ed_max = e_dir[d]; + } + float ed_sum = 0.0f; + float ed_exps[4]; + for (int d = 0; d < b0_size; d++) { + ed_exps[d] = expf(e_dir[d] - ed_max); + ed_sum += ed_exps[d]; + } + float inv_ed_sum = (ed_sum > 1e-12f) ? (1.0f / ed_sum) : 0.0f; + float entropy_actual = 0.0f; + for (int d = 0; d < b0_size; d++) { + float p = ed_exps[d] * inv_ed_sum; + if (p > 1e-12f) entropy_actual -= p * logf(p); + } + + const float alpha = isv_signals_ptr[/*HOLD_FLOOR_ALPHA_INDEX*/ 426]; + const float k = isv_signals_ptr[/*HOLD_FLOOR_K_INDEX*/ 427]; + const float eps0 = isv_signals_ptr[/*HOLD_FLOOR_EPS0_INDEX*/ 428]; + const float x_arg = k * (entropy_actual - eps0); + const float sigmoid = 1.0f / (1.0f + expf(-x_arg)); + hold_floor = alpha * sigmoid; + + /* Cooldown counter — ISV[COOLDOWN_BARS_REMAINING=435]. f32 + * representation of an integer counter; > 0 means cooldown + * still active. */ + cooldown_active = (isv_signals_ptr[/*COOLDOWN_BARS_REMAINING_INDEX*/ 435] > 0.0f); + } + for (int d = 0; d < b0_size; d++) { + q_eff_dir[d] = e_dir[d]; + } + /* Hold-floor application (no-op when cooldown supersedes below). + * DIR_HOLD = 1 by `state_layout.cuh`. */ + if (DIR_HOLD < b0_size) { + q_eff_dir[DIR_HOLD] += hold_floor; + } + /* SP10 (Fix 38, 2026-05-03): unconditional temperature-blended * Thompson selector. Read ISV-driven temperature; fall back to 1.0 * (pure Thompson) when isv_signals_ptr is unavailable (test paths, @@ -1271,11 +1353,34 @@ extern "C" __global__ void experience_action_select( if (!(thompson_temp > 0.5f)) thompson_temp = 0.5f; if (thompson_temp > 2.0f) thompson_temp = 2.0f; - { + if (cooldown_active && DIR_HOLD < b0_size) { + /* SP15 Phase 3.5.3.b — cooldown mask: when + * `ISV[COOLDOWN_BARS_REMAINING=435] > 0`, force the picked + * direction to Hold per spec §9.2 (3.5.3) post-amendment-2 fix. + * Hard short-circuit (skip Pass 2 entirely) — sidesteps the + * temperature-blend numerics where a finite-sentinel-on-non-Hold + * approach would still let a pure-Thompson (τ=1) sample dominate + * the masked direction. + * + * Also skip the rng counter advance Pass 2 would have performed — + * deterministic per-step Philox sub-stream advance is bounded + * by the contiguous mag/ord/urg path below; the direction-stream + * draws Pass 2 would have performed are unused here, which only + * shifts the mag/ord/urg streams' state by a constant amount per + * cooldown step. Reproducibility-at-fixed-seed is preserved + * because the cooldown gate is a deterministic function of the + * ISV state at this thread's bar. */ + dir_idx = DIR_HOLD; + } else { /* Pass 2 recomputes per-direction probs/atoms; the alternative — * caching across directions — would cost b0_size * 2 * MAX_ATOMS * stack floats (~4 KiB per thread) and risk local-memory spill. - * Recompute is < 1 µs per thread at n_atoms=51. */ + * Recompute is < 1 µs per thread at n_atoms=51. + * + * Uses q_eff_dir[d] (= e_dir[d] + hold_floor at DIR_HOLD) + * instead of raw e_dir[d] so the SP15 3.5.b hold_floor signal + * shapes both the deterministic centre of the temperature blend + * AND the argmax fallback at τ=0. */ float best_q_eff = -1e30f; int best_d = 0; for (int d = 0; d < b0_size; d++) { @@ -1287,7 +1392,7 @@ extern "C" __global__ void experience_action_select( /* Temperature blend: τ=0 → argmax of E[Q]; τ=1 → pure * Thompson; τ>1 → exaggerated exploration. */ - float q_eff = e_dir[d] + thompson_temp * (q_sample - e_dir[d]); + float q_eff = q_eff_dir[d] + thompson_temp * (q_sample - q_eff_dir[d]); if (q_eff > best_q_eff) { best_q_eff = q_eff; best_d = d; } } dir_idx = best_d; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index b66516b3e..cb234e15e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1490,79 +1490,18 @@ pub fn launch_sp15_regret_signal( Ok(()) } -/// SP15 Phase 3.5 (2026-05-06): cubin for the `hold_floor_kernel`. -/// -/// Per spec §8.2 (3.5) post-amendment-2 fix: confidence-aware Hold floor -/// `hold_floor = α × σ(k × (entropy − ε₀))` — action-selection level -/// (NOT reward shaping). The follow-up consumer adds `hold_floor` to the -/// Hold component of the per-action Q vector pre-argmax/Thompson -/// selection so Hold becomes uncertainty expression rather than the -/// distributional default. α / k / ε₀ are ISV-tracked (slots 426 / 427 / -/// 428); the signal-driven producers updating these from the rolling -/// 95th percentile of |Q_dir| (NOT running max — outlier-ratchet -/// vulnerable per spec second-review #6), the running variance of -/// entropy, and the 75th percentile of the running entropy distribution -/// (slot 429 = `ENTROPY_DIST_REF_INDEX`) are documented Phase 3.5 -/// follow-up sub-tasks. -/// -/// Phase 3.5 lands kernel + launcher + 4 ISV constructor seeds + 4 -/// state-reset registry entries + 4 dispatch arms only; the -/// action-selection wiring (add `hold_floor` to Q_hold pre-argmax / -/// Thompson) is deferred to a follow-up atomic commit per -/// `feedback_no_partial_refactor.md` (Phase 1.1-1.3 + 3.1-3.4 -/// precedent). -pub static SP15_HOLD_FLOOR_CUBIN: &[u8] = - include_bytes!(concat!(env!("OUT_DIR"), "/hold_floor_kernel.cubin")); - -/// SP15 Phase 3.5 (2026-05-06): launcher for the confidence-aware Hold -/// floor kernel. Free function (matches `launch_sp15_dd_penalty` / -/// `launch_sp15_regret_signal` precedent — single kernel, single cubin, -/// single launcher) so unit/oracle tests can drive the kernel directly -/// without the trainer struct; production callers invoke the same -/// launcher. -/// -/// `entropy` is the policy's per-step action entropy (nats). `isv` MUST -/// be the ISV bus (≥ `SP15_SLOT_END=443` f32 slots — slots 426 / 427 / -/// 428 are read-only; the kernel does NOT update them in this commit, -/// the signal-driven producers updating them are documented Phase 3.5 -/// follow-up sub-tasks). `floor_out` MUST point to at least 1 writable -/// f32 slot; the kernel writes the bounded sigmoid output `α × σ(k × -/// (entropy − ε₀))`. Single thread, single block (mirrors -/// `dd_penalty_kernel` / `regret_signal_kernel` — no reduction). -pub fn launch_sp15_hold_floor( - stream: &Arc, - entropy: f32, - isv: cudarc::driver::sys::CUdeviceptr, - floor_out: cudarc::driver::sys::CUdeviceptr, -) -> Result<(), MLError> { - let module = stream - .context() - .load_cubin(SP15_HOLD_FLOOR_CUBIN.to_vec()) - .map_err(|e| MLError::ModelError(format!( - "load sp15_hold_floor cubin: {e}" - )))?; - let kernel = module - .load_function("hold_floor_kernel") - .map_err(|e| MLError::ModelError(format!( - "load hold_floor_kernel function: {e}" - )))?; - unsafe { - stream - .launch_builder(&kernel) - .arg(&entropy) - .arg(&isv) - .arg(&floor_out) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), - shared_mem_bytes: 0, - }) - .map_err(|e| MLError::ModelError(format!( - "launch hold_floor_kernel: {e}" - )))?; - } - Ok(()) -} +// SP15 Phase 3.5.b (2026-05-06): hold_floor moved INLINE into +// `experience_action_select` — the standalone `hold_floor_kernel.cu` + +// `SP15_HOLD_FLOOR_CUBIN` + `launch_sp15_hold_floor` were Phase 3.5 +// scaffolding for a deferred consumer. Phase 3.5.b wires the consumer +// directly inside the production action-selection kernel reading the +// same ISV slots (426 / 427 / 428 / 429) and computing the same +// `hold_floor = α × σ(k × (entropy − ε₀))` as a single-line `__device__` +// expression on per-thread direction-policy entropy. Launching a kernel +// to write one f32 just to read it back was unnecessary; deletion is +// the correct application of `feedback_wire_everything_up.md` + +// `feedback_no_legacy_aliases.md`. ISV slots + state_reset_registry +// entries + dispatch arms remain — only the launch path is removed. /// SP15 Phase 3.5.2 (2026-05-06): cubin for the /// `dd_asymmetric_reward_kernel`. diff --git a/crates/ml/src/cuda_pipeline/hold_floor_kernel.cu b/crates/ml/src/cuda_pipeline/hold_floor_kernel.cu deleted file mode 100644 index 8979cefc5..000000000 --- a/crates/ml/src/cuda_pipeline/hold_floor_kernel.cu +++ /dev/null @@ -1,65 +0,0 @@ -// crates/ml/src/cuda_pipeline/hold_floor_kernel.cu -// -// SP15 Phase 3.5 — confidence-aware Hold floor (bounded sigmoid). -// -// hold_floor = α × σ(k × (entropy − ε₀)) -// -// Action-selection level (NOT reward shaping). The follow-up consumer -// adds `hold_floor` to the Hold component of the per-action Q vector -// pre-argmax/Thompson selection — Hold becomes uncertainty expression -// rather than the distributional default. α / k / ε₀ are ISV-tracked -// (slots 426 / 427 / 428). -// -// **Initial sentinel values** (Phase 3.5 lands the primitive only; the -// signal-driven producers updating these from a rolling 95th percentile -// of |Q_dir| (NOT running max — outlier-ratchet vulnerable per spec -// second-review #6), the running variance of entropy, and the 75th -// percentile of the running entropy distribution are documented Phase -// 3.5 follow-up sub-tasks): -// -// α = HOLD_FLOOR_ALPHA = 0.5 (sentinel — half a Q-unit at peak σ) -// k = HOLD_FLOOR_K = 10.0 (steep sigmoid — fast saturation) -// ε₀ = HOLD_FLOOR_EPS0 = 1.0 (entropy threshold — ~1 nat ≈ ln(e)) -// ENTROPY_DIST_REF (slot 429) = 0.0 (placeholder for the 75th-percentile -// producer; not consumed by this kernel — it tracks the ε₀ producer's -// running reference, landed in the follow-up sub-task). -// -// Hard rules: -// -// * `feedback_no_atomicadd` — single-thread, single-block, pure scalar -// arithmetic; no reductions, no `atomicAdd`. -// * `feedback_no_partial_refactor` — kernel + launcher land first; -// consumer wiring (action selection adds `hold_floor` to Q_hold -// pre-Thompson/argmax) deferred to a follow-up atomic commit per -// the established Phase 1.1-1.5 / 3.1-3.4 precedent. -// * `feedback_no_stubs` — kernel returns the actual sigmoid output; -// the oracle test verifies BOTH gate sides (high-entropy ≈ 0.497, -// low-entropy < 0.05). -// * `pearl_bounded_modifier_outputs_require_structural_activation` — -// the bounded sigmoid IS the structural activation; the ISV-tracked -// α anchors the bound without a downstream runtime clamp. -// -// Single-thread, single-block (mirrors `dd_penalty_kernel` / -// `regret_signal_kernel` — no reduction). - -extern "C" __global__ void hold_floor_kernel( - float entropy, - const float* __restrict__ isv, - float* __restrict__ floor_out -) { - if (threadIdx.x != 0 || blockIdx.x != 0) return; - - // ISV slot indices (mirror dd_penalty_kernel's literal-index pattern; - // matched against sp15_isv_slots.rs by the slot-layout regression - // test sp15_slot_layout_locked). - const float alpha = isv[/*HOLD_FLOOR_ALPHA_INDEX*/ 426]; - const float k = isv[/*HOLD_FLOOR_K_INDEX*/ 427]; - const float eps0 = isv[/*HOLD_FLOOR_EPS0_INDEX*/ 428]; - - const float x = k * (entropy - eps0); - // Standard logistic sigmoid: 1 / (1 + e^-x). - const float sigmoid = 1.0f / (1.0f + expf(-x)); - *floor_out = alpha * sigmoid; - - __threadfence_system(); -} diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 0e4f2cf98..b0b9c3076 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -1020,74 +1020,16 @@ mod gpu { ); } - /// SP15 Phase 3.5 — confidence-aware Hold floor: bounded sigmoid - /// `hold_floor = α × σ(k × (entropy − ε₀))` per spec §8.2 (3.5) - /// post-amendment-2 fix. With α=0.5, k=10, ε₀=1.0: - /// * entropy=1.5 (above ε₀): hold_floor ≈ 0.5 × σ(5) ≈ 0.5 × 0.993 - /// ≈ 0.497 → STRONG Hold bias. - /// * entropy=0.5 (below ε₀): hold_floor ≈ 0.5 × σ(-5) ≈ 0.5 × - /// 0.0067 ≈ 0.003 → near-zero floor. - /// Validates the bounded-sigmoid math AND the structural property - /// that low-entropy (high-confidence) states do not get a Hold - /// distributional default. - #[test] - #[ignore = "requires GPU"] - fn hold_floor_bounded_sigmoid() { - use ml::cuda_pipeline::sp15_isv_slots::{ - HOLD_FLOOR_ALPHA_INDEX, HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX, - }; - - let stream = make_test_stream(); - - let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } - .expect("alloc MappedF32Buffer for isv bus"); - let mut isv_init = vec![0.0f32; ISV_LEN]; - isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5; - isv_init[HOLD_FLOOR_K_INDEX] = 10.0; - isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0; - isv_buf.write_from_slice(&isv_init); - - let out_buf = unsafe { MappedF32Buffer::new(1) } - .expect("alloc MappedF32Buffer for hold_floor out"); - - // High entropy → strong Hold bias. - out_buf.write_from_slice(&[0.0_f32]); - ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_hold_floor( - &stream, - /*entropy*/ 1.5, - isv_buf.dev_ptr, - out_buf.dev_ptr, - ) - .expect("launch hold_floor_kernel (high-entropy)"); - stream - .synchronize() - .expect("synchronize after hold_floor_kernel launch (high-entropy)"); - let high_entropy_floor = out_buf.read_all()[0]; - assert!( - (high_entropy_floor - 0.497).abs() < 0.01, - "hold_floor at entropy=1.5 = {high_entropy_floor}, expected ~0.497" - ); - - // Low entropy → near-zero floor (sentinel non-zero output buffer - // ensures the kernel actually overwrites with the small value - // rather than skipping the store). - out_buf.write_from_slice(&[1.0_f32]); - ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_hold_floor( - &stream, - /*entropy*/ 0.5, - isv_buf.dev_ptr, - out_buf.dev_ptr, - ) - .expect("launch hold_floor_kernel (low-entropy)"); - stream - .synchronize() - .expect("synchronize after hold_floor_kernel launch (low-entropy)"); - let low_entropy_floor = out_buf.read_all()[0]; - assert!( - low_entropy_floor < 0.05, - "hold_floor at entropy=0.5 = {low_entropy_floor}, expected near 0" - ); - } + // SP15 Phase 3.5.b (2026-05-06): the standalone `hold_floor_kernel.cu` + // + `launch_sp15_hold_floor` were deleted when the consumer wiring + // landed inline inside `experience_action_select`. The previous + // `hold_floor_bounded_sigmoid` oracle test that drove the standalone + // kernel directly is removed in the same commit per + // `feedback_no_legacy_aliases.md`. The bounded-sigmoid math is now + // exercised end-to-end by the new + // `action_select_applies_hold_floor_inline` test below + // (full action_select kernel call with deterministic e_dir → known + // entropy → expected Q[Hold] bump). /// SP15 Phase 3.5.2 — gain side: asymmetric reward under DD, /// `r_adjusted = r × (1 + λ × dd_pct)` per spec §9.2 (3.5.2). With @@ -1885,6 +1827,378 @@ mod gpu { raw_sharpe ); } + + // ── SP15 Phase 3.5.b + 3.5.3.b oracle tests ────────────────────────── + // + // These tests exercise the inline hold_floor + cooldown-mask wiring + // inside the production `experience_action_select` kernel. The fixture + // mirrors `distributional_q_tests.rs::ProdActionSelectFixture` (which + // is `pub(crate)` in the trainer's test module and not reachable from + // this integration-test file) — we re-build a minimal launcher here + // to keep the oracle test self-contained. Branch sizes match the + // production 4-direction action space. + + use cudarc::driver::{CudaFunction, CudaSlice, LaunchConfig, PushKernelArg}; + use std::sync::OnceLock; + + const ACTION_SELECT_PROD_EXP_CUBIN: &[u8] = include_bytes!(concat!( + env!("OUT_DIR"), + "/experience_kernels.cubin" + )); + static ACTION_SELECT_KERNEL: OnceLock = OnceLock::new(); + + fn load_action_select(stream: &Arc) -> &'static CudaFunction { + ACTION_SELECT_KERNEL.get_or_init(|| { + let module = stream + .context() + .load_cubin(ACTION_SELECT_PROD_EXP_CUBIN.to_vec()) + .expect("load experience_kernels cubin"); + module + .load_function("experience_action_select") + .expect("load experience_action_select") + }) + } + + /// Branch sizes for the 4-direction production action space. + const AS_B0: i32 = 4; // direction (Short=0, Hold=1, Long=2, Flat=3) + const AS_B1: i32 = 3; // magnitude + const AS_B2: i32 = 3; // order + const AS_B3: i32 = 3; // urgency + const AS_Q_STRIDE: usize = (AS_B0 + AS_B1 + AS_B2 + AS_B3) as usize; + const AS_DIR_HOLD: i32 = 1; + const AS_DIR_LONG: i32 = 2; + + fn decode_dir(action_idx: i32) -> i32 { + action_idx / (AS_B1 * AS_B2 * AS_B3) + } + + /// Build C51 logits that produce a deterministic per-direction E[Q] + /// equal to `target_v[d]`. Uses a high peak logit so softmax is + /// ~delta on the closest atom; with the inverse-CDF Thompson sample + /// the `q_sample` ≈ `e_dir` (jitter < linear-support delta_z), + /// keeping `q_eff = e_dir + temp·(q_sample-e_dir)` tightly anchored + /// at e_dir for any temp ∈ [0.5, 2.0]. + fn fill_peaked_logits_action_select( + batch: usize, + n_atoms: i32, + v_min: f32, + delta_z: f32, + target_v: &[f32; 4], + peak_logit: f32, + ) -> Vec { + let mut out = vec![0.0_f32; batch * (AS_B0 as usize) * (n_atoms as usize)]; + for i in 0..batch { + for (d, tv) in target_v.iter().enumerate() { + let peak_a = (((*tv - v_min) / delta_z).round() as usize) + .min((n_atoms - 1) as usize); + let base = + i * (AS_B0 as usize) * (n_atoms as usize) + d * (n_atoms as usize); + out[base + peak_a] = peak_logit; + } + } + out + } + + fn fill_linear_support_action_select( + batch: usize, v_min: f32, v_max: f32, delta_z: f32, + ) -> Vec { + let mut out = vec![0.0_f32; batch * (AS_B0 as usize) * 3]; + for i in 0..batch { + for d in 0..(AS_B0 as usize) { + let base = i * (AS_B0 as usize) * 3 + d * 3; + out[base] = v_min; + out[base + 1] = v_max; + out[base + 2] = delta_z; + } + } + out + } + + /// Launch `experience_action_select` with deterministic eval-mode + /// inputs (eps_start = eps_end = 0). `isv_dev_ptr` MUST point to a + /// >= ISV_LEN f32 buffer (slot 339 = thompson_temp; slot 426 = α; + /// slot 427 = k; slot 428 = ε₀; slot 435 = cooldown_remaining). + /// Returns the per-sample factored action picks. + fn launch_action_select_for_test( + stream: &Arc, + batch: usize, + n_atoms: i32, + q_values: &CudaSlice, + b_logits_dir: &CudaSlice, + per_sample_support: &CudaSlice, + actions_buf: &mut CudaSlice, + q_gaps_buf: &mut CudaSlice, + intent_buf: &mut CudaSlice, + conv_buf: &mut CudaSlice, + mag_conv_buf: &mut CudaSlice, + isv_dev_ptr: u64, + ) -> Vec { + let kernel = load_action_select(stream); + let blocks = ((batch as u32).div_ceil(256)).max(1); + let cfg = LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }; + let n_i32: i32 = batch as i32; + let null_ptr: u64 = 0; + // Eval mode: cosine schedule degenerate (eps_start = eps_end = 0). + let eps_start: f32 = 0.0; + let eps_end: f32 = 0.0; + unsafe { + stream + .launch_builder(kernel) + .arg(q_values) + .arg(&mut *actions_buf) + .arg(&mut *q_gaps_buf) + .arg(&eps_start) + .arg(&eps_end) + .arg(&0_i32) // current_epoch + .arg(&1_i32) // total_epochs + .arg(&n_i32) + .arg(&AS_B0) + .arg(&AS_B1) + .arg(&AS_B2) + .arg(&AS_B3) + .arg(&0.0_f32) // q_gap_threshold + .arg(&null_ptr) // portfolio_states (NULL) + .arg(&0_i32) // min_hold_bars + .arg(&0.0_f32) // max_position + .arg(&1.0_f32) // eps_exp_mult + .arg(&1.0_f32) // eps_ord_mult + .arg(&1.0_f32) // eps_urg_mult + .arg(&0_i32) // timestep + .arg(&null_ptr) // per_sample_epsilon (NULL → cosine) + .arg(&isv_dev_ptr) // isv_signals_ptr + .arg(&0_i32) // contrarian_active + .arg(&mut *intent_buf) + .arg(&mut *conv_buf) + .arg(&mut *mag_conv_buf) + .arg(b_logits_dir) + .arg(per_sample_support) + .arg(&null_ptr) // atom_positions = NULL → linear support + .arg(&n_atoms) + .launch(cfg) + .expect("launch experience_action_select"); + } + stream.synchronize().expect("sync after action_select"); + stream.clone_dtoh(actions_buf).expect("readback actions") + } + + /// SP15 Phase 3.5.b — confidence-aware Hold floor wired inline. + /// + /// Construct a per-sample direction posterior with a near-uniform + /// shape so Shannon entropy is high (≈ ln(4) ≈ 1.386 nats > ε₀ = 1.0). + /// e_dir = [0.50, 0.45, 0.50, 0.45] yields softmax probs ≈ + /// [0.255, 0.245, 0.255, 0.245] → entropy ≈ 1.386 nats. + /// With α = 0.5, k = 10, ε₀ = 1.0: + /// x = 10 × (1.386 − 1.0) ≈ 3.86, σ(x) ≈ 0.979, hold_floor ≈ 0.490. + /// + /// Without hold_floor: argmax(e_dir) is Short(0) or Long(2) (ties to + /// the lower index → Short wins on the strictly-greater test). With + /// hold_floor: Hold's e_dir 0.45 + 0.49 = 0.94 dominates, so picked + /// dir = Hold. + /// + /// Validates: (a) the inline hold_floor computation; (b) the + /// SP15 3.5.b architectural property that the floor shifts the + /// argmax to Hold under high-entropy uncertainty. + #[test] + #[ignore = "requires GPU"] + fn action_select_applies_hold_floor_inline() { + use ml::cuda_pipeline::sp15_isv_slots::{ + COOLDOWN_BARS_REMAINING_INDEX, HOLD_FLOOR_ALPHA_INDEX, + HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX, + }; + let stream = make_test_stream(); + let batch: usize = 1; + let n_atoms: i32 = 51; + let v_min: f32 = -1.0; + let v_max: f32 = 1.0; + let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0); + + // Near-uniform e_dir: [Short=0.50, Hold=0.45, Long=0.50, Flat=0.45]. + // Softmax(near-uniform) → high entropy ≈ ln(4). + let target_v = [0.50_f32, 0.45_f32, 0.50_f32, 0.45_f32]; + let b_logits_host = fill_peaked_logits_action_select( + batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0, + ); + let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z); + let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE]; + + // ISV bus with Thompson temp at floor (0.5) so q_eff stays close + // to e_dir; hold_floor anchors at α = 0.5, k = 10, ε₀ = 1.0; + // cooldown disabled (= 0). + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for isv bus"); + let mut isv_init = vec![0.0f32; ISV_LEN]; + // ISV[ISV_EVAL_THOMPSON_TEMP_IDX = 339] — 0.5 (floor). + isv_init[339] = 0.5; + isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5; + isv_init[HOLD_FLOOR_K_INDEX] = 10.0; + isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0; + isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0; + isv_buf.write_from_slice(&isv_init); + + // GPU buffers. + let q_values = stream.clone_htod(&q_values_host).expect("upload q_values"); + let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits"); + let per_sample_support = stream + .clone_htod(&support_host) + .expect("upload support"); + let mut actions_buf = stream.alloc_zeros::(batch).expect("alloc actions"); + let mut q_gaps_buf = stream.alloc_zeros::(batch).expect("alloc q_gaps"); + let mut intent_buf = stream.alloc_zeros::(batch).expect("alloc intent"); + let mut conv_buf = stream.alloc_zeros::(batch).expect("alloc conviction"); + let mut mag_conv_buf = stream.alloc_zeros::(batch).expect("alloc mag_conv"); + + let actions = launch_action_select_for_test( + &stream, batch, n_atoms, + &q_values, &b_logits_dir, &per_sample_support, + &mut actions_buf, &mut q_gaps_buf, &mut intent_buf, + &mut conv_buf, &mut mag_conv_buf, + isv_buf.dev_ptr, + ); + + let dir_picked = decode_dir(actions[0]); + assert_eq!( + dir_picked, AS_DIR_HOLD, + "hold_floor inline must shift argmax to Hold when entropy > ε₀; \ + got dir = {dir_picked} (expected {AS_DIR_HOLD})" + ); + } + + /// SP15 Phase 3.5.3.b — cooldown mask: when COOLDOWN_BARS_REMAINING > 0, + /// the picked direction is forced to Hold regardless of e_dir + /// posterior. Sets up e_dir = [Short=0.0, Hold=0.0, Long=1.0, Flat=0.0] + /// (Long strongly preferred), then activates cooldown (= 5.0). The + /// mask short-circuits Pass 2 entirely → dir = Hold. + #[test] + #[ignore = "requires GPU"] + fn action_select_forces_hold_during_cooldown() { + use ml::cuda_pipeline::sp15_isv_slots::{ + COOLDOWN_BARS_REMAINING_INDEX, HOLD_FLOOR_ALPHA_INDEX, + HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX, + }; + let stream = make_test_stream(); + let batch: usize = 1; + let n_atoms: i32 = 51; + let v_min: f32 = -1.0; + let v_max: f32 = 1.0; + let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0); + + // Long strongly preferred over Hold (e_dir gap = 1.0); without + // cooldown the argmax would land on Long. + let target_v = [0.0_f32, 0.0_f32, 1.0_f32, 0.0_f32]; + let b_logits_host = fill_peaked_logits_action_select( + batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0, + ); + let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z); + let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE]; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for isv bus"); + let mut isv_init = vec![0.0f32; ISV_LEN]; + isv_init[339] = 0.5; + isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5; + isv_init[HOLD_FLOOR_K_INDEX] = 10.0; + isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0; + // Cooldown active — kernel must force Hold. + isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 5.0; + isv_buf.write_from_slice(&isv_init); + + let q_values = stream.clone_htod(&q_values_host).expect("upload q_values"); + let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits"); + let per_sample_support = stream + .clone_htod(&support_host) + .expect("upload support"); + let mut actions_buf = stream.alloc_zeros::(batch).expect("alloc actions"); + let mut q_gaps_buf = stream.alloc_zeros::(batch).expect("alloc q_gaps"); + let mut intent_buf = stream.alloc_zeros::(batch).expect("alloc intent"); + let mut conv_buf = stream.alloc_zeros::(batch).expect("alloc conviction"); + let mut mag_conv_buf = stream.alloc_zeros::(batch).expect("alloc mag_conv"); + + let actions = launch_action_select_for_test( + &stream, batch, n_atoms, + &q_values, &b_logits_dir, &per_sample_support, + &mut actions_buf, &mut q_gaps_buf, &mut intent_buf, + &mut conv_buf, &mut mag_conv_buf, + isv_buf.dev_ptr, + ); + + let dir_picked = decode_dir(actions[0]); + assert_eq!( + dir_picked, AS_DIR_HOLD, + "cooldown mask must force Hold regardless of e_dir; \ + got dir = {dir_picked} (expected {AS_DIR_HOLD}, even though Long had +1.0 advantage)" + ); + } + + /// SP15 Phase 3.5.3.b — when cooldown counter is 0 the mask is + /// inactive; the picked direction follows the standard Thompson + + /// hold_floor argmax. With Long strongly preferred (e_dir Long=1.0, + /// rest=0.0) AND the entropy below ε₀ (one-hot ≈ 0 entropy) so + /// hold_floor ≈ 0, picked dir should be Long. + #[test] + #[ignore = "requires GPU"] + fn action_select_no_force_hold_when_cooldown_zero() { + use ml::cuda_pipeline::sp15_isv_slots::{ + COOLDOWN_BARS_REMAINING_INDEX, HOLD_FLOOR_ALPHA_INDEX, + HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX, + }; + let stream = make_test_stream(); + let batch: usize = 1; + let n_atoms: i32 = 51; + let v_min: f32 = -1.0; + let v_max: f32 = 1.0; + let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0); + + // Long strongly preferred; near-degenerate e_dir → entropy near + // 0 (well below ε₀ = 1.0) → hold_floor ≈ α × σ(-10) ≈ 0.5 × 4.5e-5 + // ≈ 2e-5 → Hold gain negligible → argmax stays at Long. + let target_v = [0.0_f32, 0.0_f32, 1.0_f32, 0.0_f32]; + let b_logits_host = fill_peaked_logits_action_select( + batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0, + ); + let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z); + let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE]; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for isv bus"); + let mut isv_init = vec![0.0f32; ISV_LEN]; + isv_init[339] = 0.5; + isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5; + isv_init[HOLD_FLOOR_K_INDEX] = 10.0; + isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0; + // Cooldown inactive. + isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0; + isv_buf.write_from_slice(&isv_init); + + let q_values = stream.clone_htod(&q_values_host).expect("upload q_values"); + let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits"); + let per_sample_support = stream + .clone_htod(&support_host) + .expect("upload support"); + let mut actions_buf = stream.alloc_zeros::(batch).expect("alloc actions"); + let mut q_gaps_buf = stream.alloc_zeros::(batch).expect("alloc q_gaps"); + let mut intent_buf = stream.alloc_zeros::(batch).expect("alloc intent"); + let mut conv_buf = stream.alloc_zeros::(batch).expect("alloc conviction"); + let mut mag_conv_buf = stream.alloc_zeros::(batch).expect("alloc mag_conv"); + + let actions = launch_action_select_for_test( + &stream, batch, n_atoms, + &q_values, &b_logits_dir, &per_sample_support, + &mut actions_buf, &mut q_gaps_buf, &mut intent_buf, + &mut conv_buf, &mut mag_conv_buf, + isv_buf.dev_ptr, + ); + + let dir_picked = decode_dir(actions[0]); + assert_eq!( + dir_picked, AS_DIR_LONG, + "with cooldown = 0 and Long strongly preferred, picked dir \ + must be Long; got dir = {dir_picked} (expected {AS_DIR_LONG})" + ); + } } /// SP15 Phase 1.7 — verify the trainer's `set_test_data_from_slices` API diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 656e93a07..a6f18b928 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 Phase 3.5.b + 3.5.3.b — wire hold_floor (inline) + cooldown mask into experience_action_select (2026-05-06): closes the deferred-consumer gap left by Phase 3.5 (`hold_floor_kernel.cu` + `launch_sp15_hold_floor`, commit `5d36f3238`) and Phase 3.5.3 (`cooldown_kernel.cu` + `launch_sp15_cooldown`, commit `649128e73`). Both phases landed only the producer + state-machinery scaffolding; the actual action-selection consumers were deferred per `feedback_no_partial_refactor.md` to keep the kernel-by-kernel landings small. Wave 1.B bundles both deferred consumer migrations atomically. (1) **Architectural decision — hold_floor is now an INLINE `__device__` computation, NOT a separate kernel call**: the standalone `hold_floor_kernel.cu` produced one f32 (`α × σ(k × (entropy − ε₀))`) per launch into a scratch buffer for `experience_action_select` to read back. Launching a kernel to write a single scalar that the next kernel immediately consumes is unnecessary indirection. The bounded-sigmoid math is trivial (3 ISV reads + one `sigmoidf` + one multiply); inlining it eliminates a launch, eliminates a scratch buffer, and removes a parallel-path liability per `feedback_wire_everything_up.md` + `feedback_no_legacy_aliases.md`. The 4 ISV slots (`HOLD_FLOOR_ALPHA=426`, `HOLD_FLOOR_K=427`, `HOLD_FLOOR_EPS0=428`, `ENTROPY_DIST_REF=429`) and their state_reset_registry entries + dispatch arms remain — only the launch path is removed. (2) **Entropy source — per-step direction-policy entropy from `softmax(e_dir)`**: the kernel already computes `e_dir[d] = E[Q(d)]` (joint C51 expected Q) for every direction during Pass 1 of the Thompson direction selector. Pass 1.5 (new) takes the stable softmax of those 4 e_dir floats and computes Shannon entropy `−Σ p[d] log p[d]` directly in registers — no atomic, no shared memory, no cross-thread reduction (one thread per sample). This is the natural per-thread "policy uncertainty" signal: high entropy = the posterior is near-uniform (model uncertain about WHERE to trade) → hold_floor lifts Hold's q_eff so Hold becomes the principled uncertainty expression rather than the distributional default; low entropy = the posterior concentrates on one direction (model confident) → hold_floor ≈ 0 → Hold gets no boost. (3) **q_eff_dir scratch — preserves e_dir for downstream consumers**: hold_floor is added to a local `q_eff_dir[4]` (= e_dir + hold_floor at DIR_HOLD), NOT to `e_dir[]` directly. The conviction + q_gap consumers downstream (`out_conviction`, `out_magnitude_conviction`, `out_q_gaps`) keep reading raw `e_dir` so the policy-confidence signals reflect the model's expressed preference — adding hold_floor there would corrupt the Kelly-cap warmup floor with a meta-confidence "be cautious" mask that's not what the consumer is asking for. (4) **Cooldown mask — hard short-circuit, NOT temperature-blended**: when `ISV[COOLDOWN_BARS_REMAINING=435] > 0`, Pass 2 is skipped entirely and `dir_idx = DIR_HOLD` is set directly. The alternative (q_eff_dir[non-Hold] = -INFINITY pre-blend) breaks under the Thompson temperature blend `q_eff = q_eff_dir + temp·(q_sample - q_eff_dir)`: at -INFINITY the algebra produces NaN; at -1e30 the formula reduces to `q_eff = q_sample` at temp=1 (the default, pure Thompson), defeating the mask. Hard short-circuit is the only correct semantics — and matches the spec's "force Hold for M bars" language exactly. The Philox draws Pass 2 would have made for direction sampling are not advanced on the cooldown path; this shifts the mag/ord/urg streams' state by a constant amount per cooldown step but reproducibility-at-fixed-seed is preserved because the cooldown gate is a deterministic function of the ISV state at this thread's bar. (5) **Deletion of `hold_floor_kernel.cu` + `launch_sp15_hold_floor` + `SP15_HOLD_FLOOR_CUBIN` + cubin manifest entry** per `feedback_no_legacy_aliases`: scaffolding for a deferred consumer becomes dead code once the consumer lands inline. The build.rs entry is replaced with a comment block documenting the move; the launcher is replaced with a stub doc comment in gpu_dqn_trainer.rs explaining the inline migration; the standalone `hold_floor_bounded_sigmoid` oracle test is removed (the math is now exercised end-to-end through the new `action_select_applies_hold_floor_inline` test). (6) **3 new oracle tests** in `crates/ml/tests/sp15_phase1_oracle_tests.rs::mod gpu`: `action_select_applies_hold_floor_inline` (near-uniform e_dir → high entropy ≈ ln(4) → hold_floor ≈ 0.49 lifts Hold's q_eff past Long's; argmax = Hold), `action_select_forces_hold_during_cooldown` (Long-preferred e_dir + cooldown=5.0 → mask short-circuit forces Hold regardless of e_dir), `action_select_no_force_hold_when_cooldown_zero` (Long-preferred e_dir + cooldown=0 + low-entropy → standard argmax wins, picked dir = Long). All 3 tests load the production `experience_kernels.cubin`, set up minimal C51 logits with peaked atoms (peak_logit=50 → softmax ~delta on closest atom, `q_sample` ≈ `e_dir`), and call `experience_action_select` end-to-end via cudarc. Tests pass on RTX 3050 Ti. (7) **Phase 3.5 + Phase 3.5.3 deferred consumers eliminated** per `feedback_wire_everything_up.md`. The cooldown_kernel itself (which maintains the consecutive_losses streak counter and decrements COOLDOWN_BARS_REMAINING per bar) is unchanged and still wired in production — only the consumer that reads its output is now connected to action-selection logic. Touched: `crates/ml/src/cuda_pipeline/experience_kernels.cu` (inline hold_floor + cooldown short-circuit added at the start of Pass 2 in `experience_action_select`; `e_dir` preserved for downstream consumers; ~80 LOC added), `crates/ml/src/cuda_pipeline/hold_floor_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (`launch_sp15_hold_floor` + `SP15_HOLD_FLOOR_CUBIN` removed; replaced by a non-doc comment block referencing the inline migration), `crates/ml/build.rs` (`"hold_floor_kernel.cu"` cubin manifest entry removed; replaced with a comment block explaining the move), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (removed `hold_floor_bounded_sigmoid` test; added 3 new `action_select_*` tests + `ProdActionSelectFixture` re-implementation in the integration-test module since the production fixture is `pub(crate)` and not reachable). Hard rules: `feedback_no_partial_refactor` (3.5.b + 3.5.3.b consumers + hold_floor_kernel deletion + 3 oracle tests + audit doc all in this commit; no parallel paths, no feature flags), `feedback_no_cpu_compute_strict` (hold_floor + cooldown mask are GPU-resident inside `experience_action_select` — no CPU forwards, no CPU reductions), `feedback_no_stubs` (the inline computation is the actual policy bias signal end-to-end, not a return-zero placeholder), `feedback_wire_everything_up` + `feedback_no_legacy_aliases` (hold_floor_kernel + launcher deleted because the consumer wiring landed inline; ISV slots + state_reset_registry entries + dispatch arms remain because they're still consumed by the inline path), `pearl_no_host_branches_in_captured_graph` (changes are inside `experience_action_select` which is already inside the existing CUDA Graph capture; no new launches added or removed from the captured stream — only one fewer kernel launch upstream of action_select on the cooldown side because the standalone hold_floor_kernel is gone, but it was always launched OUTSIDE the experience-collection capture region). Verified: `cargo check -p ml --features cuda` clean (18 unrelated pre-existing warnings); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored` runs 6 tests green (3 cooldown + 3 new action_select); `cargo test -p ml --features cuda --lib` holds 946 pass / 13 fail = baseline (no regression). + SP15 Phase 1.6.b + 1.7.b — wire dev-eval (Q8 final-fold) + test-eval (per-fold WF test slice) into trainer (2026-05-06): closes the deferred-consumer gap left by Phase 1.6 (CLI flags + `dev_features`/`holdout_features` stash, commit `ce019c72d`) and Phase 1.7 (`set_test_data_from_slices` observer + `test_features` stash, commit `ef373c34d`). Both phases landed only the data-flow scaffolding; the actual `evaluate_dqn_graphed` consumer was deferred per `feedback_no_partial_refactor.md` because wiring a parallel `GpuBacktestEvaluator` instance plus the TLOB-sync infrastructure is its own atomic refactor. Wave 1.A bundles both consumer migrations atomically. (1) **Architectural choice — parallel evaluator instances, NOT window-swap**: per the deferred-decision note in the Phase 1.7 audit row, the `gpu_evaluator` field's lazy-init path is fundamentally tied to the `val_data` window passed at construction time (`metrics.rs::launch_validation_loss` lines 549-633 build a single window from `self.val_data` once per fold and reuse the cached evaluator across epochs for deterministic cuBLAS state + stable buffer pointers). A window-swap on the val evaluator between val and dev/test eval calls would require invalidating the val CUDA Graph (recorded device-pointer values become stale when the underlying buffers are repointed) every epoch — fragile and a direct violation of `pearl_no_host_branches_in_captured_graph.md` (the cached graph records the buffer set at capture time; touching the buffers between captures is exactly the host-driven branch the pearl forbids). Parallel instances mirror the val_evaluator construction pattern verbatim and keep each evaluator's CUDA Graph self-contained. (2) **New trainer fields** (`crates/ml/src/trainers/dqn/trainer/mod.rs:720-744`): `dev_evaluator: Option` + `test_evaluator: Option` (both `None` at construction, lazy-built on first call). `dev_ofi_offset: usize` mirrors `ofi_val_offset` for dev; `test_start_bar` (already on the struct from Phase 1.7) provides the test-side OFI offset. (3) **Lazy-init helper** (`crates/ml/src/trainers/dqn/trainer/metrics.rs::launch_extra_eval`, +280 LOC): single helper takes an `ExtraEvalKind` discriminator (`Dev` / `Test`) + optional `fold_idx` and routes to the right `Option` slot. Mirrors `launch_validation_loss` for TLOB sync (`set_tlob_from_training` on the eval-stream's per-stream cuBLAS handle), ISV signal pointer (`set_isv_signals_ptr(fused.isv_signals_dev_ptr())`), `eval_v_range()` Q-range, network dims + branch sizes from the agent, and the `DqnBacktestConfig` struct construction. The helper diverges from val in three places: (a) it uses synchronous `evaluate_dqn_graphed` (one-shot) instead of `evaluate_dqn_graphed_async` + deferred `consume_metrics_after_event` — there's no async pipelining benefit because dev/test eval doesn't fire per-epoch; (b) `val_subsample_stride: 1` instead of val's `4` — one-shot eval doesn't need the val cost optimisation, and stride=1 keeps the kernel's annualisation factor honest; (c) the eval evaluator is rebuilt every fold for `Test` (slice length varies per fold; the evaluator allocates step buffers sized for `max_len` at construction) but built once and kept for `Dev` (only one Q8 dev slice per training run, never changes mid-run). The dev/test path does NOT clobber `self.last_eval_magnitude_dist` / `self.last_eval_intent_magnitude_dist` / `self.last_val_metrics` — those are tied to the val evaluator's most-recent state and the dev/test runs would otherwise overwrite them with non-val readings. (4) **Wire-up site for test_eval — inside the fold loop** (`crates/ml/src/trainers/dqn/trainer/mod.rs` after `train_with_data_full_loop_slices` returns, before regime metrics emit): gated on `fold.test_end > fold.test_start && self.test_features.is_some()`. The `set_test_data_from_slices` call already gates on the same `fold.test_end > fold.test_start` (Phase 1.7), so when `test_features` is non-empty we know the slice is real. Errors are logged + non-fatal: a failed test_slice eval shouldn't abort the fold loop because test eval is observation-only (no Q-targets, no replay writes, no downstream consumers). (5) **Wire-up site for dev_eval — after the fold loop** (`crates/ml/src/trainers/dqn/trainer/mod.rs` immediately after the regime replay decay reset): gated on `self.dev_features.is_some()`. The `dev_ofi_offset = dev_start` is set inside the dev-features stash block (`mod.rs:1163`), so by the time the fold loop completes we have a valid offset for the global OFI features array. Sealed Q9 holdout remains untouched — the existing `debug_assert` sealed-slice guard at `mod.rs:1228` catches accidental refactors that reintroduce holdout into the training path; Phase 4.3 `argo-eval-final.sh` is the only legitimate Q9 consumer via a separate eval-only entry point that does NOT call `train_walk_forward`. (6) **HEALTH_DIAG line formats**: `HEALTH_DIAG[N]: dev_eval dev_sharpe_net= dev_calmar= dev_max_dd= dev_trades=` (emitted once after the final fold) and `HEALTH_DIAG[N]: test_slice fold=K test_sharpe_net= test_calmar= test_max_dd= test_trades=` (emitted once per fold, K = 0-based fold index). N is `self.current_epoch` (last epoch of the fold for both kinds). The `*_sharpe_net` key uses the existing fused-metrics-kernel cost-aware Sharpe (post-Phase-1.1.b split — `WindowMetrics.sharpe` already includes `tx_cost_bps + spread_cost` via the env-step PnL feed); when Phase 1.2.b cost-net sharpe lands (BLOCKED pending LobBar merge per task #358), the consumer will swap to the dedicated cost-net kernel output without changing the HEALTH_DIAG key name — preserves the aggregator-script contract `aggregate-multi-seed-metrics.py` reads from. (7) **Stream choice**: `validation_stream` when present (mirrors val_evaluator's stream), with `cuda_stream` as fallback on the legacy single-stream path. Records a fresh event on `cuda_stream` and waits on `validation_stream` before launching eval kernels — same cross-stream barrier the val path uses (`metrics.rs:532-547`) so the most-recent training kernel's weight + ISV writes are globally visible before the eval kernels read them. CPU-only build path skips silently (no streams → no-op fall-through), same as val. (8) **OFI alignment**: dev_features/test_features are at GLOBAL bar indices (dev_start..dev_end and fold.test_start..fold.test_end respectively). `init_from_fxcache` stashes the FULL ofi-features array on `self.ofi_features` (line `mod.rs:1889`), so the eval gather reads `ofi_features[ofi_offset + i]` for sample `i` — `ofi_offset = self.dev_ofi_offset` for Dev (`= dev_start` in train_walk_forward) or `self.test_start_bar` for Test (already populated by `set_test_data_from_slices`). The `Arc::from(ofi)` reassignment at line 1889 keeps the Arc pointing at the FULL pre-train_walk_forward dataset because the trainer's own `self.ofi_features` was set up at the original 9-quarter scale before `train_walk_forward` was called. (9) **Phase 1.6 + 1.7 orphan stash-fields now have consumers** per `feedback_wire_everything_up.md`: `dev_features` / `dev_targets` (Phase 1.6) + `test_features` / `test_targets` (Phase 1.7) all reach a real GPU eval kernel with an emitted HEALTH_DIAG line. Touched: `crates/ml/src/trainers/dqn/trainer/mod.rs` (+3 fields on `DQNTrainer` struct: `dev_evaluator`, `dev_ofi_offset`, `test_evaluator`; +1-line `self.dev_ofi_offset = dev_start;` write inside the existing dev-features stash block; +14-line test_eval call inside the fold loop after `train_with_data_full_loop_slices`; +14-line dev_eval call after the regime-replay-decay reset post-fold-loop), `crates/ml/src/trainers/dqn/trainer/constructor.rs` (+5 lines initialising the 3 new fields to None/0), `crates/ml/src/trainers/dqn/trainer/metrics.rs` (+1 `ExtraEvalKind` enum at top of file; +~280-line `launch_extra_eval` helper just before `estimate_avg_q_value_with_early_stopping`). Hard rules: `feedback_no_partial_refactor` (both 1.6.b and 1.7.b consumers + both new fields + both HEALTH_DIAG lines + audit doc all in this commit; no parallel paths, no feature flags), `feedback_no_cpu_compute_strict` (dev/test eval is GPU-resident through the same `evaluate_dqn_graphed` path as val — no CPU forwards, no CPU reductions), `feedback_no_stubs` (both call sites are real GPU eval invocations against real stashed slices; failure is logged + non-fatal but the kernel actually fires when the slice is non-empty and the CUDA stream is present), `feedback_isv_for_adaptive_bounds` (no hardcoded thresholds introduced; the dev/test eval inherits all ISV-driven adaptive bounds from `fused_ctx.isv_signals_dev_ptr()` via the same `set_isv_signals_ptr` call val uses), `pearl_no_host_branches_in_captured_graph` (each evaluator's CUDA Graph is self-contained; no graph-capture interaction between dev/test/val because they live in separate `Option` slots with their own buffer ptrs and capture metadata; `evaluator.invalidate_dqn_graph()` is called before each launch_extra_eval invocation to force re-capture against the live training weights). Verified: `cargo check -p ml --features cuda` clean (3 unrelated pre-existing warnings only); `cargo test -p ml --features cuda --lib` holds the 946 pass / 13 fail baseline (no regression introduced); the existing `set_test_data_from_slices_fires_observer_and_stashes` Phase 1.7 oracle test still passes (the new code didn't change the observer-firing surface); a unit-style oracle test asserting the dev/test evaluators actually get built was NOT added because the build path requires a fully-initialised `fused_ctx` with TLOB params + agent forward path + GPU walk-forward generator, which exceeds the smoke-test fixture scope — L40S is the canonical verifier per the spec's existing "L40S smoke is the validation path" pattern. The two new HEALTH_DIAG lines will appear in the next L40S run output once a fold finishes (test_slice) and once the full training run finishes (dev_eval), giving the multi-seed aggregator real `dev_*` / `test_*` keys to track. SP15 Phase 1.3.b — wire dd_state per-step launch + drop equity-recompute bug + env-0 canonical observable (2026-05-06): closes the Phase 1.3 orphan-launcher gap per `feedback_wire_everything_up.md` and atomically fixes two architectural blockers found during the wire-up investigation (Path A: minimum-scope fix; per-env redesign deferred to Phase 1.3.b-followup if smoke shows env-0-only aggregation insufficient). (1) **Double-update bug fix** (`crates/ml/src/cuda_pipeline/dd_state_kernel.cu`): the original kernel recomputed `new_equity = pos_state[PS_PREV_EQUITY] + pnl_step` and wrote it back into the same slot, but `experience_env_step` already maintains `PS_PREV_EQUITY = new_portfolio_value` (`experience_kernels.cu:3473-3475`). Wiring the kernel as-is would silently double-accumulate equity every step (1× from env_step, 1× from dd_state's recompute), corrupting the position state buffer for every downstream consumer (Kelly EMAs, dd penalty, trade-stats). Per `feedback_no_quickfixes.md` the recompute is removed outright (no `if (already_updated)` guards, no sentinel skip-stores) — the kernel is now READ-ONLY on `pos_state[PS_PREV_EQUITY]` and `pos_state[PS_PEAK_EQUITY]`. The `pnl_step` parameter is dropped from both the CUDA kernel signature and the `launch_sp15_dd_state` launcher signature accordingly. (2) **Env-0 canonical observable** (`crates/ml/src/cuda_pipeline/dd_state_kernel.cu`): kernel grid is `[1,1,1]` (single thread, single block) and the 6 DD ISV slots [401..407) are scalars, but production has N envs in the portfolio_states buffer. The kernel now reads `pos_state[0 * PS_STRIDE + PS_PREV_EQUITY]` and `pos_state[0 * PS_STRIDE + PS_PEAK_EQUITY]` — env 0 specifically — documented in the kernel header as "env 0 as canonical DD observable; per-env redesign deferred to Phase 1.3.b-followup." This is the minimum-scope fix that produces a coherent per-step DD signal for downstream consumers; it accepts the limitation that the DD slots reflect env-0's history rather than aggregate DD across all envs. The recovery / persistence counter logic is also adapted: zero on `current_dd ≤ 0` (new HWM or pre-trade flat) instead of the original `new_equity > peak` branch (which referenced the now-gone local variable). The two conditions are equivalent at the env-0 reading because env_step has just written `peak = max(prev_peak, prev_equity)` ≥ `prev_equity`, so `current_dd = max(0, (peak - prev_equity)/peak)` evaluates to exactly 0 iff env_step set a new HWM this step. (3) **Wire-up location**: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` step "5b" inside `launch_timestep_loop`, immediately after the `experience_env_step` launch (line ~4378, after the env_step kernel block ends at line ~4379) and before the step counter advance kernel (line ~4385). Same stream as env_step → CUDA serializes the two on-stream so the dd_state read sees env_step's writes (no event sync required). OUTSIDE the exp-fwd graph capture region (which is bounded `begin_capture` line ~3734 → `end_capture` line ~3829, well before env_step) so per `pearl_no_host_branches_in_captured_graph.md` no graph-capture interaction. Launcher invoked with `(stream, config.dd_threshold, self.isv_signals_dev_ptr, self.portfolio_states.raw_ptr())` — `dd_budget = config.dd_threshold` keeps the DD_PCT denominator in sync with the `w_dd` penalty term env_step applies on the same `dd_threshold` knob. (4) **Oracle test migration** (`crates/ml/tests/sp15_phase1_oracle_tests.rs::dd_state_kernel_tracks_drawdown_correctly`): test was written against the OLD kernel signature so it WOULD have failed to compile after the kernel signature change — atomic migration per `feedback_no_partial_refactor.md`. The test no longer passes `pnl_step` to the launcher; instead it pre-sums the per-step PnL series into cumulative equity values, sets `pos_state[PS_PREV_EQUITY] = equity` and `pos_state[PS_PEAK_EQUITY] = running_max(equity)` directly into the env-0 row of the portfolio buffer before each `launch_sp15_dd_state` call (mirrors what env_step does in production). The same six-step equity curve (100, 110, 90, 105, 95, 100) drives the kernel and the same expected results hold (max_dd ≈ 0.182, current_dd ≈ 0.091, recovery_bars ≥ 4, dd_pct ∈ (0,1]). Test passes on RTX 3050 Ti. (5) **Phase 1.3 orphan launcher closed** per `feedback_wire_everything_up.md` — `launch_sp15_dd_state` now has a production caller in the per-step training-loop hot path. Downstream consumers receive live values when their consumer wiring lands: Phase 3.3 `dd_penalty` (reads DD_CURRENT/DD_MAX), Phase 3.5.2 `dd_asymmetric_reward` (reads DD_PCT), Phase 3.5.4 `plasticity_injection` persistence (reads DD_PERSISTENCE), Phase 3.5.5 `dd_trajectory` (reads DD_PCT). All four consumer kernels are still orphan launchers themselves at the time of this commit; their wire-up is independent of this one. (6) **Per-env DD redesign deferral** (Phase 1.3.b-followup, task #360): if L40S smoke shows the env-0 single-observable aggregation produces stale or biased DD signals when envs diverge significantly, the followup will introduce a per-env DD tile + reduction kernel that aggregates DD across envs (e.g., max_dd over all envs, or weighted-mean DD by `|position|`). The current min-scope fix is sufficient for representative-env semantics that the DD-aware reward terms expect (one DD signal per training instance). Touched: `crates/ml/src/cuda_pipeline/dd_state_kernel.cu` (kernel — pnl_step param dropped, equity recompute removed, env-0 reads via `0 * PS_STRIDE`, recovery-counter branch adapted to `current_dd ≤ 0`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launcher — `pnl_step` param dropped from `launch_sp15_dd_state` signature; docstring rewritten to reflect read-only contract + env-0 canonical observable + post-env_step ordering requirement), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (production caller — new step "5b" block in `launch_timestep_loop` after env_step launch, calls `launch_sp15_dd_state(&self.stream, config.dd_threshold, self.isv_signals_dev_ptr, self.portfolio_states.raw_ptr())`), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (oracle migration — equity series pre-summed and written into `pos_state` directly per the new contract). Hard rules: `feedback_no_partial_refactor` (kernel signature change + oracle test migration + production wire-up land in this commit; per-env redesign is a separate atomic followup), `feedback_no_quickfixes` (recompute removed outright; no skip-store guards), `feedback_wire_everything_up` (closes the Phase 1.3 orphan launcher), `pearl_no_host_branches_in_captured_graph` (verified launch site is outside the exp-fwd graph capture region). Verified: `cargo check -p ml --features cuda` clean, oracle test `dd_state_kernel_tracks_drawdown_correctly` green, ml lib suite holds 946 pass / 13 fail = baseline.