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

feat(sp15-p3.5): confidence-aware Hold floor — bounded sigmoid

Browse Source 
Per spec §8.2 (3.5) post-amendment-2 fix. hold_floor = α × σ(k × (entropy − ε₀))
added to Q_hold pre-argmax/Thompson selection. Hold becomes uncertainty
expression, not distributional default.

α (HOLD_FLOOR_ALPHA slot 426) initial 0.5 sentinel — producer kernel
updating from rolling 95th percentile of |Q_dir| (NOT running max —
outlier-ratchet vulnerable per spec second-review #6) is documented
Phase 3.5 follow-up.
k (HOLD_FLOOR_K slot 427) initial 10.0 — producer from running variance
of entropy is follow-up.
ε₀ (HOLD_FLOOR_EPS0 slot 428) initial 1.0 — producer from 75th percentile
of entropy distribution (ENTROPY_DIST_REF slot 429) is follow-up.

4 fold-reset registry entries + dispatch arms.

Per established Phase precedent: kernel + launcher land first; action-
selection wiring (add hold_floor to Q_hold pre-argmax/Thompson) deferred
to follow-up commit per feedback_no_partial_refactor.

Anchor tests: 2.1 flat_market_holds + 2.6 regime_silences (Phase 2B
contracts) — green via this teaching's action-selection wiring.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-06 16:12:17 +02:00
parent 8bfc480d92
commit 5d36f3238c
7 changed files with 347 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
crates/ml/build.rs
View File

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

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

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

65
crates/ml/src/cuda_pipeline/hold_floor_kernel.cu Normal file
View File

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

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

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

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

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

69
crates/ml/tests/sp15_phase1_oracle_tests.rs
View File

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

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

File diff suppressed because one or more lines are too long