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

feat(sp11): A2 — controller kernel + SimHash novelty buffer

Browse Source 
reward_subsystem_controller_kernel: 5 canaries → 10 outputs, true Z-score
(delta_ema/sqrt(var_ema)), sigmoid blending, weight renormalization to Σ=1,
saboteur post-clamp, curiosity permanent floor (0.2 × bound). Pearls A+D
chained on outputs per spec §3.4.1.

novelty_simhash_kernel: 42×16 random projection → 16-bit SimHash code,
1M-slot bucket count table for novelty signal `1/sqrt(1+count)`. Race-
tolerated update per feedback_no_atomicadd (under-counts bias novelty
UPWARD — safe direction).

novelty_simhash_proj_init_kernel: Philox-seeded GPU init for the
projection matrix (CPU is read-only per feedback_no_cpu_forwards).

HEALTH_DIAG `sp11_reward` line emits 10 outputs + improvement_z each
epoch. Reset registry: novelty hash table reset arm wired (closes the
A0 deferral); projection matrix is frozen at trainer init for run
lifetime, not reset.

All 20 SP11 slots populate every step. No consumer reads them yet —
training behavior unchanged from A1. 3 new GPU oracle tests pass on
RTX 3050 Ti (controller midpoint, weight renorm, saboteur clamp).

Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md §3.4 §3.5.2

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-04 02:26:05 +02:00
parent 66f5fd8f00
commit 25eba79ad5
9 changed files with 1318 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
crates/ml/build.rs
View File

