From 25eba79ad5c4a7b0c5e326d04a8de75e1ebe937e Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 4 May 2026 02:26:05 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp11):=20A2=20=E2=80=94=20controller=20ker?= =?UTF-8?q?nel=20+=20SimHash=20novelty=20buffer?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- crates/ml/build.rs | 16 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 457 +++++++++++++++++- .../cuda_pipeline/novelty_simhash_kernel.cu | 104 ++++ .../novelty_simhash_proj_init_kernel.cu | 50 ++ .../reward_subsystem_controller_kernel.cu | 211 ++++++++ .../src/trainers/dqn/state_reset_registry.rs | 35 +- .../src/trainers/dqn/trainer/training_loop.rs | 115 +++++ crates/ml/tests/sp11_producer_unit_tests.rs | 300 +++++++++++- docs/isv-slots.md | 57 ++- 9 files changed, 1318 insertions(+), 27 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu create mode 100644 crates/ml/src/cuda_pipeline/novelty_simhash_proj_init_kernel.cu create mode 100644 crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index c81abd405..ba78c6819 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -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 diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index df8951e33..0f2760049 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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::() 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, diff --git a/crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu b/crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu new file mode 100644 index 000000000..bfc8c1021 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu @@ -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 +#include + +/* 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; +} diff --git a/crates/ml/src/cuda_pipeline/novelty_simhash_proj_init_kernel.cu b/crates/ml/src/cuda_pipeline/novelty_simhash_proj_init_kernel.cu new file mode 100644 index 000000000..82d6aaee2 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/novelty_simhash_proj_init_kernel.cu @@ -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 + +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; +} diff --git a/crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu b/crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu new file mode 100644 index 000000000..3dc49425c --- /dev/null +++ b/crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu @@ -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 +#include + +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(); +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 424ae7399..037c49c5f 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -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, diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 48be6d748..52a82eb51 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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 diff --git a/crates/ml/tests/sp11_producer_unit_tests.rs b/crates/ml/tests/sp11_producer_unit_tests.rs index d5bf26abb..cce342bd1 100644 --- a/crates/ml/tests/sp11_producer_unit_tests.rs +++ b/crates/ml/tests/sp11_producer_unit_tests.rs @@ -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, + 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] + ); +} diff --git a/docs/isv-slots.md b/docs/isv-slots.md index b179cd934..14fddcfa6 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -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`