diff --git a/crates/ml/build.rs b/crates/ml/build.rs index c468821e6..f016b2875 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -939,6 +939,30 @@ 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", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 06682dcb2..7a097aa87 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1477,6 +1477,80 @@ 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(()) +} + /// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty signal — lookup + /// update kernels sharing one cubin. Lookup reads /// `1/sqrt(1+count)` ∈ [0, 1] for each (state, action) bucket; update @@ -20584,6 +20658,43 @@ impl GpuDqnTrainer { *sig_ptr.add(LAMBDA_REGRET_INDEX) = 1.0_f32; *sig_ptr.add(REGRET_GRAD_NORM_INDEX) = 0.0_f32; + // SP15 Phase 3.5 (2026-05-06): confidence-aware Hold + // floor anchors (per spec §8.2 (3.5) post-amendment-2 + // fix). `hold_floor = α × σ(k × (entropy − ε₀))` — + // 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 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. Until the + // producers land the structural cold-start anchors + // hold the Hold floor at a finite scale per + // `feedback_isv_for_adaptive_bounds.md`: + // HOLD_FLOOR_ALPHA = 0.5 (sentinel — half a Q-unit + // at peak σ) + // HOLD_FLOOR_K = 10.0 (steep sigmoid — fast + // saturation) + // HOLD_FLOOR_EPS0 = 1.0 (entropy threshold — + // ~1 nat ≈ ln(e)) + // ENTROPY_DIST_REF = 0.0 (Pearl A sentinel — first + // non-zero observation from + // the deferred ε₀ producer + // replaces directly per + // `pearl_first_observation_bootstrap`) + use crate::cuda_pipeline::sp15_isv_slots::{ + ENTROPY_DIST_REF_INDEX, HOLD_FLOOR_ALPHA_INDEX, + HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX, + }; + *sig_ptr.add(HOLD_FLOOR_ALPHA_INDEX) = 0.5_f32; + *sig_ptr.add(HOLD_FLOOR_K_INDEX) = 10.0_f32; + *sig_ptr.add(HOLD_FLOOR_EPS0_INDEX) = 1.0_f32; + *sig_ptr.add(ENTROPY_DIST_REF_INDEX) = 0.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 diff --git a/crates/ml/src/cuda_pipeline/hold_floor_kernel.cu b/crates/ml/src/cuda_pipeline/hold_floor_kernel.cu new file mode 100644 index 000000000..8979cefc5 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/hold_floor_kernel.cu @@ -0,0 +1,65 @@ +// 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/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 541cf22e5..4daf52e3e 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1178,6 +1178,43 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[REGRET_GRAD_NORM_INDEX=425] — SP15 Phase 3.4 (2026-05-06) stateful EMA of the regret-penalty gradient L2 norm. Read by the follow-up λ_regret grad-balance producer kernel (deferred per `feedback_no_partial_refactor.md`); FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md`. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.4).", }, + // SP15 Phase 3.5 (2026-05-06) — confidence-aware Hold-floor + // anchors (per spec §8.2 (3.5) post-amendment-2 fix). + // `hold_floor = α × σ(k × (entropy − ε₀))` — 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. The signal-driven producers that + // overwrite α / k / ε₀ 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; FoldReset + // rewrites the structural cold-start anchors so the + // absorbers re-engage on each new fold per + // `feedback_isv_for_adaptive_bounds.md`. ENTROPY_DIST_REF + // is the running 75th-percentile reference the deferred ε₀ + // producer reads — Pearl A sentinel 0. + RegistryEntry { + name: "sp15_hold_floor_alpha", + category: ResetCategory::FoldReset, + description: "ISV[HOLD_FLOOR_ALPHA_INDEX=426] — SP15 Phase 3.5 (2026-05-06) confidence-aware Hold-floor scale factor (per spec §8.2 (3.5) post-amendment-2 fix). FoldReset rewrites the structural cold-start anchor 0.5 per `feedback_isv_for_adaptive_bounds.md` — the runtime value will be signal-driven from a follow-up producer kernel that tracks the rolling 95th percentile of |Q_dir| (NOT running max — outlier-ratchet vulnerable per spec second-review #6); the 0.5 anchor re-engages on each new fold so the cold-start absorber holds α at a finite scale until the producer accumulates observations. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.5).", + }, + RegistryEntry { + name: "sp15_hold_floor_k", + category: ResetCategory::FoldReset, + description: "ISV[HOLD_FLOOR_K_INDEX=427] — SP15 Phase 3.5 (2026-05-06) confidence-aware Hold-floor sigmoid steepness (per spec §8.2 (3.5)). FoldReset rewrites the structural cold-start anchor 10.0 per `feedback_isv_for_adaptive_bounds.md` — the runtime value will be signal-driven from a follow-up producer kernel that tracks the running variance of entropy; the 10.0 anchor re-engages on each new fold. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.5).", + }, + RegistryEntry { + name: "sp15_hold_floor_eps0", + category: ResetCategory::FoldReset, + description: "ISV[HOLD_FLOOR_EPS0_INDEX=428] — SP15 Phase 3.5 (2026-05-06) confidence-aware Hold-floor entropy threshold ε₀ (per spec §8.2 (3.5)). FoldReset rewrites the structural cold-start anchor 1.0 per `feedback_isv_for_adaptive_bounds.md` — the runtime value will be signal-driven from a follow-up producer kernel that tracks the 75th percentile of the running entropy distribution (`ENTROPY_DIST_REF_INDEX=429`); the 1.0 anchor re-engages on each new fold. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.5).", + }, + RegistryEntry { + name: "sp15_entropy_dist_ref", + category: ResetCategory::FoldReset, + description: "ISV[ENTROPY_DIST_REF_INDEX=429] — SP15 Phase 3.5 (2026-05-06) running 75th-percentile reference for the entropy distribution. Read by the follow-up ε₀ producer kernel (deferred per `feedback_no_partial_refactor.md`); FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first non-zero observation per `pearl_first_observation_bootstrap.md`. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.5).", + }, // SP5 Task A1: Wiener-state companion reset. The wiener_state_buf // covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) + // SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 49e530bb8..ed4385368 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -8250,6 +8250,45 @@ impl DQNTrainer { fused.trainer().write_isv_signal_at(REGRET_GRAD_NORM_INDEX, 0.0); } } + // SP15 Phase 3.5 (2026-05-06): confidence-aware Hold-floor + // anchors (per spec §8.2 (3.5) post-amendment-2 fix). + // HOLD_FLOOR_ALPHA / K / EPS0 are structural cold-start + // anchors per `feedback_isv_for_adaptive_bounds.md` — the + // runtime values are signal-driven from the follow-up + // producers updating them 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; the + // anchors re-engage on each new fold so the cold-start + // absorbers hold the Hold floor at a finite scale until the + // producers accumulate observations. ENTROPY_DIST_REF is a + // stateful EMA — FoldReset sentinel 0 (Pearl A) so the + // deferred ε₀ producer's first-observation replacement + // fires on the new fold's first non-zero entropy. + "sp15_hold_floor_alpha" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::HOLD_FLOOR_ALPHA_INDEX; + fused.trainer().write_isv_signal_at(HOLD_FLOOR_ALPHA_INDEX, 0.5); + } + } + "sp15_hold_floor_k" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::HOLD_FLOOR_K_INDEX; + fused.trainer().write_isv_signal_at(HOLD_FLOOR_K_INDEX, 10.0); + } + } + "sp15_hold_floor_eps0" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::HOLD_FLOOR_EPS0_INDEX; + fused.trainer().write_isv_signal_at(HOLD_FLOOR_EPS0_INDEX, 1.0); + } + } + "sp15_entropy_dist_ref" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::ENTROPY_DIST_REF_INDEX; + fused.trainer().write_isv_signal_at(ENTROPY_DIST_REF_INDEX, 0.0); + } + } // SP11 Task A2 (2026-05-04): novelty visit-count hash table // reset arm — closes the A0 deferral per the registry entry's // docstring. 1M f32 slots, zero-filled via mapped-pinned host diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 39fd9729a..133fa823b 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -992,6 +992,75 @@ mod gpu { "case B: regret = {regret_b}, expected 0 (ideal <= cost+trail)" ); } + + /// 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 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 5af9f2324..4d95100ca 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 — confidence-aware Hold floor (bounded sigmoid) (2026-05-06): per spec §8.2 (3.5) post-amendment-2 fix. New single-kernel cubin `hold_floor_kernel.cu` 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. Computes `hold_floor = α × σ(k × (entropy − ε₀))` from ISV[HOLD_FLOOR_ALPHA_INDEX=426] + ISV[HOLD_FLOOR_K_INDEX=427] + ISV[HOLD_FLOOR_EPS0_INDEX=428] (read-only — kernel does NOT update them in this commit; the signal-driven producers updating them 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); writes a single f32 `hold_floor` value to a caller-supplied output buffer. 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. 4 new ISV slots: `HOLD_FLOOR_ALPHA_INDEX=426` (initial 0.5; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up producer kernel that tracks the rolling 95th percentile of |Q_dir| (NOT running max — outlier-ratchet vulnerable); the 0.5 anchor re-engages on each new fold), `HOLD_FLOOR_K_INDEX=427` (initial 10.0; structural cold-start anchor — runtime value signal-driven from follow-up running-variance-of-entropy producer; 10.0 anchor re-engages each fold), `HOLD_FLOOR_EPS0_INDEX=428` (initial 1.0; structural cold-start anchor — runtime value signal-driven from follow-up 75th-percentile-of-entropy producer reading `ENTROPY_DIST_REF_INDEX=429`; 1.0 anchor re-engages each fold), `ENTROPY_DIST_REF_INDEX=429` (initial 0.0; stateful EMA, FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md`). 4 new `state_reset_registry` entries (`sp15_hold_floor_alpha` / `sp15_hold_floor_k` / `sp15_hold_floor_eps0` rewrite the 0.5 / 10.0 / 1.0 anchors at fold boundary so the cold-start absorbers re-engage on each new fold; `sp15_entropy_dist_ref` resets to 0). 4 new dispatch arms in `training_loop.rs` enforce the registry-arm contract (`every_fold_and_soft_reset_entry_has_dispatch_arm` regression test). Anchor tests 2.1 flat_market_holds + 2.6 regime_silences (Phase 2B contracts) — green via the deferred action-selection wiring (add `hold_floor` to Q_hold pre-argmax/Thompson); this commit lands the Hold-floor primitive. **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 — kernel + launcher verify in isolation first via the GPU oracle test below). Touched: `crates/ml/src/cuda_pipeline/hold_floor_kernel.cu` (new — single kernel, single-thread/single-block, mirrors `dd_penalty_kernel.cu`'s literal-index pattern with `__threadfence_system()` after the write; the kernel's ISV pointer is `const float*` because the producers updating α / k / ε₀ are deferred follow-up sub-tasks), `crates/ml/build.rs` (+1 cubin manifest entry in `kernels_with_common`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_HOLD_FLOOR_CUBIN` cubin slot + 1 `pub fn launch_sp15_hold_floor` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_HOLD_FLOOR_CUBIN.to_vec())` and resolving `get_function("hold_floor_kernel")` — single-thread/single-block grid mirroring `launch_sp15_dd_penalty` precedent + 4 ISV constructor writes (HOLD_FLOOR_ALPHA=0.5, HOLD_FLOOR_K=10.0, HOLD_FLOOR_EPS0=1.0, ENTROPY_DIST_REF=0.0) appended to the SP15 Phase 3.4 anchor block), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+4 `RegistryEntry` records — `sp15_hold_floor_alpha` rewrites 0.5 anchor; `sp15_hold_floor_k` rewrites 10.0 anchor; `sp15_hold_floor_eps0` rewrites 1.0 anchor; `sp15_entropy_dist_ref` Pearl A sentinel 0), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+4 dispatch arms for the 4 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 GPU oracle test in `mod gpu`: `hold_floor_bounded_sigmoid` validates BOTH gate sides sharing one ISV buffer with α=0.5, k=10, ε₀=1.0 — high-entropy `entropy=1.5` yields `hold_floor ≈ 0.5 × σ(5) ≈ 0.497` (strong Hold bias); low-entropy `entropy=0.5` yields `hold_floor ≈ 0.5 × σ(-5) < 0.05` (near-zero floor); sentinel non-zero output buffer between launches ensures the kernel actually overwrites with the small low-entropy value rather than skipping the store). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + registry entries + dispatch arms land atomically; action-selection wiring follows as its own atomic commit), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel actually computes the bounded sigmoid; the oracle test verifies BOTH gate sides — high-entropy strong Hold bias AND low-entropy near-zero floor with sentinel-overwrite semantics rather than skip-store), `feedback_no_htod_htoh_only_mapped_pinned` (oracle test uses `MappedF32Buffer` for both ISV bus and hold_floor output buffer; `entropy` scalar is by-value at the launcher boundary), `feedback_isv_for_adaptive_bounds` (the 0.5 / 10.0 / 1.0 anchors are structural cold-start absorbers rewritten at fold boundary so the cold-start absorbers re-engage per phase precedent; runtime α / k / ε₀ will be signal-driven from the deferred producers), `pearl_bounded_modifier_outputs_require_structural_activation` (the bounded sigmoid IS the structural activation per the spec — α anchors the bound without a downstream runtime clamp; the producer that updates α from rolling 95th-percentile-of-|Q_dir| keeps the bound coupled to the Q-scale and avoids outlier ratchet from running-max — this is the load-bearing pearl mechanism for the post-amendment-2 fix), `pearl_first_observation_bootstrap` (ENTROPY_DIST_REF sentinel 0 so the deferred ε₀ producer's Pearl A first-observation replacement fires on the new fold's first non-zero entropy). + SP15 Phase 3.4 — regret signal = r_discipline content + first-observation bootstrap EMA (2026-05-06): per spec §8.2 (3.4). New single-kernel cubin `regret_signal_kernel.cu` with one `extern "C" __global__` symbol (`regret_signal_kernel`) → one cubin → one launcher (`launch_sp15_regret_signal`) — mirrors the Phase 3.3 `dd_penalty_kernel.cu` single-kernel-per-cubin pattern. Computes `regret_bar = ideal_pnl_missed − cost_t` when `action == Hold (0)` AND `ideal_pnl_missed > cost_t + trail_dist`, else 0; updates ISV[REGRET_EMA_INDEX=423] (read-modify-write) via first-observation bootstrap (per `pearl_first_observation_bootstrap` — sentinel 0 replaced directly on the first non-zero observation) + simple exponential decay α=0.05 (Wiener-α adaptive smoothing per `pearl_wiener_optimal_adaptive_alpha` is a documented follow-up). Asymmetric gate: regret fires only when policy was Hold AND the ideal trade cleared the slippage+trail floor (so the lost opportunity is real after overhead). 3 new ISV slots: `REGRET_EMA_INDEX=423` (initial 0.0; FoldReset sentinel 0 so the kernel's first-observation bootstrap re-engages on the new fold's first non-zero `regret_bar`), `LAMBDA_REGRET_INDEX=424` (initial 1.0; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up grad-balance producer once `REGRET_GRAD_NORM_INDEX=425` accumulates observations), `REGRET_GRAD_NORM_INDEX=425` (initial 0.0; stateful EMA, FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation). 3 new `state_reset_registry` entries (`sp15_regret_ema` resets to 0; `sp15_lambda_regret` rewrites the 1.0 anchor at fold boundary so the cold-start absorber re-engages on each new fold; `sp15_regret_grad_norm` resets to 0). 3 new dispatch arms in `training_loop.rs` enforce the registry-arm contract (`every_fold_and_soft_reset_entry_has_dispatch_arm` regression test). Anchor test 2.6 regime_silences (Phase 2B contract) — green via Phase 3.4 + 3.5 combined; this commit lands the regret primitive. **Phase 3.4 lands kernel + launcher + 3 ISV anchor seeds + 3 registry entries + 3 dispatch arms only**; the host-side reward composer that subtracts `LAMBDA_REGRET × REGRET_EMA` from r_discipline is deferred to a follow-up atomic commit per `feedback_no_partial_refactor.md` (matching Phase 1.1-1.3 + 3.1-3.3 precedent — kernel + launcher verify in isolation first via the GPU oracle tests below). Touched: `crates/ml/src/cuda_pipeline/regret_signal_kernel.cu` (new — single kernel, single-thread/single-block, mirrors `dd_penalty_kernel.cu`'s literal-index pattern with `__threadfence_system()` after the EMA write; the kernel's ISV pointer is mutable `float*` because it read-modify-writes REGRET_EMA, contrasted with `dd_penalty_kernel`'s `const float*` read-only ISV), `crates/ml/build.rs` (+1 cubin manifest entry in `kernels_with_common`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_REGRET_SIGNAL_CUBIN` cubin slot + 1 `pub fn launch_sp15_regret_signal` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_REGRET_SIGNAL_CUBIN.to_vec())` and resolving `get_function("regret_signal_kernel")` — single-thread/single-block grid mirroring `launch_sp15_dd_penalty` precedent + 3 ISV constructor writes (REGRET_EMA=0.0, LAMBDA_REGRET=1.0, REGRET_GRAD_NORM=0.0) appended to the SP15 Phase 3.3 anchor block), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+3 `RegistryEntry` records — `sp15_regret_ema` Pearl A sentinel 0; `sp15_lambda_regret` rewrites 1.0 anchor; `sp15_regret_grad_norm` Pearl A sentinel 0), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+3 dispatch arms for the 3 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+2 GPU oracle tests in `mod gpu`: `regret_logged_on_missed_profitable_trade` validates `action=Hold(0), ideal=5.0, cost=1.0, trail=0.5 → regret_bar = 5.0 − 1.0 = 4.0` AND first-observation bootstrap fires sentinel 0 → 4.0 in REGRET_EMA; `regret_zero_when_action_active_or_ideal_below_threshold` validates the asymmetric gate via 2 sub-cases sharing one ISV buffer — Case A `action=1 (active), ideal=5.0` produces zero regardless, Case B `action=Hold, ideal=1.0 ≤ cost+trail=1.5` produces zero). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic + read-modify-write of one EMA slot, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + registry entries + dispatch arms land atomically; reward-composer subtraction follows as its own atomic commit), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel computes the asymmetric regret gate AND writes the EMA back to ISV; the second oracle test verifies BOTH gate paths actually produce zero rather than skipping the store), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both ISV bus and regret-bar output buffer; scalar args `action`/`ideal_pnl_missed`/`cost_t`/`trail_dist` are by-value at the launcher boundary), `feedback_isv_for_adaptive_bounds` (the 1.0 LAMBDA_REGRET anchor is a structural cold-start absorber rewritten at fold boundary so the cold-start absorber re-engages per phase precedent; runtime LAMBDA_REGRET will be signal-driven from the deferred grad-balance producer), `pearl_first_observation_bootstrap` (REGRET_EMA sentinel 0 with explicit `if (prev_ema == 0.0f && regret_bar > 0.0f) { new_ema = regret_bar; }` branch in the kernel — the new fold's first non-zero `regret_bar` replaces the sentinel directly so the EMA is not anchored to its boot-time zero), `pearl_wiener_optimal_adaptive_alpha` (DOCUMENTED FOLLOW-UP — initial commit uses fixed α=0.05 per the task spec's deferral; the swap to Wiener-α adaptive smoothing using the existing SP4 wiener_state companion buffer pattern lands as a follow-up atomic commit and does NOT block this Phase's primitive landing). SP15 Phase 3.3 — quadratic DD penalty + ISV-driven λ + threshold (2026-05-06): per spec §8.2 (3.3). New single-kernel cubin `dd_penalty_kernel.cu` with one `extern "C" __global__` symbol (`dd_penalty_kernel`) → one cubin → one launcher (`launch_sp15_dd_penalty`) — mirrors the Task 1.3 `dd_state_kernel.cu` single-kernel-per-cubin pattern (NOT the Task 3.1 multi-kernel pattern; Phase 3.3 has a single kernel). Computes `penalty = λ_dd × max(0, dd_current − dd_threshold)²` from ISV[DD_CURRENT_INDEX=401] (written by Task 1.3's `dd_state_kernel`, already landed) + ISV[LAMBDA_DD_INDEX=420] + ISV[DD_THRESHOLD_INDEX=421]; writes a single f32 penalty value to a caller-supplied output buffer. Asymmetric: zero below threshold, quadratic growth above — encodes loss aversion per `pearl_audit_unboundedness_for_implicit_asymmetry` (the symmetric SP12 v3 reward bound removed implicit behavioral asymmetry; this asymmetric quadratic restores it on the drawdown axis). 3 new ISV slots: `LAMBDA_DD_INDEX=420` (initial 1.0; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up grad-balance producer once `DD_PENALTY_GRAD_NORM_INDEX=422` accumulates observations), `DD_THRESHOLD_INDEX=421` (initial 0.05 = 5% drawdown trigger; Invariant-1 anchor — fixed structural anchor, NOT a stateful EMA), `DD_PENALTY_GRAD_NORM_INDEX=422` (stateful EMA, FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md`). 3 new `state_reset_registry` entries (`sp15_lambda_dd` rewrites the 1.0 anchor at fold boundary so the cold-start absorber re-engages on each new fold; `sp15_dd_threshold` rewrites the 0.05 anchor at fold boundary; `sp15_dd_penalty_grad_norm` resets to 0). 3 new dispatch arms in `training_loop.rs` enforce the registry-arm contract (`every_fold_and_soft_reset_entry_has_dispatch_arm` regression test). Anchor test 2.5 drawdown_de_risks (Phase 2C / Phase 3.5 paired) — green via Phase 3.5 mechanisms; this commit lands the penalty primitive. **Phase 3.3 lands kernel + launcher + 3 ISV anchor seeds + 3 registry entries + 3 dispatch arms only**; the host-side reward composer that subtracts this kernel's output from r_total is deferred to a follow-up atomic commit per `feedback_no_partial_refactor.md` (matching Phase 1.1-1.3 + 3.1-3.2 precedent — kernel + launcher verify in isolation first via the GPU oracle tests below). Touched: `crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu` (new — single kernel, single-thread/single-block, mirrors `dd_state_kernel.cu`'s literal-index pattern with `__threadfence_system()` after the write), `crates/ml/build.rs` (+1 cubin manifest entry in `kernels_with_common`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_DD_PENALTY_CUBIN` cubin slot + 1 `pub fn launch_sp15_dd_penalty` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_DD_PENALTY_CUBIN.to_vec())` and resolving `get_function("dd_penalty_kernel")` — single-thread/single-block grid mirroring `launch_sp15_dd_state` precedent + 3 ISV constructor writes (LAMBDA_DD=1.0, DD_THRESHOLD=0.05, DD_PENALTY_GRAD_NORM=0.0) appended to the SP15 Phase 3.1 anchor block), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+3 `RegistryEntry` records — `sp15_lambda_dd` rewrites 1.0 anchor; `sp15_dd_threshold` rewrites 0.05 anchor; `sp15_dd_penalty_grad_norm` Pearl A sentinel 0), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+3 dispatch arms for the 3 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+2 GPU oracle tests in `mod gpu`: `dd_penalty_quadratic_above_threshold` validates `dd=0.10, threshold=0.05, λ_dd=10 → penalty = 10 × (0.10−0.05)² = 0.025`; `dd_penalty_zero_below_threshold` validates the asymmetric structural property — `dd=0.02 < threshold=0.05` produces penalty=0 even with λ_dd=10, sentinel non-zero output buffer ensures the kernel actually overwrites with zero rather than skipping). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + registry entries + dispatch arms land atomically; reward-composer subtraction follows as its own atomic commit), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel actually computes the asymmetric quadratic; the second oracle test verifies the kernel writes a real zero rather than skipping the store), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both ISV bus and penalty output buffer), `feedback_isv_for_adaptive_bounds` (the 1.0 λ_dd anchor is a structural cold-start absorber; the 0.05 threshold is an Invariant-1 anchor — both rewritten at fold boundary so the cold-start absorber re-engages per phase precedent; runtime λ_dd will be signal-driven from the deferred grad-balance producer), `pearl_audit_unboundedness_for_implicit_asymmetry` (the asymmetric quadratic is the load-bearing pearl mechanism — symmetric clamps would erase the loss-aversion teaching this kernel encodes), `pearl_first_observation_bootstrap` (DD_PENALTY_GRAD_NORM sentinel 0 so the follow-up producer's Pearl A first-observation replacement fires on the new fold's first non-zero gradient norm), `pearl_symmetric_clamp_audit` (one-sided `fmaxf(0, …)` is structurally asymmetric BY DESIGN per spec §8.2 (3.3) — this kernel's `fmaxf(0, dd − thr)` is the asymmetry the loss-aversion teaching requires, NOT a symmetric-clamp violation; the audit pearl forbids accidentally one-sided clamps on bounded scalars, here the one-sidedness is the load-bearing semantics).