@@ -564,6 +564,22 @@ fn main() {
"val_sharpe_delta_compute_kernel.cu",
"saboteur_engagement_compute_kernel.cu",
"reward_component_mag_ratio_compute_kernel.cu",
// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller +
// SimHash novelty buffer. Three new kernels in this layer:
// 1. reward_subsystem_controller — main producer (5 canaries → 10
// outputs in scratch). Chained `apply_pearls_ad_kernel` smooths
// into ISV[340..350) per spec §3.4.1 (output smoothing).
// 2. novelty_simhash_proj_init — one-shot Philox-driven init of
// the 42×16 projection matrix at trainer construct (CPU is
// read-only per `feedback_no_cpu_forwards`).
// 3. novelty_simhash — lookup + update (race-tolerated
// per `feedback_no_atomicadd`; under-counts bias novelty UP,
// the safe direction). Lookup/update launchers exist in A2 but
// are wired into the replay path in B1 — A2 just builds the
// kernels and reset registry contract end-to-end.
"reward_subsystem_controller_kernel.cu",
"novelty_simhash_proj_init_kernel.cu",
"novelty_simhash_kernel.cu",
// SP7 (2026-05-03): GPU dispatch for cql/c51 budget consumer.
// Eliminates capture-time host-branch freeze in
// `compute_adaptive_budgets`. The dispatch kernel resolves

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

@@ -553,6 +553,43 @@ static SP11_SABOTEUR_ENGAGEMENT_COMPUTE_CUBIN: &[u8] =
static SP11_MAG_RATIO_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/reward_component_mag_ratio_compute_kernel.cubin"));
/// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller. Single
/// block, 10 threads. Reads 5 canary ISV slots [350..360) (val-sharpe Δ +
/// variance, 6 mag-ratios, saboteur engagement, PnL magnitude EMA) and
/// writes 10 outputs to `scratch[SCRATCH_SP11_CONTROLLER_BASE..+10)`.
/// Downstream `apply_pearls_ad_kernel` (Pearls A+D output smoothing per
/// spec §3.4.1, n_slots=10) smooths the controller outputs into the
/// 10 contiguous component weight + controller scalar slots at
/// ISV[340..350) — 6 weights + curiosity_pressure + saboteur_mult +
/// weight_floor + curiosity_bound.
static SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/reward_subsystem_controller_kernel.cubin"));
/// SP11 Fix 39 (2026-05-04, Task A2): one-shot GPU init for the SimHash
/// projection matrix. Philox-driven (deterministic from a host-passed seed
/// derived from the SP11 salt) writes ±1 values into a 42×16 = 672-element
/// f32 buffer. Per `feedback_no_cpu_forwards.md` ("CPU is read-only"), the
/// projection matrix MUST be populated by a GPU kernel. Launched ONCE at
/// trainer construct; never re-initialised; never written from host. The
/// matrix outlives every fold (it's the hash function, not state) so it
/// has NO entry in the fold-reset registry.
static SP11_NOVELTY_SIMHASH_PROJ_INIT_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/novelty_simhash_proj_init_kernel.cubin"));
/// 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
/// non-atomically increments the bucket count (race-tolerated per
/// `feedback_no_atomicadd.md`; under-counts bias novelty UPWARD, the safe
/// direction). 1M-slot hash table + 42×16 SimHash projection.
///
/// A2 builds + registers the cubin and provides Rust launchers; the
/// production wire-up to the replay path lives in B1 — A2 just owns the
/// reset registry contract end-to-end (the novelty_hash_buf reset arm
/// lands here alongside the buffer field).
static SP11_NOVELTY_SIMHASH_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/novelty_simhash_kernel.cubin"));
/// SP7 (2026-05-03): GPU dispatch for cql/c51 budget consumer + dev-ptr
/// SAXPY/scale variants. Consolidates three kernels into one cubin so they
/// share a CUmodule with the SP7 dispatch buffer. See
@@ -1234,8 +1271,26 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
/// 7 floats for reward_component_mag_ratio (SCRATCH_SP11_MAG_RATIO_BASE=269, [269..276))
/// First 6 = component ratios → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6);
/// 7th = popart EMA mirror → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359].
/// Combined: 259 + 6 + 1 + 10 = 276 scratch slots [0..276).
pub const SP5_SCRATCH_TOTAL: usize = 276;
/// SP11 (Fix 39, 2026-05-04, A2) adds:
/// 10 floats for reward_subsystem_controller (SCRATCH_SP11_CONTROLLER_BASE=276, [276..286))
/// Layout matches ISV[340..350) one-for-one:
/// [276..282) = 6 component weights → ISV[340..346)
/// [282] = curiosity_pressure → ISV[346]
/// [283] = saboteur_intensity → ISV[347]
/// [284] = reward_weight_floor → ISV[348]
/// [285] = curiosity_bound → ISV[349]
/// Smoothed by `apply_pearls_ad_kernel` per spec §3.4.1 (Pearls A+D
/// applied to controller outputs). Slots 276..282 are contiguous in
/// scratch; the chained Pearls launch can cover them with `n_slots=6`
/// in one shot. Slots 282..286 each get a singleton Pearls launch
/// because the corresponding ISV slots [346..350) are also 4
/// contiguous singletons in the absolute ISV layout — actually they
/// ARE contiguous, but split into 1+1+1+1 launches in the launcher
/// body for explicit per-output traceability with the SP11 spec
/// table; functionally equivalent to a single n_slots=4 launch.
/// See `launch_sp11_reward_subsystem_controller` for the dispatch.
/// Combined: 259 + 6 + 1 + 10 + 10 = 286 scratch slots [0..286).
pub const SP5_SCRATCH_TOTAL: usize = 286;
/// SP5 Layer D Task D1 (rewrite, 2026-05-02): scratch index base for
/// `pnl_aggregation_update`.
@@ -1361,6 +1416,22 @@ pub const SCRATCH_SP11_VAL_SHARPE_DELTA_BASE: usize = 266; // [266..268) → I
pub const SCRATCH_SP11_SABOTEUR_ENGAGEMENT: usize = 268; // [268..269) → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358]
pub const SCRATCH_SP11_MAG_RATIO_BASE: usize = 269; // [269..276) — first 6 → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6); 7th → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
/// SP11 (Fix 39, 2026-05-04, Task A2): scratch slots for the
/// `reward_subsystem_controller` producer. 10 contiguous floats
/// at [276..286) match the ISV[340..350) layout one-for-one:
/// [276..282) = 6 component weights (popart, cf, trail, micro,
/// opp_cost, bonus) → ISV[REWARD_*_WEIGHT_INDEX..+6)
/// [282] = curiosity_pressure → ISV[CURIOSITY_PRESSURE_INDEX=346]
/// [283] = saboteur_intensity_mult → ISV[SABOTEUR_INTENSITY_MULT_INDEX=347]
/// [284] = reward_weight_floor → ISV[REWARD_WEIGHT_FLOOR_INDEX=348]
/// [285] = curiosity_bound → ISV[CURIOSITY_BOUND_INDEX=349]
///
/// The chained `apply_pearls_ad_kernel` (Pearls A+D output smoothing per
/// spec §3.4.1) reads the entire 10-slot block in one launch with
/// `n_slots=10` since both scratch and ISV layouts are contiguous in
/// matched-stride order.
pub const SCRATCH_SP11_CONTROLLER_BASE: usize = 276; // [276..286) → ISV[340..350)
/// SP5 Task A7: scratch index base for pearl_8_trail_update trail_dist[4] output block.
/// Slots [199..203): per-direction trail-stop distance (Short=0, Hold=1, Long=2, Flat=3).
/// Written by `pearl_8_trail_update`; consumed by apply_pearls_ad_kernel →
@@ -4742,6 +4813,56 @@ pub struct GpuDqnTrainer {
/// (alloc_episodes × alloc_timesteps). Cached on the trainer because
/// the engagement kernel takes it as an i32 arg.
pub(crate) sp11_saboteur_delta_reward_len: usize,
/// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller
/// kernel. Single-block, 10-thread producer reading 5 canary ISV slots
/// [350..360) and writing 10 floats to scratch[SCRATCH_SP11_CONTROLLER_BASE..+10).
/// Followed by 1 chained `apply_pearls_ad_kernel` launch (n_slots=10)
/// → ISV[340..350) (6 component weights + 4 controller scalars).
/// Loaded from `reward_subsystem_controller_kernel.cubin`.
sp11_reward_subsystem_controller_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, Task A2): one-shot GPU init kernel for the
/// 42×16 SimHash projection matrix. Philox-driven, deterministic from a
/// host-passed seed derived from `SP11_PROJ_SEED_SALT`. Launched ONCE
/// at trainer construct (see the constructor's allocation block) and
/// never re-launched — the projection matrix is frozen for the run
/// lifetime. Loaded from `novelty_simhash_proj_init_kernel.cubin`.
sp11_novelty_simhash_proj_init_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty lookup kernel.
/// Reads (state, action) bucket count from the 1M-slot hash table and
/// returns `1/sqrt(1+count)` ∈ [0, 1] per replay sample. Loaded from
/// `novelty_simhash_kernel.cubin`.
///
/// **A2 builds + registers; B1 wires into the replay path.** The
/// launcher `launch_sp11_novelty_simhash_lookup` is provided in A2 but
/// not invoked from `training_loop.rs` yet — the replay-time consumer
/// migration lands atomically in Layer B per
/// `feedback_no_partial_refactor.md`.
sp11_novelty_simhash_lookup_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty update kernel.
/// Race-tolerated non-atomic bucket increment per
/// `feedback_no_atomicadd.md`; under-counts bias novelty UP (safe).
/// **A2 builds + registers; B1 wires into the replay path.** Loaded
/// from `novelty_simhash_kernel.cubin` (same cubin as the lookup
/// kernel — both share the SimHash code-derivation device-inline).
sp11_novelty_simhash_update_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, Task A2): novelty visit-count hash table.
/// 1M slots × f32 = 4 MB mapped-pinned. Bucket index =
/// `(action_id × 65536 + simhash_code) & TABLE_MASK`. Reset at fold
/// boundary via the `sp11_novelty_hash` registry entry's dispatch arm.
/// The buffer is mapped-pinned even though the lookup/update kernels
/// are pure GPU paths — the host-side reset is a `host_slice_mut().fill(0.0)`
/// (CPU writing a literal value, NOT compute) per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write rule.
pub(crate) novelty_hash_buf: super::mapped_pinned::MappedF32Buffer,
/// SP11 Fix 39 (2026-05-04, Task A2): SimHash projection matrix.
/// 42 × 16 = 672 f32 mapped-pinned. Populated ONCE at trainer
/// construct by `launch_novelty_simhash_proj_init` (Philox-driven from
/// `SP11_PROJ_SEED_SALT` per `sp11_isv_slots.rs`). Never re-initialised;
/// never written from host. Per `feedback_no_cpu_forwards.md` ("CPU is
/// read-only"), the matrix MUST be populated on-device. The matrix
/// outlives every fold (it's the hash function, not state) — it has
/// NO entry in the fold-reset registry.
pub(crate) novelty_simhash_proj: super::mapped_pinned::MappedF32Buffer,
/// SP7 (2026-05-03): GPU dispatch kernel for the cql/c51 budget consumer.
/// 8 threads (2 heads × 4 branches), single block. Reads ISV activation
/// flag + controller budget per (head, branch); resolves
@@ -12661,6 +12782,222 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP11 Fix 39 (2026-05-04, Task A2): launch the reward-subsystem
/// controller producer + chained Pearls A+D output smoothing.
///
/// Single-block, 10-thread producer reads 5 canary slots from ISV
/// [350..360) (val-sharpe Δ + variance, 6 mag-ratios, saboteur
/// engagement, PnL magnitude EMA) and writes 10 outputs to
/// `scratch[SCRATCH_SP11_CONTROLLER_BASE..+10)`. A single chained
/// `apply_pearls_ad_kernel` launch with `n_slots=10` smooths into
/// ISV[REWARD_POPART_WEIGHT_INDEX=340..CURIOSITY_BOUND_INDEX+1=350) —
/// the controller's full output block. Both scratch and ISV layouts
/// are contiguous and matched-stride, so the n_slots=10 single-launch
/// path is correct (vs splitting into two contiguous-block launches
/// the way `launch_sp11_mag_ratio_compute` does for its 6+1 split).
///
/// Per spec §3.4.1 (output smoothing) the smoothed outputs are what
/// downstream consumers (B1) read from ISV — not the raw scratch
/// values. Pearl A's sentinel-bootstrap fires on the first per-fold
/// observation; Pearl D's Wiener-α adapts smoothing strength to each
/// output's own noise level (curiosity may be smoother than weights
/// etc., adaptively).
///
/// Pre-conditions (must be satisfied before this launch in the same
/// stream): all 5 canary slots [350..360) are populated. In
/// `training_loop.rs` the call site fires AFTER A1's three canary
/// launches and BEFORE the HEALTH_DIAG emit so the controller reads
/// fresh canary state and the operator sees current outputs.
pub(crate) fn launch_sp11_reward_subsystem_controller(&self) -> Result<(), MLError> {
use crate::cuda_pipeline::sp4_wiener_ema::{launch_apply_pearls, ALPHA_META};
use crate::cuda_pipeline::sp5_isv_slots::SP5_SLOT_BASE;
use crate::cuda_pipeline::sp11_isv_slots::{
REWARD_POPART_WEIGHT_INDEX,
VAL_SHARPE_DELTA_EMA_INDEX, VAL_SHARPE_VAR_EMA_INDEX,
REWARD_COMPONENT_MAG_RATIO_BASE,
SABOTEUR_ENGAGEMENT_RATE_INDEX,
PNL_REWARD_MAGNITUDE_EMA_INDEX,
};
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_sp11_reward_subsystem_controller: isv_signals_dev_ptr must be allocated");
let isv_dev = self.isv_signals_dev_ptr;
let scratch_dev = self.producer_step_scratch_buf.dev_ptr;
let wiener_dev = self.wiener_state_buf.dev_ptr;
// Producer writes scratch[SCRATCH_SP11_CONTROLLER_BASE..+10).
// We pass the offset device pointer as the scratch arg so the
// kernel's view starts at scratch_out[0] = scratch[base] — same
// pointer-arithmetic pattern as A1's val_sharpe_delta launcher.
let scratch_offset_bytes = (SCRATCH_SP11_CONTROLLER_BASE as u64)
* std::mem::size_of::<f32>() as u64;
let scratch_out_dev: u64 = scratch_dev + scratch_offset_bytes;
let delta_ema_slot_i32: i32 = VAL_SHARPE_DELTA_EMA_INDEX as i32;
let var_ema_slot_i32: i32 = VAL_SHARPE_VAR_EMA_INDEX as i32;
let mag_ratio_base_slot_i32: i32 = REWARD_COMPONENT_MAG_RATIO_BASE as i32;
let saboteur_engagement_slot_i32: i32 = SABOTEUR_ENGAGEMENT_RATE_INDEX as i32;
let pnl_mag_ema_slot_i32: i32 = PNL_REWARD_MAGNITUDE_EMA_INDEX as i32;
// Step 1: producer kernel — single block, 10 threads (one per output).
unsafe {
self.stream
.launch_builder(&self.sp11_reward_subsystem_controller_kernel)
.arg(&isv_dev)
.arg(&scratch_out_dev)
.arg(&delta_ema_slot_i32)
.arg(&var_ema_slot_i32)
.arg(&mag_ratio_base_slot_i32)
.arg(&saboteur_engagement_slot_i32)
.arg(&pnl_mag_ema_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (10, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("sp11 reward_subsystem_controller: {e}")))?;
}
// Step 2: chained apply_pearls_ad_kernel (n_slots=10) per spec
// §3.4.1. Both scratch [SCRATCH_SP11_CONTROLLER_BASE..+10) and
// ISV [REWARD_POPART_WEIGHT_INDEX..+10) are contiguous and
// matched-stride, so a single n_slots=10 launch covers them.
//
// Wiener-state offset for the first output slot:
// wiener_off = SP4_PRODUCER_COUNT*3 + (isv_slot - SP5_SLOT_BASE)*3
// The applicator advances by 3 floats per slot internally.
let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3;
let isv_idx_base_i32 = REWARD_POPART_WEIGHT_INDEX as i32;
let scratch_idx_base_i32: i32 = SCRATCH_SP11_CONTROLLER_BASE as i32;
let wiener_off = base_wiener_offset
+ (isv_idx_base_i32 - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, scratch_idx_base_i32,
isv_dev, isv_idx_base_i32,
wiener_dev, wiener_off,
10,
ALPHA_META,
)?;
}
Ok(())
}
/// SP11 Fix 39 (2026-05-04, Task A2): launch the SimHash novelty
/// LOOKUP kernel — reads (state, action) bucket count from
/// `novelty_hash_buf` and writes `1/sqrt(1+count)` ∈ [0, 1] per
/// sample to the caller-supplied output buffer.
///
/// **A2 builds + registers; B1 wires into the replay path.** This
/// launcher is provided in A2 alongside the kernel registration so
/// that B1 (consumer migration) is a pure consumer change with no
/// hot-side wiring. Per `feedback_wire_everything_up.md`'s
/// "build OR don't build" rule, this scaffolding is sanctioned —
/// the production wire-up to the replay path lands atomically in
/// Layer B per `feedback_no_partial_refactor.md`.
///
/// Inputs (all caller-owned device buffers; no copies):
/// - states : `[batch_size × 42]` f32 (replay-batch states)
/// - actions : `[batch_size]` i32 (replay-batch actions)
/// - novelty_out : `[batch_size]` f32 (per-sample novelty signal)
/// Reads `self.novelty_simhash_proj` (frozen at trainer init) and
/// `self.novelty_hash_buf` (fold-reset registry-managed).
#[allow(dead_code)] // Wired in B1; A2 just makes the launcher available.
pub(crate) fn launch_sp11_novelty_simhash_lookup(
&self,
states_dev: u64,
actions_dev: u64,
novelty_out_dev: u64,
batch_size: usize,
) -> Result<(), MLError> {
debug_assert!(states_dev != 0,
"launch_sp11_novelty_simhash_lookup: states_dev must be a valid device pointer");
debug_assert!(actions_dev != 0,
"launch_sp11_novelty_simhash_lookup: actions_dev must be a valid device pointer");
debug_assert!(novelty_out_dev != 0,
"launch_sp11_novelty_simhash_lookup: novelty_out_dev must be a valid device pointer");
let proj_dev = self.novelty_simhash_proj.dev_ptr;
let table_dev = self.novelty_hash_buf.dev_ptr;
let bs_i32: i32 = i32::try_from(batch_size).unwrap_or(i32::MAX);
let block_dim: u32 = 256;
let grid_dim: u32 = (batch_size as u32 + block_dim - 1) / block_dim;
unsafe {
self.stream
.launch_builder(&self.sp11_novelty_simhash_lookup_kernel)
.arg(&states_dev)
.arg(&actions_dev)
.arg(&proj_dev)
.arg(&table_dev)
.arg(&novelty_out_dev)
.arg(&bs_i32)
.launch(LaunchConfig {
grid_dim: (grid_dim, 1, 1),
block_dim: (block_dim, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("sp11 novelty_simhash_lookup: {e}")))?;
}
Ok(())
}
/// SP11 Fix 39 (2026-05-04, Task A2): launch the SimHash novelty
/// UPDATE kernel — non-atomically increments the (state, action)
/// bucket count in `novelty_hash_buf` per sample.
///
/// **A2 builds + registers; B1 wires into the replay path.**
///
/// Race-tolerated update per `feedback_no_atomicadd.md`: concurrent
/// writers to the same bucket race; last-writer-wins loses
/// increments. This biases counts DOWNWARD, which biases
/// `novelty = 1/sqrt(1+count)` UPWARD — the SAFE direction (policy
/// over-explores rather than masking visited state-actions as
/// novel). NEVER replace with atomicAdd; the safety argument is
/// preserved in the kernel header.
///
/// Spec §3.5.2 explicitly accepts this trade-off as the price of
/// avoiding atomicAdd.
#[allow(dead_code)] // Wired in B1; A2 just makes the launcher available.
pub(crate) fn launch_sp11_novelty_simhash_update(
&self,
states_dev: u64,
actions_dev: u64,
batch_size: usize,
) -> Result<(), MLError> {
debug_assert!(states_dev != 0,
"launch_sp11_novelty_simhash_update: states_dev must be a valid device pointer");
debug_assert!(actions_dev != 0,
"launch_sp11_novelty_simhash_update: actions_dev must be a valid device pointer");
let proj_dev = self.novelty_simhash_proj.dev_ptr;
let table_dev = self.novelty_hash_buf.dev_ptr;
let bs_i32: i32 = i32::try_from(batch_size).unwrap_or(i32::MAX);
let block_dim: u32 = 256;
let grid_dim: u32 = (batch_size as u32 + block_dim - 1) / block_dim;
unsafe {
self.stream
.launch_builder(&self.sp11_novelty_simhash_update_kernel)
.arg(&states_dev)
.arg(&actions_dev)
.arg(&proj_dev)
.arg(&table_dev)
.arg(&bs_i32)
.launch(LaunchConfig {
grid_dim: (grid_dim, 1, 1),
block_dim: (block_dim, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("sp11 novelty_simhash_update: {e}")))?;
}
Ok(())
}
/// SP7 (2026-05-03): launch the GPU dispatch kernel that resolves
/// `(active >= 0.5) ? max(controller, EPS_DIV) : bootstrap` for every
/// (head ∈ {CQL, C51}, branch ∈ [0..4)) on-device.
@@ -15999,6 +16336,111 @@ impl GpuDqnTrainer {
format!("SP11 val_sharpe_history_pinned alloc: {e}")
))?;
// SP11 Fix 39 (2026-05-04, Task A2): load reward-subsystem
// controller + SimHash novelty kernels.
//
// Controller is the main A2 producer: 5 canaries → 10 outputs in
// scratch[SCRATCH_SP11_CONTROLLER_BASE..+10) → Pearls A+D output
// smoothing → ISV[340..350). SimHash kernels (proj-init, lookup,
// update) build the novelty buffer infrastructure that B1 wires
// into the replay path.
let sp11_reward_subsystem_controller_kernel = {
let module = stream.context()
.load_cubin(SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp11 reward_subsystem_controller cubin load: {e}")))?;
module.load_function("reward_subsystem_controller_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 reward_subsystem_controller load: {e}")))?
};
let sp11_novelty_simhash_proj_init_kernel = {
let module = stream.context()
.load_cubin(SP11_NOVELTY_SIMHASH_PROJ_INIT_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp11 novelty_simhash_proj_init cubin load: {e}")))?;
module.load_function("novelty_simhash_proj_init_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 novelty_simhash_proj_init load: {e}")))?
};
// The lookup + update kernels share one cubin (one CUmodule).
let sp11_novelty_simhash_module = stream.context()
.load_cubin(SP11_NOVELTY_SIMHASH_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp11 novelty_simhash cubin load: {e}")))?;
let sp11_novelty_simhash_lookup_kernel = sp11_novelty_simhash_module
.load_function("novelty_simhash_lookup_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 novelty_simhash_lookup load: {e}")))?;
let sp11_novelty_simhash_update_kernel = sp11_novelty_simhash_module
.load_function("novelty_simhash_update_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 novelty_simhash_update load: {e}")))?;
// SP11 Fix 39 (2026-05-04, A2): allocate the novelty visit-count
// hash table (1M slots × f32 = 4 MB mapped-pinned). Constructor-
// zero-initialised by `MappedF32Buffer::new` (cuMemHostAlloc is
// followed by `write_bytes(0, len)` per the buffer's invariants
// doc) which is the desired cold-start state for the novelty
// signal: every (state, action) bucket reads count=0 →
// novelty=1/sqrt(1)=1.0. The fold-reset arm `sp11_novelty_hash`
// re-zeros the buffer at fold boundaries.
let novelty_hash_buf = unsafe {
super::mapped_pinned::MappedF32Buffer::new(1usize << 20)
}.map_err(|e| MLError::ModelError(
format!("SP11 novelty_hash_buf alloc: {e}")
))?;
// SP11 Fix 39 (2026-05-04, A2): allocate the SimHash projection
// matrix (42 × 16 = 672 f32 mapped-pinned) and immediately launch
// the GPU init kernel to populate it with deterministic ±1 values
// derived from `SP11_PROJ_SEED_SALT`. Per `feedback_no_cpu_forwards.md`
// ("CPU is read-only") + `feedback_no_cpu_compute_strict.md`, the
// projection MUST be populated by a GPU kernel. The buffer is
// mapped-pinned so the kernel can write through `dev_ptr` without
// a separate device alloc — the host never reads or writes the
// matrix after this point.
//
// Per `feedback_wire_everything_up.md`: the init kernel is wired
// to the constructor that allocates the buffer. No orphan
// additions. The matrix is frozen for the run lifetime; no
// fold-reset entry exists for it (it's the hash function, not
// state).
let novelty_simhash_proj = unsafe {
super::mapped_pinned::MappedF32Buffer::new(42 * 16)
}.map_err(|e| MLError::ModelError(
format!("SP11 novelty_simhash_proj alloc: {e}")
))?;
{
use crate::cuda_pipeline::sp11_isv_slots::SP11_PROJ_SEED_SALT;
// Salt-only seed: the project does not have a global config.seed;
// the SP11 spec just requires a deterministic seed reproducible
// across runs. The salt itself is a stable 64-bit constant in
// `sp11_isv_slots.rs` so every run gets the same projection
// matrix → reproducible novelty hashing across reruns.
let seed_u64: u64 = SP11_PROJ_SEED_SALT;
// Philox-uniform takes an i32 seed — fold the high + low halves.
let seed_i32: i32 = ((seed_u64 ^ (seed_u64 >> 32)) as u32) as i32;
let proj_dev = novelty_simhash_proj.dev_ptr;
let n_floats: i32 = (42 * 16) as i32;
let block_dim: u32 = 32;
let n_blocks: u32 = (672 + block_dim - 1) / block_dim; // 21 blocks
unsafe {
stream
.launch_builder(&sp11_novelty_simhash_proj_init_kernel)
.arg(&proj_dev)
.arg(&seed_i32)
.arg(&n_floats)
.launch(LaunchConfig {
grid_dim: (n_blocks, 1, 1),
block_dim: (block_dim, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(
format!("sp11 novelty_simhash_proj_init launch: {e}")
))?;
}
// Sync so the matrix is fully populated before any subsequent
// construct-time work (or the first lookup launch in B1) reads
// it. One-shot init at construct → no hot-path penalty.
stream.synchronize()
.map_err(|e| MLError::ModelError(
format!("sp11 novelty_simhash_proj_init sync: {e}")
))?;
}
// SP7 (2026-05-03): load consume_lb_budget_kernel.cubin — three kernels
// (lb_budget_dispatch + dqn_scale_f32_dev_ptr_kernel +
// dqn_saxpy_f32_dev_ptr_kernel) sharing one CUmodule. Separate from
@@ -19277,6 +19719,17 @@ impl GpuDqnTrainer {
// is constructed (see training_loop.rs init path).
sp11_saboteur_delta_reward_dev_ptr: 0,
sp11_saboteur_delta_reward_len: 0,
// SP11 Fix 39 (Task A2): controller + SimHash novelty kernels
// and their backing buffers. Projection matrix was populated
// on-device above by `launch_novelty_simhash_proj_init`; hash
// table is constructor-zero and reset at fold boundary via
// the `sp11_novelty_hash` registry arm.
sp11_reward_subsystem_controller_kernel,
sp11_novelty_simhash_proj_init_kernel,
sp11_novelty_simhash_lookup_kernel,
sp11_novelty_simhash_update_kernel,
novelty_hash_buf,
novelty_simhash_proj,
lb_budget_dispatch_kernel,
saxpy_f32_dev_ptr_aux,
scale_f32_dev_ptr_aux,

104
crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu Normal file
View File

@@ -0,0 +1,104 @@
// crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu
//
// SP11 Fix 39 Task A2 (2026-05-04) — SimHash novelty signal.
//
// Two functions sharing the simhash_code device-inline helper:
// 1. novelty_simhash_lookup_kernel — reads bucket count, returns
// 1/sqrt(1+count) ∈ [0, 1]
// 2. novelty_simhash_update_kernel — increments bucket count after replay
// sample (race-tolerated, no atomicAdd)
//
// State space is continuous 42-dim. We project state through a fixed
// 42×16 random projection matrix (init'd once at trainer construct by
// `novelty_simhash_proj_init_kernel`), sign-quantize to a 16-bit code, and
// bucket by `(action_id × 65536 + code) & TABLE_MASK`. TABLE_SIZE = 1<<20 =
// 1M slots; linear-probe collisions are tolerated and bias novelty UPWARD
// (the safe direction).
//
// Race-tolerance contract (forbids atomicAdd per `feedback_no_atomicadd.md`):
// concurrent threads writing to the same bucket race; last-writer-wins
// loses increments. This biases counts DOWNWARD, which biases
// `novelty = 1/sqrt(1+count)` UPWARD. Over-novelty = the policy explores
// MORE than strictly required, never LESS than required → safe direction
// (an already-visited state-action will not be masked as novel by under-
// counting; it just gets revisited with slightly elevated curiosity until
// the next race-free update lands a count). Spec §3.5.2 explicitly accepts
// this trade-off as the price of avoiding atomicAdd.
//
// These kernels are NOT wired into training_loop.rs in A2 — they're called
// at REPLAY time in B1. A2 just builds + registers them so B1 wiring is a
// pure consumer change. Per `feedback_wire_everything_up.md` the launchers
// land in this commit alongside the kernels (so the cubin and Rust handles
// are paired) but the production wire-up to the replay path is B1's scope.
//
// Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md §3.5.2
// Plan: docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md A2.2 Step 7
#include <cuda_runtime.h>
#include <math.h>
/* 1M-slot hash table; index mask must be (TABLE_SIZE - 1) which equals the
* low-20-bits mask. */
#define SP11_NOVELTY_TABLE_SIZE (1u << 20)
#define SP11_NOVELTY_TABLE_MASK ((1u << 20) - 1u)
/* Locality-sensitive 16-bit SimHash from a 42-dim state row through a
* 42×16 projection matrix. Sign quantization: bit b is set iff dot(state,
* proj[:, b]) > 0. Returns a u32 with the high 16 bits zero. */
__device__ __forceinline__ unsigned int simhash_code_42_16(
const float* __restrict__ state, /* [42] */
const float* __restrict__ proj_matrix) /* [42 × 16] row-major */
{
unsigned int code = 0u;
for (int b = 0; b < 16; b++) {
float dot = 0.0f;
#pragma unroll 8
for (int d = 0; d < 42; d++) {
dot += state[d] * proj_matrix[d * 16 + b];
}
if (dot > 0.0f) code |= (1u << b);
}
return code;
}
extern "C" __global__ void novelty_simhash_lookup_kernel(
const float* __restrict__ states, /* [batch_size × 42] */
const int* __restrict__ actions, /* [batch_size] */
const float* __restrict__ proj_matrix, /* [42 × 16] */
const float* __restrict__ hash_table, /* [TABLE_SIZE] */
float* __restrict__ novelty_out, /* [batch_size] */
const int batch_size)
{
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= batch_size) return;
const unsigned int code = simhash_code_42_16(&states[i * 42], proj_matrix);
/* Cast action to u32 first so the multiplication is unsigned; negative
* action sentinels (e.g. -1) would otherwise wrap to a huge bucket
* index. action ∈ [0, 108) in production; defensive cast costs nothing. */
const unsigned int act = (unsigned int)actions[i];
const unsigned int bucket = (act * 65536u + code) & SP11_NOVELTY_TABLE_MASK;
const float count = hash_table[bucket];
novelty_out[i] = 1.0f / sqrtf(1.0f + count);
}
extern "C" __global__ void novelty_simhash_update_kernel(
const float* __restrict__ states,
const int* __restrict__ actions,
const float* __restrict__ proj_matrix,
float* __restrict__ hash_table, /* [TABLE_SIZE] in/out */
const int batch_size)
{
/* Race-tolerated update — see header for the safety argument. The
* under-count → novelty-up direction is the safe one. */
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= batch_size) return;
const unsigned int code = simhash_code_42_16(&states[i * 42], proj_matrix);
const unsigned int act = (unsigned int)actions[i];
const unsigned int bucket = (act * 65536u + code) & SP11_NOVELTY_TABLE_MASK;
/* Non-atomic RMW: read-modify-write loses increments under races, which
* biases counts down → novelty up (safe). NEVER replace this with
* atomicAdd; see `feedback_no_atomicadd.md`. */
hash_table[bucket] = hash_table[bucket] + 1.0f;
}

50
crates/ml/src/cuda_pipeline/novelty_simhash_proj_init_kernel.cu Normal file
View File

@@ -0,0 +1,50 @@
// crates/ml/src/cuda_pipeline/novelty_simhash_proj_init_kernel.cu
//
// SP11 Fix 39 Task A2 (2026-05-04) — one-shot GPU init for the SimHash
// projection matrix.
//
// Per `feedback_no_cpu_forwards.md` ("CPU is read-only") and
// `feedback_no_cpu_compute_strict.md`, the projection matrix MUST be
// populated by a GPU kernel using the project's deterministic Philox-like
// hash (see `philox_uniform` in `common_device_functions.cuh`). Host-side
// `StdRng` writes through `host_slice_mut()` would violate the rule.
//
// The matrix is 42 × 16 = 672 floats with sign-quantized entries ∈ {-1, +1}.
// Generated once at trainer construct from a u64 seed (config-derived +
// SP11_PROJ_SEED_SALT salt) and never re-initialised; the same seed always
// produces the same projection so reproducible runs share their hash
// function.
//
// `philox_uniform(id, counter, seed)` returns a uniform float in [0, 1)
// derived from a stateless Philox-like hash; we map the low bit of an
// integer derivation to the {-1, +1} sign. Concretely we read
// `philox_uniform(i, 0, seed) * (1 << 24)` as a scaled u32, take the low
// bit, and emit ±1.
//
// Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md §3.5.2
// Plan: docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md A2.2 Step 6
#include <cuda_runtime.h>
extern "C" __global__ void novelty_simhash_proj_init_kernel(
/* Output: 42 × 16 = 672 floats. */
float* __restrict__ proj,
/* Seed (host-passed, derived from config + SP11_PROJ_SEED_SALT). */
const int seed,
/* Element count (42 × 16 = 672). Passed explicitly so the kernel binary
* is independent of state_dim/code_bits choices. */
const int n_floats)
{
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= n_floats) return;
/* Derive a deterministic uniform in [0,1) from (i, seed). Map to {-1, +1}
* via the low bit of the scaled-u32 representation.
*
* The Philox-like hash in `common_device_functions.cuh` mixes (id,
* counter, seed) so different proj matrices (different seeds) get
* statistically independent ±1 patterns. Same seed always reproduces. */
const float u = philox_uniform(i, /*counter=*/0, seed);
const unsigned int u_scaled = (unsigned int)(u * 16777216.0f); // 2^24 = uniform's bit width
proj[i] = (u_scaled & 1u) ? 1.0f : -1.0f;
}

