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

feat(sp15-p3.5.b+3.5.3.b): wire hold_floor (inline) + cooldown mask into experience_action_select

Browse Source 
Phase 3.5 (hold_floor_kernel) + Phase 3.5.3 (cooldown_kernel) landed
the producer + state machinery; both deferred the action-selection
consumer wiring. This task wires both atomically.

Architectural decision: hold_floor is now an INLINE __device__
computation inside experience_action_select reading ISV slots
426/427/428/429 directly. The standalone hold_floor_kernel.cu +
launch_sp15_hold_floor + HOLD_FLOOR_CUBIN are deleted — launching a
kernel to write one f32 just to read it back was unnecessary. ISV
slots + state_reset_registry entries remain; only the launch path is
removed per feedback_wire_everything_up + feedback_no_legacy_aliases.

Entropy source: per-step Shannon entropy of softmax(e_dir) computed
inline from the 4 e_dir floats already in registers (Pass 1 of the
Thompson direction selector). High entropy = uncertain policy → Hold
gets the floor lift; low entropy = confident policy → floor ≈ 0.

q_eff_dir scratch preserves e_dir for downstream consumers
(out_conviction, out_q_gaps, out_magnitude_conviction) — adding
hold_floor there would corrupt the Kelly-cap warmup floor with a
meta-confidence mask.

cooldown mask: when ISV[COOLDOWN_BARS_REMAINING=435] > 0,
action_select hard short-circuits to dir_idx = DIR_HOLD before
Pass 2 — sidesteps the temperature-blend numerics where a
finite-sentinel-on-non-Hold approach would let pure-Thompson (τ=1)
samples dominate the masked direction. Cooldown supersedes
hold_floor — when forcing Hold the floor is moot.

3 new oracle tests:
  - action_select_applies_hold_floor_inline (no cooldown)
  - action_select_forces_hold_during_cooldown
  - action_select_no_force_hold_when_cooldown_zero

Atomic per feedback_no_partial_refactor: action_select changes +
hold_floor_kernel deletion + cubin manifest update + 3 oracle tests +
audit doc all in this commit. No parallel paths, no feature flags.

Eliminates Phase 3.5 + Phase 3.5.3 deferred consumers. The
cooldown_kernel itself remains (it maintains the consecutive_losses
streak + decrements COOLDOWN_BARS_REMAINING per bar); only its
consumer is now wired.

Verified: cargo check -p ml --features cuda clean; ml lib suite
holds 946 pass / 13 fail = baseline; all 6 oracle tests pass on
RTX 3050 Ti (3 pre-existing cooldown + 3 new action_select).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-06 21:18:31 +02:00
parent d7f60d4dd7
commit f01a292f6f
6 changed files with 515 additions and 233 deletions
Download Patch File Download Diff File Expand all files Collapse all files

35
crates/ml/build.rs
View File

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

111
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

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

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

@@ -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<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(())
}
// 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`.

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

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

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

@@ -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<CudaFunction> = OnceLock::new();
fn load_action_select(stream: &Arc<CudaStream>) -> &'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<f32> {
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<f32> {
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<CudaStream>,
batch: usize,
n_atoms: i32,
q_values: &CudaSlice<f32>,
b_logits_dir: &CudaSlice<f32>,
per_sample_support: &CudaSlice<f32>,
actions_buf: &mut CudaSlice<i32>,
q_gaps_buf: &mut CudaSlice<f32>,
intent_buf: &mut CudaSlice<i32>,
conv_buf: &mut CudaSlice<f32>,
mag_conv_buf: &mut CudaSlice<f32>,
isv_dev_ptr: u64,
) -> Vec<i32> {
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::<i32>(batch).expect("alloc actions");
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conviction");
let mut mag_conv_buf = stream.alloc_zeros::<f32>(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::<i32>(batch).expect("alloc actions");
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conviction");
let mut mag_conv_buf = stream.alloc_zeros::<f32>(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::<i32>(batch).expect("alloc actions");
let mut q_gaps_buf = stream.alloc_zeros::<f32>(batch).expect("alloc q_gaps");
let mut intent_buf = stream.alloc_zeros::<i32>(batch).expect("alloc intent");
let mut conv_buf = stream.alloc_zeros::<f32>(batch).expect("alloc conviction");
let mut mag_conv_buf = stream.alloc_zeros::<f32>(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

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

File diff suppressed because one or more lines are too long