211
crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu Normal file
View File

@@ -0,0 +1,211 @@
// crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu
//
// SP11 Fix 39 Task A2 (2026-05-04) — Reward subsystem controller.
//
// Single block, 10 threads. Reads 5 canary ISV slots [350..360) (val-sharpe
// trend + variance, 6 per-component magnitude ratios, saboteur engagement,
// PnL magnitude EMA) and writes 10 outputs to scratch[0..10) which a chained
// `apply_pearls_ad_kernel` then smooths into ISV[340..350) (6 component
// weights + 4 controller scalars).
//
// Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md §3.4
// Plan: docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md A2.1
//
// Critical math (per spec §3.4):
// - True Z-score: improvement_z = delta_ema / sqrtf(var_ema + EPS_DIV²)
// using `fmaxf(sqrtf(delta_var), EPS_DIV)` — sqrt of variance, not the
// variance directly. The bound z to [-1,+1] form would prevent sigmoid
// saturation, collapsing the controller's effective range.
// - improving = sigmoid(dz); stagnant_or_worse = 1 - improving.
// - Per-component blend (winner × improving + diversifier × stagnant)
// pre-floor, then floor-clamped, then RENORMALIZED to Σ=1. Renormalization
// is non-negotiable: PopArt's reward-magnitude EMA assumes a consistent
// reward-scale; sum(weights) drifting > 1 inflates running |reward|.
// - Curiosity uses the permanent-floor pattern (max(), NOT blend):
// `pressure = max(stagnant × bound, CURIOSITY_PERMANENT_FRACTION × bound)`.
// Per `pearl_blend_formulas_must_have_permanent_floor.md` — blend
// deadlocks when real signal hits 0 at maturity=1; max() form keeps
// permanent minimum.
// - Saboteur: post-clamp is required so the engagement floor (ENGAGEMENT_FLOOR
// = 0.1) doesn't silently pull the modulated output below SABOTEUR_MIN.
//
// Invariant-1 anchors (rate-limiters, NOT regime thresholds — every adaptive
// bound otherwise lives on ISV per `feedback_isv_for_adaptive_bounds.md`):
// the seven `const float` values are documented in the SP11 spec §3.4.2
// Invariant-1 carve-out table; values must remain in lockstep with the
// Rust-side constants in `crate::cuda_pipeline::sp11_isv_slots`. The kernel
// can't import the Rust consts so it re-declares them as device-side `const
// float` locals.
#include <cuda_runtime.h>
#include <math.h>
extern "C" __global__ void reward_subsystem_controller_kernel(
/* ISV signal bus (read-only). */
const float* __restrict__ isv,
/* Producer scratch buffer (write-only): 10 floats — first 6 are the
* component weights (popart/cf/trail/micro/opp_cost/bonus), then
* curiosity_pressure, saboteur_intensity_mult, weight_floor,
* curiosity_bound. */
float* __restrict__ scratch_out,
/* Slot indices (i32, threaded as kernel args so the same kernel binary
* works regardless of which absolute ISV slot the canary lives in;
* matches the SP9 producer-launcher pattern at
* `intent_eval_divergence_compute_kernel`). */
const int delta_ema_slot, // VAL_SHARPE_DELTA_EMA_INDEX
const int var_ema_slot, // VAL_SHARPE_VAR_EMA_INDEX
const int mag_ratio_base_slot, // REWARD_COMPONENT_MAG_RATIO_BASE
const int saboteur_engagement_slot, // SABOTEUR_ENGAGEMENT_RATE_INDEX
const int pnl_mag_ema_slot) // PNL_REWARD_MAGNITUDE_EMA_INDEX
{
if (blockIdx.x != 0) return;
if (threadIdx.x >= 10) return;
/* Invariant-1 anchors — re-declared from sp11_isv_slots.rs. Keep in
* lockstep with the Rust-side `SP11_*` constants. */
const float EPS_DIV = 1e-6f;
const float SABOTEUR_MIN = 0.5f;
const float SABOTEUR_MAX = 2.0f;
const float WEIGHT_HARD_FLOOR = 0.01f;
const float CURIOSITY_PERMANENT_FRACTION = 0.2f;
const float ENGAGEMENT_FLOOR = 0.1f;
const float CURIOSITY_BOUND_FRACTION = 0.3f;
const float WEIGHT_FLOOR_FRACTION = 0.5f;
/* Block-shared scratch — block-loaded once by threads 0..5, then read by
* all 10 threads after a single __syncthreads barrier. */
__shared__ float ratios[6];
__shared__ float improving;
__shared__ float stagnant_or_worse;
__shared__ float adaptive_floor;
__shared__ float blends[6];
__shared__ float blend_sum;
__shared__ float curiosity_bound;
/* Load 6 mag-ratio canaries into shared mem (one read per slot, vs 6
* redundant ISV reads across the 10 threads). */
if (threadIdx.x < 6) {
ratios[threadIdx.x] = isv[mag_ratio_base_slot + threadIdx.x];
}
__syncthreads();
/* Thread 0: compute improving, stagnant_or_worse, adaptive_floor,
* curiosity_bound — values consumed by every per-output thread below. */
if (threadIdx.x == 0) {
const float delta_ema = isv[delta_ema_slot];
const float delta_var = isv[var_ema_slot];
/* True Z-score: dz = mean Δ / std-of-residuals.
* sqrtf(delta_var) is the standard deviation; fmaxf with EPS_DIV
* guards the (numerical) zero-variance case. The result is unbounded
* so sigmoid can saturate at the regime extremes — collapsing the
* controller's range to [0.27, 0.73] would defeat the design. */
const float dz = delta_ema / fmaxf(sqrtf(delta_var), EPS_DIV);
improving = 1.0f / (1.0f + expf(-dz));
stagnant_or_worse = 1.0f - improving;
/* Adaptive component-weight floor: signal-driven from the worst-
* suppressed component's mag-ratio. WEIGHT_HARD_FLOOR=0.01 is the
* absolute numerical floor (Invariant 1). The fraction
* WEIGHT_FLOOR_FRACTION=0.5 is the rate-limiter on how aggressively
* the floor tracks min_grad_ratio. */
float min_grad_ratio = ratios[0];
for (int c = 1; c < 6; c++) {
min_grad_ratio = fminf(min_grad_ratio, ratios[c]);
}
adaptive_floor = fmaxf(WEIGHT_HARD_FLOOR,
WEIGHT_FLOOR_FRACTION * min_grad_ratio);
/* Curiosity bound = signal-relative cap on the per-step exploration
* bonus — keeps curiosity scaled to the run's PnL magnitude rather
* than a tuned constant. */
curiosity_bound = isv[pnl_mag_ema_slot] * CURIOSITY_BOUND_FRACTION;
}
__syncthreads();
/* Threads 0..5: compute per-component blend (pre-renormalization).
* Each thread reads its own ratios[i] from shared mem; no cross-thread
* data hazard because writes go to per-thread blends[i]. */
if (threadIdx.x < 6) {
const float winner = ratios[threadIdx.x];
const float diversifier = (1.0f - winner) / 5.0f;
float w = improving * winner + stagnant_or_worse * diversifier;
/* Floor enforcement before renorm — preserves Σ ≥ 6 × adaptive_floor
* lower bound that the renorm can normalize against without producing
* NaN. The 1.0 upper clamp is defensive (per-component blend should
* never exceed 1, but algebraic edge cases at extreme z values are
* possible). */
w = fmaxf(adaptive_floor, fminf(1.0f, w));
blends[threadIdx.x] = w;
}
__syncthreads();
/* Thread 0: reduce blend_sum once before the renorm pass. Single-threaded
* because the reduction is trivially small (6 floats) and avoids the
* synchronization overhead of a tree-reduce. */
if (threadIdx.x == 0) {
blend_sum = 0.0f;
for (int c = 0; c < 6; c++) blend_sum += blends[c];
}
__syncthreads();
/* Threads 0..5: renormalize so Σweights = 1 exactly. The fmaxf with
* EPS_DIV guards the (algebraically impossible) all-floor-zero case. */
if (threadIdx.x < 6) {
scratch_out[threadIdx.x] = blends[threadIdx.x]
/ fmaxf(blend_sum, EPS_DIV);
}
/* Thread 6: curiosity_pressure — permanent-floor pattern (max(), NOT
* blend). Per `pearl_blend_formulas_must_have_permanent_floor.md`:
* `maturity*real + (1-maturity)*floor` deadlocks when real signal hits 0
* at maturity=1. The max() form ensures the floor is a permanent minimum
* even when the policy converges (improving=1.0 → curiosity_dynamic=0,
* but `max(0, floor)=floor` so curiosity never zeros out). */
if (threadIdx.x == 6) {
const float curiosity_floor = CURIOSITY_PERMANENT_FRACTION
* curiosity_bound;
const float curiosity_dynamic = stagnant_or_worse * curiosity_bound;
scratch_out[6] = fmaxf(curiosity_dynamic, curiosity_floor);
}
/* Thread 7: saboteur_intensity_mult — z-driven blend modulated by
* engagement, post-clamped to [SABOTEUR_MIN, SABOTEUR_MAX].
*
* Without the post-clamp the engagement floor (0.1) would silently pull
* the modulated output below SABOTEUR_MIN at low improvement. Per spec
* §3.4 the post-clamp is required:
* intensity_z_term × max(engagement, ENGAGEMENT_FLOOR)
* THEN clamp(SABOTEUR_MIN, SABOTEUR_MAX).
*
* Behaviour: improving=1 → intensity_z_term=2.0 (max curriculum load);
* improving=0 → intensity_z_term=0.5 (relaxed curriculum,
* regression-mitigation). Engagement self-
* correction follows pearl_engagement_rate_self_correction. */
if (threadIdx.x == 7) {
const float engagement = isv[saboteur_engagement_slot];
const float intensity_z_term = SABOTEUR_MIN
+ (SABOTEUR_MAX - SABOTEUR_MIN) * improving;
const float modulated = intensity_z_term
* fmaxf(engagement, ENGAGEMENT_FLOOR);
scratch_out[7] = fmaxf(SABOTEUR_MIN, fminf(SABOTEUR_MAX, modulated));
}
/* Thread 8: weight_floor — already computed in shared mem; thread 8 just
* mirrors it to scratch so apply_pearls_ad_kernel can smooth the slot. */
if (threadIdx.x == 8) {
scratch_out[8] = adaptive_floor;
}
/* Thread 9: curiosity_bound — same mirror pattern as thread 8. */
if (threadIdx.x == 9) {
scratch_out[9] = curiosity_bound;
}
/* Threadfence after all per-thread writes — same pattern as the SP9
* producer kernel. Mapped-pinned coherence requires the system fence for
* host-side reads to observe writes (only relevant for tests; the
* apply_pearls_ad_kernel chained downstream reads same-stream-after, so
* device-side is sequentially consistent without a fence — the system
* fence is a defensive measure). */
__threadfence_system();
}

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

@@ -795,16 +795,31 @@ impl StateResetRegistry {
// of the Pearl A sentinel contract reset together per
// `feedback_no_partial_refactor.md`.
//
// The novelty-hash device buffer (`novelty_hash_buf`) is NOT an
// ISV slot but a separate device-side hash table sized for ~1M
// 16-bit codes. Its registry entry + dispatch arm
// (`fill_zero_async` at fold boundary) lands in Task A2
// alongside the buffer field's introduction on `GpuDqnTrainer`.
// Adding the entry here would be an orphan against a
// non-existent buffer field, violating
// `feedback_wire_everything_up.md`. A2 owns the contract end-
// to-end: buffer field, registry entry, dispatch arm, and
// launcher invocation in the same commit.
// SP11 Task A2 (2026-05-04): novelty visit-count hash table
// reset arm — closes the A0 deferral. The buffer is a 1M-slot
// f32 mapped-pinned table on `GpuDqnTrainer.novelty_hash_buf`;
// the dispatch arm (`reset_named_state` in `training_loop.rs`)
// zero-fills the host-side mapped slice via
// `host_slice_mut().fill(0.0)` — the kernel reading via
// `dev_ptr` observes zeros after the next stream-sync barrier.
//
// Cross-fold persistence is wrong here: the policy at fold N+1
// is a different model than fold N, so visit counts from N
// would mis-classify N+1's state-action distribution as
// already-explored. FoldReset is the correct semantics — the
// hash function (the projection matrix) stays frozen for
// reproducibility, but the visit-count table re-cold-starts.
//
// The 42×16 SimHash projection matrix on `novelty_simhash_proj`
// is intentionally NOT in this registry — it's the hash
// function, not state, and must outlive every fold for
// reproducibility (different runs with same seed get same
// hash function).
RegistryEntry {
name: "sp11_novelty_hash",
category: ResetCategory::FoldReset,
description: "GpuDqnTrainer.novelty_hash_buf — SP11 Fix 39 SimHash novelty visit-count hash table (1M slots × f32 = 4 MB mapped-pinned). Cleared at fold boundary so the new fold's policy starts with a clean visit-count → 1.0 novelty signal everywhere → maximal exploration pressure during the fold's cold-start window. The dispatch arm zero-fills via `host_slice_mut().fill(0.0)` — CPU writing a literal zero (NOT compute) per `feedback_no_htod_htoh_only_mapped_pinned.md`. The companion 42×16 SimHash projection matrix at `novelty_simhash_proj` is NOT in this registry — it's the hash function (frozen at trainer init from `SP11_PROJ_SEED_SALT`) and must outlive every fold for reproducibility.",
},
RegistryEntry {
name: "sp11_reward_popart_weight",
category: ResetCategory::FoldReset,

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

@@ -3447,6 +3447,36 @@ impl DQNTrainer {
tracing::warn!(error = %e,
"SP11 saboteur_engagement_compute launch failed");
}
// SP11 Fix 39 (2026-05-04, Task A2): reward-
// subsystem controller. Reads all 5 canary
// ISV slots [350..360) (val-sharpe Δ +
// variance, 6 mag-ratios, saboteur engagement,
// PnL magnitude EMA) and writes 10 outputs
// through chained Pearls A+D into ISV[340..350).
//
// Pre-conditions (must run AFTER, on the same
// stream): the 3 A1 canary launches above
// populate slots [350..360) for THIS step.
// The controller depends on canary outputs
// being current, so this launch lands AFTER
// them in the same per-epoch metrics block
// and BEFORE the HEALTH_DIAG emit so the
// operator sees fresh controller outputs.
//
// No consumer reads ISV[340..350) yet — Layer
// A is additive; consumer migration lands
// atomically in B1 per
// `feedback_no_partial_refactor.md`. Training
// behaviour is unchanged from A1; the only
// observable change is the new HEALTH_DIAG
// `sp11_reward` line emitted below.
if let Err(e) = fused.trainer()
.launch_sp11_reward_subsystem_controller()
{
tracing::warn!(error = %e,
"SP11 reward_subsystem_controller launch failed");
}
}
}
}
@@ -4247,6 +4277,75 @@ impl DQNTrainer {
);
}
// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem
// controller observability. Emits the 10 controller outputs
// (6 component weights + curiosity_pressure +
// saboteur_intensity_mult + reward_weight_floor +
// curiosity_bound) and the inline-computed `improvement_z`
// canary derivation so operators can see whether the
// controller is acting on a fresh val-sharpe trend signal or
// a Pearl-A-sentinel-bootstrap state.
//
// `improvement_z` is NOT in ISV — the controller consumes it
// internally (sigmoid mapping → improving/stagnant). For
// HEALTH_DIAG visibility we re-derive it here from the same
// canary slots the kernel reads:
// z = delta_ema / max(sqrt(var_ema), 1e-6)
// which is the kernel's `dz` formula verbatim per spec §3.4.
// Per `feedback_trust_code_not_docs.md` the formula stays in
// lockstep with the kernel — if the kernel changes, this line
// updates in the same commit.
//
// No consumer reads ISV[340..350) yet (Layer A is additive),
// so all weights show their FoldReset sentinel zero on the
// first-emit epoch then trend toward populated values via
// Pearl A bootstrap. The chained `apply_pearls_ad_kernel`
// smoothing means the line shows the SMOOTHED outputs
// (consumer-side reads), not raw scratch values, per spec
// §3.4.1.
{
use crate::cuda_pipeline::sp5_isv_slots::{
REWARD_POPART_WEIGHT_INDEX, REWARD_CF_WEIGHT_INDEX,
REWARD_TRAIL_WEIGHT_INDEX, REWARD_MICRO_WEIGHT_INDEX,
REWARD_OPP_COST_WEIGHT_INDEX, REWARD_BONUS_WEIGHT_INDEX,
CURIOSITY_PRESSURE_INDEX, SABOTEUR_INTENSITY_MULT_INDEX,
REWARD_WEIGHT_FLOOR_INDEX, CURIOSITY_BOUND_INDEX,
VAL_SHARPE_DELTA_EMA_INDEX, VAL_SHARPE_VAR_EMA_INDEX,
};
let (w_pop, w_cf, w_tr, w_mi, w_oc, w_bn,
curiosity_p, sab_mult, w_floor, curiosity_b,
delta_ema, var_ema) =
if let Some(ref fused) = self.fused_ctx {
let trainer = fused.trainer();
(
trainer.read_isv_signal_at(REWARD_POPART_WEIGHT_INDEX),
trainer.read_isv_signal_at(REWARD_CF_WEIGHT_INDEX),
trainer.read_isv_signal_at(REWARD_TRAIL_WEIGHT_INDEX),
trainer.read_isv_signal_at(REWARD_MICRO_WEIGHT_INDEX),
trainer.read_isv_signal_at(REWARD_OPP_COST_WEIGHT_INDEX),
trainer.read_isv_signal_at(REWARD_BONUS_WEIGHT_INDEX),
trainer.read_isv_signal_at(CURIOSITY_PRESSURE_INDEX),
trainer.read_isv_signal_at(SABOTEUR_INTENSITY_MULT_INDEX),
trainer.read_isv_signal_at(REWARD_WEIGHT_FLOOR_INDEX),
trainer.read_isv_signal_at(CURIOSITY_BOUND_INDEX),
trainer.read_isv_signal_at(VAL_SHARPE_DELTA_EMA_INDEX),
trainer.read_isv_signal_at(VAL_SHARPE_VAR_EMA_INDEX),
)
} else {
(0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32,
0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32,
0.0_f32, 0.0_f32)
};
let improvement_z = delta_ema / var_ema.sqrt().max(1e-6_f32);
tracing::info!(
"HEALTH_DIAG[{}]: sp11_reward [w_pop={:.3} w_cf={:.3} w_tr={:.3} w_mi={:.3} w_oc={:.3} w_bn={:.3}] [curiosity_p={:.3} sab_mult={:.3} w_floor={:.3} curiosity_b={:.3}] [improvement_z={:.3}]",
epoch,
w_pop, w_cf, w_tr, w_mi, w_oc, w_bn,
curiosity_p, sab_mult, w_floor, curiosity_b,
improvement_z,
);
}
// SP6 Pearl 2: per-branch budget HEALTH_DIAG lines — enables Layer C debugging
// by showing how much each branch's loss budget differs from the trunk mean.
if let Some(arr) = self.last_c51_budget_per_branch {
@@ -7178,6 +7277,22 @@ impl DQNTrainer {
fused.trainer().write_isv_signal_at(PNL_REWARD_MAGNITUDE_EMA_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
// slice (CPU writes a literal zero — allowed under
// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write
// rule, NOT a host-side compute).
//
// The companion projection matrix `novelty_simhash_proj` is
// NOT reset here — it's the hash function (frozen at trainer
// init from `SP11_PROJ_SEED_SALT`) and must outlive every
// fold for reproducibility. See the registry entry header.
"sp11_novelty_hash" => {
if let Some(ref mut fused) = self.fused_ctx {
fused.trainer_mut().novelty_hash_buf.host_slice_mut().fill(0.0);
}
}
"sp5_pnl_aggregation" => {
// Layer D Task D1: zero ISV[286..290) — total/mean/var/max_dd —
// so the new fold's first `launch_sp5_pnl_aggregation` triggers

300
crates/ml/tests/sp11_producer_unit_tests.rs
View File

@@ -34,7 +34,9 @@ use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
use ml::cuda_pipeline::sp11_isv_slots::{
PNL_REWARD_MAGNITUDE_EMA_INDEX, VAL_SHARPE_DELTA_EMA_INDEX, VAL_SHARPE_VAR_EMA_INDEX,
PNL_REWARD_MAGNITUDE_EMA_INDEX, REWARD_COMPONENT_MAG_RATIO_BASE,
SABOTEUR_ENGAGEMENT_RATE_INDEX,
VAL_SHARPE_DELTA_EMA_INDEX, VAL_SHARPE_VAR_EMA_INDEX,
};
/// ISV total dimension (must match `gpu_dqn_trainer::ISV_TOTAL_DIM`).
@@ -57,6 +59,9 @@ const SP11_SABOTEUR_ENGAGEMENT_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/saboteur_engagement_compute_kernel.cubin"));
const SP11_MAG_RATIO_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/reward_component_mag_ratio_compute_kernel.cubin"));
/// SP11 A2 (2026-05-04): controller cubin for the 3 oracle tests below.
const SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/reward_subsystem_controller_kernel.cubin"));
/// Build a CUDA stream against device 0. Mirrors `make_test_stream` in
/// `sp4_producer_unit_tests.rs` and `distributional_q_tests.rs`.
@@ -273,3 +278,296 @@ fn reward_component_mag_ratio_normalises_and_mirrors_popart() {
out[6]
);
}
// ── SP11 Task A2 — reward-subsystem controller GPU oracle tests ──────────
/// Populate ISV with the 5 controller canary inputs the kernel reads.
/// `mag_ratios` are written contiguously at REWARD_COMPONENT_MAG_RATIO_BASE
/// (slots [352..358) in the production layout); the val-sharpe Δ + var
/// canaries land at slots 350/351; saboteur engagement at 358; PnL EMA at
/// 359. Each test calls this with synthetic-but-physically-meaningful
/// values so the assertion targets are derivable from the kernel's
/// algebra in spec §3.4 — no CPU-reference oracle is required (and per
/// `feedback_no_cpu_test_fallbacks.md` would be forbidden anyway).
fn populate_controller_isv(
isv: &mut MappedF32Buffer,
delta: f32,
var: f32,
mag_ratios: [f32; 6],
saboteur_engagement: f32,
pnl_mag: f32,
) {
let h = isv.host_slice_mut();
h[VAL_SHARPE_DELTA_EMA_INDEX] = delta;
h[VAL_SHARPE_VAR_EMA_INDEX] = var;
for c in 0..6 {
h[REWARD_COMPONENT_MAG_RATIO_BASE + c] = mag_ratios[c];
}
h[SABOTEUR_ENGAGEMENT_RATE_INDEX] = saboteur_engagement;
h[PNL_REWARD_MAGNITUDE_EMA_INDEX] = pnl_mag;
}
/// Launch the controller producer with the standard launch config (1
/// block, 10 threads). Callers pre-populate `isv` via
/// `populate_controller_isv`; outputs land in `scratch[0..10)` per the
/// kernel signature (the launcher does NOT apply the kernel-side
/// scratch-base offset because the test allocates a fresh 10-element
/// scratch — the production launcher's offset arithmetic is exercised in
/// the integration-tier smoke tests).
fn launch_controller_kernel(
stream: &Arc<CudaStream>,
f: &CudaFunction,
isv: &MappedF32Buffer,
scratch: &MappedF32Buffer,
) {
let isv_dev = isv.dev_ptr;
let scratch_dev = scratch.dev_ptr;
let delta_ema_slot_i32: i32 = VAL_SHARPE_DELTA_EMA_INDEX as i32;
let var_ema_slot_i32: i32 = VAL_SHARPE_VAR_EMA_INDEX as i32;
let mag_ratio_base_slot_i32: i32 = REWARD_COMPONENT_MAG_RATIO_BASE as i32;
let saboteur_engagement_slot_i32: i32 = SABOTEUR_ENGAGEMENT_RATE_INDEX as i32;
let pnl_mag_ema_slot_i32: i32 = PNL_REWARD_MAGNITUDE_EMA_INDEX as i32;
unsafe {
stream
.launch_builder(f)
.arg(&isv_dev)
.arg(&scratch_dev)
.arg(&delta_ema_slot_i32)
.arg(&var_ema_slot_i32)
.arg(&mag_ratio_base_slot_i32)
.arg(&saboteur_engagement_slot_i32)
.arg(&pnl_mag_ema_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (10, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch reward_subsystem_controller_kernel");
}
stream.synchronize().expect("sync after controller launch");
}
/// SP11 A2 — controller midpoint behavior at z=0.
///
/// Inputs:
/// delta = 0, var = 1 → dz = 0 / sqrt(1) = 0 → improving = sigmoid(0) = 0.5,
/// stagnant = 0.5
/// mag_ratios = [1/6, 1/6, 1/6, 1/6, 1/6, 1/6] → uniform, min_grad_ratio = 1/6
/// adaptive_floor = max(0.01, 0.5 × 1/6) = 0.0833...
/// pnl_mag = 1.0 → curiosity_bound = 1.0 × 0.3 = 0.3
/// curiosity_floor = 0.2 × 0.3 = 0.06
/// curiosity_dynamic = 0.5 × 0.3 = 0.15
/// curiosity_pressure = max(0.15, 0.06) = 0.15
/// saboteur_engagement = 0.5 → intensity_z_term = 0.5 + 1.5 × 0.5 = 1.25
/// modulated = 1.25 × max(0.5, 0.1) = 0.625
/// post-clamp [0.5, 2.0] = 0.625 (in range)
///
/// Per-component winner_weight = 1/6, diversifier = (1 - 1/6)/5 = 1/6;
/// blend = 0.5 × 1/6 + 0.5 × 1/6 = 1/6 each. After floor (0.0833 < 1/6) and
/// renorm to Σ=1, each weight stays at 1/6.
#[test]
#[ignore = "requires GPU"]
fn controller_z_score_at_zero_yields_midpoint_outputs() {
let stream = make_test_stream();
let f = load_kernel(
&stream,
SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN,
"reward_subsystem_controller_kernel",
);
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
populate_controller_isv(
&mut isv,
/*delta*/ 0.0,
/*var*/ 1.0,
/*mag_ratios*/ [1.0_f32 / 6.0; 6],
/*saboteur_engagement*/ 0.5,
/*pnl_mag*/ 1.0,
);
let scratch = unsafe { MappedF32Buffer::new(10) }
.expect("alloc scratch (10 f32)");
launch_controller_kernel(&stream, &f, &isv, &scratch);
let s = scratch.read_all();
// Σweights = 1.0 ± 1e-5 — renormalization invariant per spec §3.4.
let weight_sum: f32 = s[0..6].iter().sum();
assert!(
(weight_sum - 1.0).abs() < 1e-5,
"weights sum = {weight_sum}, expected 1.0"
);
// Each weight at 1/6 (uniform input → uniform output through the
// 50/50 blend → renorm preserves uniformity).
for c in 0..6 {
assert!(
(s[c] - 1.0_f32 / 6.0).abs() < 1e-4,
"weight[{c}] = {}, expected 1/6 (≈ 0.16667)",
s[c]
);
}
// Curiosity pressure at midpoint: dynamic dominates floor.
assert!(
(s[6] - 0.15).abs() < 1e-4,
"curiosity_pressure = {}, expected 0.15 (= 0.5 × 0.3)",
s[6]
);
// Saboteur intensity: 0.625 (in-range, post-clamp is identity).
assert!(
(s[7] - 0.625).abs() < 1e-3,
"saboteur_intensity = {}, expected 0.625",
s[7]
);
// Adaptive floor = max(WEIGHT_HARD_FLOOR=0.01, 0.5 × 1/6) = 0.0833...
let expected_floor = 0.5_f32 * (1.0_f32 / 6.0);
assert!(
(s[8] - expected_floor).abs() < 1e-4,
"adaptive_floor = {}, expected {expected_floor}",
s[8]
);
// Curiosity bound = 0.3 (= pnl_mag × CURIOSITY_BOUND_FRACTION).
assert!(
(s[9] - 0.3).abs() < 1e-5,
"curiosity_bound = {}, expected 0.3",
s[9]
);
}
/// SP11 A2 — weight renormalization works after the per-component floor.
///
/// Pathological mag_ratios: one large (0.95), five small (0.01). The
/// pre-floor blend with improving=stagnant=0.5 yields heavily skewed
/// weights; the floor (`0.5 × 0.01 = 0.005` ≤ WEIGHT_HARD_FLOOR=0.01,
/// so adaptive_floor = 0.01) clamps the small weights to 0.01; renorm to
/// Σ=1 then redistributes the residual mass.
///
/// The test asserts Σweights = 1 ± 1e-5 (the renormalization invariant)
/// and that no weight drops below the WEIGHT_HARD_FLOOR after renorm.
#[test]
#[ignore = "requires GPU"]
fn controller_weights_renormalize_after_floor() {
let stream = make_test_stream();
let f = load_kernel(
&stream,
SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN,
"reward_subsystem_controller_kernel",
);
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
populate_controller_isv(
&mut isv,
/*delta*/ 1.0, // any non-zero
/*var*/ 1.0,
/*mag_ratios*/ [0.95_f32, 0.01, 0.01, 0.01, 0.01, 0.01],
/*saboteur_engagement*/ 0.5,
/*pnl_mag*/ 1.0,
);
let scratch = unsafe { MappedF32Buffer::new(10) }
.expect("alloc scratch (10 f32)");
launch_controller_kernel(&stream, &f, &isv, &scratch);
let s = scratch.read_all();
// Renormalization invariant: Σweights = 1.0 ± 1e-5. PopArt's reward-
// magnitude EMA assumes a consistent reward-scale; this is the
// non-negotiable guarantee per spec §3.4.
let weight_sum: f32 = s[0..6].iter().sum();
assert!(
(weight_sum - 1.0).abs() < 1e-5,
"weights sum = {weight_sum} after renormalization, expected 1.0"
);
// No weight below WEIGHT_HARD_FLOOR=0.01 after renorm. Renorm scales
// by 1/Σ_pre_floor_blend; when the pre-floor blend already had each
// weight ≥ 0.01 (because the floor pass clamped at adaptive_floor =
// 0.01), the post-renorm weight is `0.01 / blend_sum` which is at
// least slightly above 0.01 only if blend_sum < 1. With one weight
// dominating (0.95-derived blend ≈ 0.5) and five at floor, blend_sum
// ≈ 0.5 + 5 × 0.01 = 0.55, so the floored weights post-renorm are
// 0.01 / 0.55 ≈ 0.018 — comfortably above 0.01. The lower bound check
// `>= 0.01 - tolerance` validates the floor's structural guarantee.
for c in 0..6 {
assert!(
s[c] >= 0.01 - 1e-5,
"weight[{c}] = {} below WEIGHT_HARD_FLOOR=0.01 after renorm",
s[c]
);
}
}
/// SP11 A2 — saboteur post-clamp holds SABOTEUR_MIN at extreme regression.
///
/// Inputs:
/// delta = -100, var = 1 → dz = -100 / 1 = -100 → improving ≈ 0,
/// stagnant ≈ 1
/// intensity_z_term = 0.5 + 1.5 × 0 = 0.5
/// saboteur_engagement = 0.0 → max(0.0, 0.1) = 0.1 (engagement floor)
/// modulated = 0.5 × 0.1 = 0.05
/// post-clamp = max(0.5, min(2.0, 0.05)) = 0.5 ← SABOTEUR_MIN
///
/// Without the post-clamp the engagement floor would silently pull the
/// output to 0.05, well below SABOTEUR_MIN=0.5. The post-clamp is the
/// non-negotiable guarantee per spec §3.4 — this test fails if the
/// kernel ever drops the post-clamp.
///
/// Curiosity at extreme regression: stagnant ≈ 1 → curiosity_dynamic ≈
/// 0.3; curiosity_floor = 0.06; pressure = max(0.3, 0.06) = 0.3 — the
/// permanent-floor pattern protects the high-stagnation peak from
/// numerical dips.
#[test]
#[ignore = "requires GPU"]
fn controller_saboteur_post_clamp_holds_min() {
let stream = make_test_stream();
let f = load_kernel(
&stream,
SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN,
"reward_subsystem_controller_kernel",
);
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
populate_controller_isv(
&mut isv,
/*delta*/ -100.0, // extreme regression
/*var*/ 1.0,
/*mag_ratios*/ [1.0_f32 / 6.0; 6], // uniform — keeps the test focused on saboteur
/*saboteur_engagement*/ 0.0, // tests engagement floor
/*pnl_mag*/ 1.0,
);
let scratch = unsafe { MappedF32Buffer::new(10) }
.expect("alloc scratch (10 f32)");
launch_controller_kernel(&stream, &f, &isv, &scratch);
let s = scratch.read_all();
// Saboteur post-clamp at SABOTEUR_MIN=0.5. Without the post-clamp the
// engagement floor (0.1) would multiply intensity_z_term=0.5 to give
// 0.05 — well below 0.5.
assert!(
(s[7] - 0.5).abs() < 1e-4,
"saboteur_intensity = {}, expected SABOTEUR_MIN=0.5 after post-clamp",
s[7]
);
// Curiosity at max regression: stagnant ≈ 1.0 → curiosity_dynamic ≈
// curiosity_bound = 0.3. Permanent floor 0.06 below that → max() =
// 0.3.
assert!(
(s[6] - 0.3).abs() < 1e-3,
"curiosity_pressure at max regression = {}, expected ≈ 0.3 (curiosity_bound)",
s[6]
);
}

57
docs/isv-slots.md
View File

@@ -184,7 +184,7 @@ gradient contribution = `iqn_budget` (same as SP5 Layer B). The averaged
**Files changed:** `crates/ml/src/cuda_pipeline/gpu_iqn_head.rs`,
`crates/ml/src/trainers/dqn/fused_training.rs`.
## SP11: Reward as controlled subsystem (Tasks A0 + A1)
## SP11: Reward as controlled subsystem (Tasks A0 + A1 + A2)
SP11 Task A0 (Fix 39, 2026-05-04) extends `ISV_TOTAL_DIM` from 340 → 360 by
allocating 20 contiguous slots at `[340..360)` for the reward-subsystem
@@ -195,11 +195,32 @@ and are re-exported via `cuda_pipeline::sp5_isv_slots::*`.
chained with `apply_pearls_ad_kernel` for Pearls A+D smoothing. Slots
[350..360) now have producers populating them every epoch (val_sharpe
emit cadence for the Δ canary; per-epoch metrics block for mag-ratio +
saboteur engagement). **Layer A is still additive — no consumer reads
any [340..360) slot in this commit. Training behaviour remains
identical to A0; the controller (A2) and atomic consumer migration
(Layer B) wire downstream.** Slots [340..350) (controller outputs)
remain zero-initialized until A2.
saboteur engagement).
**Task A2 (2026-05-04): controller producer + SimHash novelty buffer**.
The main `reward_subsystem_controller_kernel` reads all 5 canary slots
[350..360), runs the spec §3.4 control law (true Z-score → sigmoid blend
of winner/diversifier weights, post-floor renormalize Σ=1, permanent-
floor curiosity, post-clamp saboteur), and writes 10 outputs to scratch
which a chained `apply_pearls_ad_kernel` (n_slots=10) smooths into the
remaining slots [340..350). After A2 **all 20 SP11 slots populate every
step**.
A2 also lands the SimHash novelty infrastructure for B1's replay-time
curiosity bonus: `novelty_simhash_proj_init_kernel` (one-shot Philox-
driven init at trainer construct, populates a 42×16 ±1 projection
matrix on-device per `feedback_no_cpu_forwards.md`),
`novelty_simhash_kernel` (lookup + update, race-tolerated update per
`feedback_no_atomicadd.md` — under-counts bias novelty UPWARD, the safe
direction). The 1M-slot hash table at `GpuDqnTrainer.novelty_hash_buf`
gets a fold-reset registry entry (`sp11_novelty_hash`) that closes the
A0 deferral; the projection matrix is frozen for the run lifetime and
intentionally has NO registry entry.
**Layer A is still additive — no consumer reads any [340..360) slot in
this commit. Training behaviour remains identical to A1/A0; only the
HEALTH_DIAG `sp11_reward` line is observably new. Atomic consumer
migration (Layer B) wires downstream.**
A1 producers:
- `val_sharpe_delta_compute_kernel.cu` — two-pass: raw Δ + squared
@@ -215,9 +236,12 @@ A1 producers:
popart magnitude into scratch[6] as a side-output. Chain (n_slots=6)
→ ISV[352..358); chain (n_slots=1) → ISV[359].
3 GPU oracle tests in `crates/ml/tests/sp11_producer_unit_tests.rs`
cover first-observation behaviour, threshold classification, and
normalisation invariants. All MappedF32Buffer fixtures (zero
6 GPU oracle tests in `crates/ml/tests/sp11_producer_unit_tests.rs`
cover (A1) first-observation behaviour, threshold classification, and
normalisation invariants on the canary kernels; plus (A2) the
controller midpoint-at-z=0 invariant, weight-renormalization-after-floor
invariant, and saboteur-post-clamp-at-extreme-regression invariant.
All MappedF32Buffer fixtures (zero
htod_copy/dtoh_sync_copy/alloc_zeros per
`feedback_no_htod_htoh_only_mapped_pinned`).
@@ -249,11 +273,16 @@ launcher's seed derivation.
All 20 slots are FoldReset (sentinel 0; Pearl A bootstraps from first
producer launch on each fold per `pearl_first_observation_bootstrap.md`).
Producers and consumers wire in subsequent A1/A2/Layer-B tasks. The
novelty-hash device buffer (separate from the ISV bus, sized for ~1M 16-bit
codes) lands in A2 alongside its registry entry and dispatch arm — adding
the entry in A0 would orphan it against a non-existent buffer field, so
A2 owns the contract end-to-end.
After A2 all 20 slots have producers; consumers wire in Layer B.
The novelty-hash device buffer (`GpuDqnTrainer.novelty_hash_buf`,
1M slots × f32 mapped-pinned) lands in A2 alongside its registry entry
(`sp11_novelty_hash`) and dispatch arm. The 42×16 SimHash projection
matrix (`GpuDqnTrainer.novelty_simhash_proj`) is populated on-device by
`launch_novelty_simhash_proj_init` at trainer construct time (Philox-
driven from `SP11_PROJ_SEED_SALT`); it is **frozen for the run
lifetime** and intentionally has NO registry entry — it is the hash
function, not state.
**Spec:** `docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md`
**Plan:** `docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md`