From e0d3abd9d2c54e11683a80af35ea993beccb97b6 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 4 May 2026 00:40:36 +0200 Subject: [PATCH] plan(sp11): patch all 5 violations + 2 gaps from review MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit feedback_no_htod_htoh_only_mapped_pinned (tests not exempt): - All test fixtures converted from htod_copy/dtoh_sync_copy to MappedF32Buffer with host_slice / host_slice_mut access. - Novelty hash buffer + projection matrix changed from cudarc::CudaSlice + alloc_zeros to MappedF32Buffer. feedback_no_cpu_forwards (CPU is read-only): - Projection matrix initialization changed from host-side StdRng + host_slice_mut writes to a one-shot GPU init kernel (novelty_simhash_proj_init_kernel) using Philox seeded from config.seed. No host RNG, no host writes. feedback_no_cpu_compute_strict (saboteur multiplication): - B1 step 5 reverted from Rust-side `read_isv_slot * scale` to GPU-side: pass base scale + ISV pointer + slot index to the saboteur perturbation kernel; multiplication happens on-device. feedback_trust_code_not_docs (grad-ratio terminology): - Spec §3.3.1 "per-component grad EMA" was wrong — SP4 grad-balancer is per-branch (4 slots), not per-reward-component. Renamed: reward_component_grad_ratio_compute_kernel → reward_component_mag_ratio_compute_kernel REWARD_COMPONENT_GRAD_RATIO_BASE → REWARD_COMPONENT_MAG_RATIO_BASE Source: existing REWARD_POPART_EMA_INDEX..+6 (per-component reward magnitude EMAs from SP4 reward_component_ema_kernel). Semantic equivalent for the controller's exploit/diversify blend. feedback_no_stubs (dead parameter): - Removed `eps_div_idx_unused` from mag_ratio kernel signature. Saboteur engagement: missing producer specified - Spec §3.3.1's two-reward-arrays formulation replaced with single `saboteur_delta_reward_buf` produced by the saboteur perturbation kernel itself (single reward computation, diff emitted as side output). Engagement kernel signature simplified to one input array. PNL_REWARD_MAGNITUDE_EMA_INDEX (slot 359): producer wired - mag-ratio kernel mirrors `isv[REWARD_POPART_EMA_INDEX]` to scratch_out[6]; chained apply_pearls_ad targets slot 359. Replay sample kernel location specified - graph_utility_kernels.cu:71 (gather_f32_scalar). New sibling kernel `gather_replay_reward_with_curiosity` defined; replaces the existing scalar gather (no legacy alias per feedback_no_legacy_aliases). novelty_simhash_lookup runs before, novelty_simhash_update after. A2 controller test placeholders → full GPU oracle assertions - Three controller tests (z=0 midpoint, weight renorm, saboteur clamp) have full mapped-pinned fixtures with assertions on weight sum, individual values, post-clamp bounds. Plan now passes: - feedback_no_htod_htoh_only_mapped_pinned (tests + production) - feedback_no_cpu_forwards (CPU never writes/computes for GPU) - feedback_no_cpu_compute_strict (all multiplications GPU-side) - feedback_no_atomicadd (race-tolerated non-atomic, safety documented) - feedback_no_partial_refactor (Layer B atomic; saboteur kernel sig change touches all callers in the same commit) - feedback_no_stubs (no dead parameters) - feedback_trust_code_not_docs (corrected spec terminology) - feedback_wire_everything_up (every new field has init + producer) - feedback_no_legacy_aliases (old gather_f32_scalar replaced, deleted) 1447 lines, +268 from previous version. --- ...-04-sp11-reward-as-controlled-subsystem.md | 482 ++++++++++++++---- 1 file changed, 375 insertions(+), 107 deletions(-) diff --git a/docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md b/docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md index c44a5b9db..ba109c975 100644 --- a/docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md +++ b/docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md @@ -4,7 +4,7 @@ **Goal:** Promote every reward-path degree of freedom (component weights, shaping multipliers, exploration bonus, saboteur curriculum) to ISV-driven slots controlled by a unified Z-score-based controller. Replace hardcoded `cf_weight=0.3f` and similar shaping constants with controller outputs. Add curiosity bonus as a NEW reward dimension applied at *replay time* (not collect-time) to escape the ep1-overfitting fixed point observed in T10 train-multi-seed-xkjkb. -**Architecture:** 20 new ISV slots [340..360) populated by 4 new GPU producer kernels (`val_sharpe_delta_compute_kernel`, `saboteur_engagement_compute_kernel`, `reward_component_grad_ratio_compute_kernel`, `reward_subsystem_controller_kernel`). All producers chained on their stream with `apply_pearls_ad_kernel` for Pearls A+D smoothing per `feedback_no_cpu_compute_strict`. State-action novelty via SimHash 42×16 → 16-bit code, 1M-slot device hash table reset per fold. Three layers: Layer A (3 additive commits — slots, canaries, controller), Layer B (1 atomic consumer-migration commit), Layer C (smoke + T10 validation + close-out). +**Architecture:** 20 new ISV slots [340..360) populated by 4 new GPU producer kernels (`val_sharpe_delta_compute_kernel`, `saboteur_engagement_compute_kernel`, `reward_component_mag_ratio_compute_kernel`, `reward_subsystem_controller_kernel`). All producers chained on their stream with `apply_pearls_ad_kernel` for Pearls A+D smoothing per `feedback_no_cpu_compute_strict`. State-action novelty via SimHash 42×16 → 16-bit code, 1M-slot device hash table reset per fold. Three layers: Layer A (3 additive commits — slots, canaries, controller), Layer B (1 atomic consumer-migration commit), Layer C (smoke + T10 validation + close-out). **Tech Stack:** Rust 1.85+, CUDA 12.4, cudarc, mapped-pinned memory (`MappedF32Buffer`/`MappedI32Buffer`), block-shared-memory tree-reduce (no atomicAdd per `feedback_no_atomicadd`), GPU-only Pearls A+D via `apply_pearls_ad_kernel` (no host compute per `feedback_no_cpu_compute_strict`), L40S smoke validation via Argo Workflows. @@ -33,7 +33,7 @@ | `crates/ml/src/cuda_pipeline/sp11_isv_slots.rs` | A0 | Define 20 SP11 ISV slot index constants + Invariant-1 fraction constants | | `crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu` | A1 | Two-pass delta_ema + delta_var_ema producer | | `crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu` | A1 | Per-bar engagement reduce + EMA | -| `crates/ml/src/cuda_pipeline/reward_component_grad_ratio_compute_kernel.cu` | A1 | Read 6 grad EMAs → normalized ratios | +| `crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu` | A1 | Read 6 grad EMAs → normalized ratios | | `crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu` | A2 | Main controller: 5 canaries → 10 outputs (scratch) | | `crates/ml/src/cuda_pipeline/novelty_simhash_kernel.cu` | A2 | SimHash 42×16 projection → bucket count update + lookup | | `crates/ml/tests/sp11_producer_unit_tests.rs` | A1+A2 | 6+ GPU-gated unit tests | @@ -89,7 +89,7 @@ //! - 6 component weights [340..346) //! - 4 controller scalar outputs [346..350) //! - 2 val-sharpe canaries [350..352) -//! - 6 grad-ratio canaries [352..358) +//! - 6 mag-ratio canaries [352..358) //! - 1 saboteur engagement (358) //! - 1 PnL magnitude EMA (359) //! @@ -114,9 +114,16 @@ pub const CURIOSITY_BOUND_INDEX: usize = 349; pub const VAL_SHARPE_DELTA_EMA_INDEX: usize = 350; pub const VAL_SHARPE_VAR_EMA_INDEX: usize = 351; -// Canaries — per-component grad ratios (6 contiguous slots) -pub const REWARD_COMPONENT_GRAD_RATIO_BASE: usize = 352; -pub const REWARD_COMPONENT_GRAD_RATIO_COUNT: usize = 6; +// Canaries — per-component reward-magnitude ratios (6 contiguous slots). +// Source: REWARD_POPART_EMA_INDEX..+6 (existing per-component reward +// magnitude EMAs from reward_component_ema_kernel). Spec §3.3.1 said +// "grad EMA" but the SP4 grad-balancer is per-branch (4 slots), not +// per-reward-component. Per `feedback_trust_code_not_docs`: code wins, +// terminology corrected. Semantics for the controller are equivalent — +// "which component contributes most to total reward" maps the same way +// in the winner_weight / diversifier_weight blend. +pub const REWARD_COMPONENT_MAG_RATIO_BASE: usize = 352; +pub const REWARD_COMPONENT_MAG_RATIO_COUNT: usize = 6; // Canaries — saboteur + PnL signal scale pub const SABOTEUR_ENGAGEMENT_RATE_INDEX: usize = 358; @@ -158,11 +165,11 @@ entries.push(ResetEntry { slot: REWARD_POPART_WEIGHT_INDEX, category: ResetCategory::FoldReset, sentinel_f32: 0.0, - description: "ISV[REWARD_POPART_WEIGHT_INDEX=340] — SP11 popart reward component weight. Produced by reward_subsystem_controller_kernel from grad-ratio canaries and improvement_z. Consumed in reward composition. FoldReset sentinel 0; Pearl A bootstraps from first controller emit (per pearl_first_observation_bootstrap.md).", + description: "ISV[REWARD_POPART_WEIGHT_INDEX=340] — SP11 popart reward component weight. Produced by reward_subsystem_controller_kernel from mag-ratio canaries and improvement_z. Consumed in reward composition. FoldReset sentinel 0; Pearl A bootstraps from first controller emit (per pearl_first_observation_bootstrap.md).", }); ``` -Add equivalent entries for all 20 slots: 6 component weights, 4 controller outputs (curiosity_pressure, saboteur_intensity_mult, reward_weight_floor, curiosity_bound), 2 val-sharpe canaries, 6 grad-ratio canaries (loop), saboteur_engagement_rate, pnl_reward_magnitude_ema. +Add equivalent entries for all 20 slots: 6 component weights, 4 controller outputs (curiosity_pressure, saboteur_intensity_mult, reward_weight_floor, curiosity_bound), 2 val-sharpe canaries, 6 mag-ratio canaries (loop), saboteur_engagement_rate, pnl_reward_magnitude_ema. - [ ] **Step 5: Add novelty-hash reset arm** @@ -209,12 +216,12 @@ EOF ### Task A1: Three canary producer kernels -**Goal:** Implement 3 canary kernels that populate the 9 canary slots (delta_ema, delta_var_ema, 6 grad-ratios, engagement_rate, pnl_magnitude_ema). Each kernel chains `apply_pearls_ad_kernel` on its stream for Pearls A+D smoothing. Wire launchers; the controller (A2) will read these. Existing training continues to behave identically because no consumer reads canary slots in production paths. +**Goal:** Implement 3 canary kernels that populate the 9 canary slots (delta_ema, delta_var_ema, 6 mag-ratios, engagement_rate, pnl_magnitude_ema). Each kernel chains `apply_pearls_ad_kernel` on its stream for Pearls A+D smoothing. Wire launchers; the controller (A2) will read these. Existing training continues to behave identically because no consumer reads canary slots in production paths. **Files:** - Create: `crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu` - Create: `crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu` -- Create: `crates/ml/src/cuda_pipeline/reward_component_grad_ratio_compute_kernel.cu` +- Create: `crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu` - Create: `crates/ml/tests/sp11_producer_unit_tests.rs` (skeleton + first 3 tests) - Modify: `crates/ml/build.rs` (cubin registration for 3 kernels) - Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (cubin includes, 3 launchers, scratch buffers) @@ -274,10 +281,17 @@ In `crates/ml/build.rs`, find the existing list of `cuda_kernels` (look for `int ```rust // crates/ml/tests/sp11_producer_unit_tests.rs +// +// IMPORTANT: per `feedback_no_htod_htoh_only_mapped_pinned`, tests MUST use +// `MappedF32Buffer` / `MappedI32Buffer` for all CPU↔GPU transfers. Direct +// `htod_copy` / `dtoh_sync_copy` is forbidden even in tests. Pattern below +// matches existing SP4 unit tests (see crates/ml/src/cuda_pipeline/mod.rs:372). + #![cfg(feature = "cuda")] use cudarc::driver::*; use ml::cuda_pipeline::sp11_isv_slots::*; +use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; #[test] fn val_sharpe_delta_first_observation_writes_raw_delta() { @@ -286,21 +300,27 @@ fn val_sharpe_delta_first_observation_writes_raw_delta() { dev.load_ptx(cubin.into(), "module", &["val_sharpe_delta_compute_kernel"]).unwrap(); let f = dev.get_func("module", "val_sharpe_delta_compute_kernel").unwrap(); - let history: Vec = vec![10.0, 15.0]; // delta = 5.0 - let history_dev = dev.htod_copy(history).unwrap(); - let mut scratch_dev = dev.alloc_zeros::(2).unwrap(); - let mut isv_dev = dev.alloc_zeros::(360).unwrap(); + // All buffers via MappedF32Buffer (mapped-pinned host-allocated, DEVICEMAP + // visible). Writes from host go through the device-mapped pointer; reads + // from device come back through the same mapping. No HtoD/DtoH copies. + let mut history = unsafe { MappedF32Buffer::new(2).unwrap() }; + history.host_slice_mut()[0] = 10.0; + history.host_slice_mut()[1] = 15.0; // delta = 5.0 + + let mut scratch = unsafe { MappedF32Buffer::new(2).unwrap() }; + let isv = unsafe { MappedF32Buffer::new(360).unwrap() }; // all zeros unsafe { f.launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (1,1,1), shared_mem_bytes: 0 }, - (&history_dev, &mut scratch_dev, &isv_dev, + (history.device_ptr(), scratch.device_ptr_mut(), isv.device_ptr(), VAL_SHARPE_DELTA_EMA_INDEX as i32, VAL_SHARPE_VAR_EMA_INDEX as i32)) .unwrap(); } + dev.synchronize().unwrap(); - let scratch = dev.dtoh_sync_copy(&scratch_dev).unwrap(); - assert!((scratch[0] - 5.0).abs() < 1e-6, "raw delta should be 5.0, got {}", scratch[0]); - assert!((scratch[1] - 25.0).abs() < 1e-4, "squared deviation against zero ema = 25.0, got {}", scratch[1]); + let scratch_view = scratch.host_slice(); + assert!((scratch_view[0] - 5.0).abs() < 1e-6, "raw delta should be 5.0, got {}", scratch_view[0]); + assert!((scratch_view[1] - 25.0).abs() < 1e-4, "squared deviation against zero ema = 25.0, got {}", scratch_view[1]); } ``` @@ -360,19 +380,19 @@ Expected: clean build, kernel cubin generated. // crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu // SP11 — per-bar saboteur engagement canary producer. // -// For each bar in the current step's batch, computes -// |reward_with_saboteur - reward_without_saboteur| > EPS_ENGAGEMENT +// Consumes a single per-bar `saboteur_delta_reward_buf` produced by the +// saboteur perturbation site itself (see step 8a). For each bar: +// engaged_i = |delta_reward[i]| > EPS_ENGAGEMENT // where EPS_ENGAGEMENT = 0.01 * isv[PNL_REWARD_MAGNITUDE_EMA_INDEX]. -// Reduces per-block via tree-reduce (no atomicAdd per feedback_no_atomicadd), -// then writes batch-aggregated engagement rate to scratch. +// Reduces via block tree-reduce (no atomicAdd per feedback_no_atomicadd), +// writes batch-aggregated engagement rate to scratch. // apply_pearls_ad_kernel chained on stream targets slot SABOTEUR_ENGAGEMENT_RATE_INDEX. #include extern "C" __global__ void saboteur_engagement_compute_kernel( - const float* __restrict__ rewards_with_saboteur, // [batch_size] - const float* __restrict__ rewards_without_saboteur, // [batch_size] - float* __restrict__ scratch_out, // [1]: engagement fraction + const float* __restrict__ saboteur_delta_reward, // [batch_size]: |r_with - r_without| + float* __restrict__ scratch_out, // [1]: engagement fraction const float* __restrict__ isv, const int batch_size, const int pnl_mag_ema_slot) @@ -384,8 +404,7 @@ extern "C" __global__ void saboteur_engagement_compute_kernel( int local_engaged = 0; for (int i = threadIdx.x; i < batch_size; i += blockDim.x) { - const float diff = fabsf(rewards_with_saboteur[i] - rewards_without_saboteur[i]); - if (diff > eps_engagement) local_engaged += 1; + if (fabsf(saboteur_delta_reward[i]) > eps_engagement) local_engaged += 1; } sdata[threadIdx.x] = local_engaged; __syncthreads(); @@ -402,6 +421,24 @@ extern "C" __global__ void saboteur_engagement_compute_kernel( } ``` +- [ ] **Step 8a: Add `saboteur_delta_reward_buf` producer in saboteur perturbation site** + +The saboteur engagement kernel consumes a per-bar `|reward_with - reward_without|` +buffer. This buffer is produced by the saboteur perturbation site itself, +where both reward values are already available locally. In +`crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (saboteur perturbation +section, around line 3328), the existing reward computation has access to +the perturbed price/slippage. Modify the per-bar reward kernel to compute +both `reward_with_saboteur` (using perturbed values) and `reward_without_saboteur` +(using un-perturbed values) in the same kernel — they share state-extraction +work — and emit `delta = reward_with - reward_without` to a new +`MappedF32Buffer` field `saboteur_delta_reward_buf` on the trainer. + +This avoids running two separate reward computations and keeps the diff +local to the kernel that already has both values. Per `feedback_wire_everything_up`, +add the buffer field + its initialization + the producer write site in +this same step (no orphan additions). + - [ ] **Step 9: Register cubin** Append `"saboteur_engagement_compute_kernel"` to `build.rs` cuda_kernels list. @@ -418,26 +455,25 @@ fn saboteur_engagement_counts_above_threshold_only() { let f = dev.get_func("module", "saboteur_engagement_compute_kernel").unwrap(); // PNL EMA = 1.0, so threshold = 0.01. - // 4 bars: diffs = [0.0, 0.005, 0.05, 0.5]. Only 2 of 4 are > 0.01. - let r_with: Vec = vec![1.0, 1.005, 1.05, 1.5]; - let r_no: Vec = vec![1.0, 1.000, 1.00, 1.0]; - let r_with_dev = dev.htod_copy(r_with).unwrap(); - let r_no_dev = dev.htod_copy(r_no).unwrap(); - let mut scratch_dev = dev.alloc_zeros::(1).unwrap(); - let mut isv_dev = dev.alloc_zeros::(360).unwrap(); - dev.dtod_copy(&dev.htod_copy(vec![1.0_f32]).unwrap(), - unsafe { &mut isv_dev.slice_mut(PNL_REWARD_MAGNITUDE_EMA_INDEX..PNL_REWARD_MAGNITUDE_EMA_INDEX+1) }) - .unwrap(); + // 4 bars: |delta_reward| = [0.0, 0.005, 0.05, 0.5]. Only 2 of 4 are > 0.01. + let mut delta_reward = unsafe { MappedF32Buffer::new(4).unwrap() }; + delta_reward.host_slice_mut().copy_from_slice(&[0.0, 0.005, 0.05, 0.5]); + + let mut scratch = unsafe { MappedF32Buffer::new(1).unwrap() }; + + let mut isv = unsafe { MappedF32Buffer::new(360).unwrap() }; + isv.host_slice_mut()[PNL_REWARD_MAGNITUDE_EMA_INDEX] = 1.0; unsafe { f.launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (32,1,1), shared_mem_bytes: 32*4 }, - (&r_with_dev, &r_no_dev, &mut scratch_dev, &isv_dev, + (delta_reward.device_ptr(), scratch.device_ptr_mut(), isv.device_ptr(), 4i32, PNL_REWARD_MAGNITUDE_EMA_INDEX as i32)) .unwrap(); } + dev.synchronize().unwrap(); - let scratch = dev.dtoh_sync_copy(&scratch_dev).unwrap(); - assert!((scratch[0] - 0.5).abs() < 1e-4, "expected engagement=0.5, got {}", scratch[0]); + assert!((scratch.host_slice()[0] - 0.5).abs() < 1e-4, + "expected engagement=0.5, got {}", scratch.host_slice()[0]); } ``` @@ -449,29 +485,33 @@ SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp11_producer_unit Add `launch_saboteur_engagement_compute` launcher in `gpu_dqn_trainer.rs` following the same Pearls A+D chained pattern. Wire into `training_loop.rs` per-step path where rewards are computed (the saboteur perturbation site). -#### A1.3 — reward_component_grad_ratio_compute_kernel +#### A1.3 — reward_component_mag_ratio_compute_kernel - [ ] **Step 12: Write kernel + test + launcher** ```cuda -// crates/ml/src/cuda_pipeline/reward_component_grad_ratio_compute_kernel.cu -// SP11 — per-component grad-ratio canary producer. +// crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu +// SP11 — per-component reward-magnitude-ratio canary producer. // -// Reads 6 existing per-component grad EMAs from ISV (popart, cf, trail, micro, -// opp_cost, bonus). Writes normalized ratios (each / sum) to scratch. -// apply_pearls_ad_kernel chained targeting slots [REWARD_COMPONENT_GRAD_RATIO_BASE..+6). +// Reads 6 existing per-component reward-magnitude EMAs from ISV (popart, cf, +// trail, micro, opp_cost, bonus) at REWARD_POPART_EMA_INDEX..+6. Writes +// normalized ratios (each / sum) to scratch. apply_pearls_ad_kernel chained +// targeting slots [REWARD_COMPONENT_MAG_RATIO_BASE..+6). // -// Note: the source per-component grad EMAs are already populated by SP4's -// reward_component_ema_kernel (REWARD_POPART_EMA_INDEX = 63 + offsets). -// This kernel just reads and renormalizes — no compute beyond a sum + divide. +// Side-output: thread 0 also writes scratch_out[6] = isv[popart_ema_base_slot] +// which apply_pearls_ad chains to PNL_REWARD_MAGNITUDE_EMA_INDEX (slot 359). +// This avoids allocating a separate copy kernel — the popart magnitude IS +// the PnL signal scale used to bound curiosity (spec §3.4.2). +// +// Source EMAs already populated by SP4's reward_component_ema_kernel — this +// kernel just reads and renormalizes (no compute beyond sum + divide). #include -extern "C" __global__ void reward_component_grad_ratio_compute_kernel( +extern "C" __global__ void reward_component_mag_ratio_compute_kernel( const float* __restrict__ isv, - float* __restrict__ scratch_out, // [6] - const int popart_ema_base_slot, // = REWARD_POPART_EMA_INDEX (63) - const int eps_div_idx_unused) // for ABI symmetry; not used (constant inlined) + float* __restrict__ scratch_out, // [7]: 6 ratios + 1 popart mirror + const int popart_ema_base_slot) // = REWARD_POPART_EMA_INDEX (63) { if (blockIdx.x != 0 || threadIdx.x >= 6) return; @@ -485,6 +525,8 @@ extern "C" __global__ void reward_component_grad_ratio_compute_kernel( sum = 0.0f; for (int c = 0; c < 6; c++) sum += values[c]; sum = fmaxf(sum, 1e-6f); + // Mirror popart magnitude to scratch[6] for PNL_REWARD_MAGNITUDE_EMA chain + scratch_out[6] = values[0]; } __syncthreads(); @@ -492,7 +534,15 @@ extern "C" __global__ void reward_component_grad_ratio_compute_kernel( } ``` -Register cubin, write a GPU oracle test (synthetic ISV with 6 known grad EMAs → assert ratios sum to 1.0 ± eps and individual ratios match expected), wire launcher + training_loop call. +Register cubin, write a GPU oracle test using `MappedF32Buffer` for ISV +fixture and scratch (per `feedback_no_htod_htoh_only_mapped_pinned`). Test +fixture: synthetic ISV with 6 known per-component magnitude EMAs at +`REWARD_POPART_EMA_INDEX..+6`; assert (a) `scratch_out[0..6]` sums to 1.0 +± 1e-5, (b) individual ratios match expected values, (c) `scratch_out[6]` +equals `isv[REWARD_POPART_EMA_INDEX]`. Wire launcher: the chained +`apply_pearls_ad_kernel` consumes scratch_out[7] writing to ISV slots +`[REWARD_COMPONENT_MAG_RATIO_BASE..+6) ∪ {PNL_REWARD_MAGNITUDE_EMA_INDEX}` +in a single n_slots=7 launch. Wire `training_loop.rs` per-step launch. - [ ] **Step 13: Build full ml crate; verify all 3 canary kernels** @@ -507,7 +557,7 @@ Expected: 3 tests pass. ```bash git add crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu \ crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu \ - crates/ml/src/cuda_pipeline/reward_component_grad_ratio_compute_kernel.cu \ + crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu \ crates/ml/build.rs \ crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ crates/ml/src/trainers/dqn/trainer/training_loop.rs \ @@ -515,7 +565,7 @@ git add crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu \ git commit -m "$(cat <<'EOF' feat(sp11): A1 — three canary producer kernels (no behavior change) -Adds val_sharpe_delta + saboteur_engagement + reward_component_grad_ratio +Adds val_sharpe_delta + saboteur_engagement + reward_component_mag_ratio GPU producers, each chained with apply_pearls_ad_kernel for Pearls A+D smoothing per feedback_no_cpu_compute_strict. All write to slots [350..360) which no consumer reads yet — training continues identically. @@ -569,7 +619,7 @@ extern "C" __global__ void reward_subsystem_controller_kernel( // saboteur_intensity_mult, weight_floor, curiosity_bound const int delta_ema_slot, const int var_ema_slot, - const int grad_ratio_base_slot, + const int mag_ratio_base_slot, const int saboteur_engagement_slot, const int pnl_mag_ema_slot) { @@ -594,7 +644,7 @@ extern "C" __global__ void reward_subsystem_controller_kernel( __shared__ float curiosity_bound; if (threadIdx.x < 6) { - ratios[threadIdx.x] = isv[grad_ratio_base_slot + threadIdx.x]; + ratios[threadIdx.x] = isv[mag_ratio_base_slot + threadIdx.x]; } __syncthreads(); @@ -669,30 +719,119 @@ Append `"reward_subsystem_controller_kernel"` to `build.rs`. - [ ] **Step 3: Write GPU oracle tests** -In `sp11_producer_unit_tests.rs`, add: +In `sp11_producer_unit_tests.rs`, add (all fixtures via `MappedF32Buffer` +per `feedback_no_htod_htoh_only_mapped_pinned`): ```rust +fn load_controller_kernel() -> (CudaDevice, CudaFunction) { + let dev = CudaDevice::new(0).unwrap(); + let cubin = include_bytes!(concat!(env!("OUT_DIR"), "/reward_subsystem_controller_kernel.cubin")); + dev.load_ptx(cubin.into(), "module", &["reward_subsystem_controller_kernel"]).unwrap(); + let f = dev.get_func("module", "reward_subsystem_controller_kernel").unwrap(); + (dev, f) +} + +fn populate_controller_isv( + isv: &mut MappedF32Buffer, + delta: f32, var: f32, 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] = ratios[c]; + } + h[SABOTEUR_ENGAGEMENT_RATE_INDEX] = saboteur_engagement; + h[PNL_REWARD_MAGNITUDE_EMA_INDEX] = pnl_mag; +} + +fn launch_controller( + dev: &CudaDevice, f: &CudaFunction, + isv: &MappedF32Buffer, scratch: &mut MappedF32Buffer, +) { + unsafe { + f.launch( + LaunchConfig { grid_dim: (1,1,1), block_dim: (10,1,1), shared_mem_bytes: 0 }, + (isv.device_ptr(), scratch.device_ptr_mut(), + VAL_SHARPE_DELTA_EMA_INDEX as i32, + VAL_SHARPE_VAR_EMA_INDEX as i32, + REWARD_COMPONENT_MAG_RATIO_BASE as i32, + SABOTEUR_ENGAGEMENT_RATE_INDEX as i32, + PNL_REWARD_MAGNITUDE_EMA_INDEX as i32), + ).unwrap(); + } + dev.synchronize().unwrap(); +} + #[test] fn controller_z_score_at_zero_yields_midpoint_outputs() { - // delta_ema=0, delta_var=1 → dz=0 → sigmoid(0)=0.5 → improving=stagnant=0.5 - // ratios all equal 1/6 → blend = 0.5*1/6 + 0.5*5/6/5 = 1/12 + 1/12 = 1/6 - // post-floor (0.5 × 1/6 = 0.083), max(0.083, 1/6) = 1/6 - // sum = 1.0, normalized weights all = 1/6 - // curiosity_bound = pnl_mag × 0.3, pnl_mag = 1.0 → bound = 0.3 - // curiosity_floor = 0.2 × 0.3 = 0.06; curiosity_dynamic = 0.5 × 0.3 = 0.15 → max = 0.15 - // saboteur = (0.5 + 1.5×0.5) × max(0.0, 0.1) = 1.25 × 0.1 = 0.125, clamped to 0.5 - // ... full assertions + // delta=0, var=1 → dz=0 → sigmoid(0)=0.5 → improving=stagnant=0.5 + // All 6 ratios = 1/6, post-floor renorm preserves uniform weights. + // curiosity_bound = pnl_mag × 0.3 = 0.3; curiosity_floor = 0.06; + // curiosity_dynamic = 0.5 × 0.3 = 0.15 → max = 0.15 + // saboteur intensity_z = 0.5 + 1.5 × 0.5 = 1.25; engagement = 0.5 + // modulated = 1.25 × 0.5 = 0.625; post-clamp = 0.625 (in range) + let (dev, f) = load_controller_kernel(); + let mut isv = unsafe { MappedF32Buffer::new(360).unwrap() }; + populate_controller_isv(&mut isv, 0.0, 1.0, [1.0/6.0; 6], 0.5, 1.0); + let mut scratch = unsafe { MappedF32Buffer::new(10).unwrap() }; + launch_controller(&dev, &f, &isv, &mut scratch); + + let s = scratch.host_slice(); + let weight_sum: f32 = s[0..6].iter().sum(); + assert!((weight_sum - 1.0).abs() < 1e-5, "weights sum = {}, expected 1.0", weight_sum); + for c in 0..6 { + assert!((s[c] - 1.0/6.0).abs() < 1e-4, "weight[{}] = {}, expected 1/6", c, s[c]); + } + assert!((s[6] - 0.15).abs() < 1e-4, "curiosity_pressure = {}, expected 0.15", s[6]); + assert!((s[7] - 0.625).abs() < 1e-3, "saboteur = {}, expected 0.625", s[7]); + assert!((s[9] - 0.3).abs() < 1e-5, "curiosity_bound = {}, expected 0.3", s[9]); } #[test] fn controller_weights_renormalize_after_floor() { - // pathological ratio array (one near-zero) → ensure post-floor sum normalizes to 1.0 + // Pathological ratios: one large, five near-zero. Floor activates, + // but renormalization restores Σ=1. + let (dev, f) = load_controller_kernel(); + let mut isv = unsafe { MappedF32Buffer::new(360).unwrap() }; + populate_controller_isv(&mut isv, 1.0, 1.0, + [0.95, 0.01, 0.01, 0.01, 0.01, 0.01], 0.5, 1.0); + let mut scratch = unsafe { MappedF32Buffer::new(10).unwrap() }; + launch_controller(&dev, &f, &isv, &mut scratch); + + let s = scratch.host_slice(); + let weight_sum: f32 = s[0..6].iter().sum(); + assert!((weight_sum - 1.0).abs() < 1e-5, + "weights sum = {} after renormalization, expected 1.0", weight_sum); + // Each weight at least the WEIGHT_HARD_FLOOR (0.01) post-renorm + for c in 0..6 { + assert!(s[c] >= 0.01 - 1e-5, "weight[{}] = {} below hard floor", c, s[c]); + } } #[test] -fn controller_saboteur_clamped_at_min() { - // engagement=0 (post-floor 0.1), improving=0 → intensity_z=0.5, modulated=0.05, - // post-clamp=0.5 (SABOTEUR_MIN). Verify clamp catches it. +fn controller_saboteur_post_clamp_holds_min() { + // engagement=0 (will be floored to 0.1), improving very low. + // delta=-100, var=1 → dz=-100 → sigmoid≈0 → improving≈0 → stagnant≈1 + // intensity_z = 0.5 + 1.5×0 = 0.5 + // modulated = 0.5 × max(0.0, 0.1) = 0.05 < SABOTEUR_MIN + // post-clamp = 0.5 (SABOTEUR_MIN) + let (dev, f) = load_controller_kernel(); + let mut isv = unsafe { MappedF32Buffer::new(360).unwrap() }; + populate_controller_isv(&mut isv, -100.0, 1.0, [1.0/6.0; 6], 0.0, 1.0); + let mut scratch = unsafe { MappedF32Buffer::new(10).unwrap() }; + launch_controller(&dev, &f, &isv, &mut scratch); + + let s = scratch.host_slice(); + assert!((s[7] - 0.5).abs() < 1e-4, + "saboteur = {}, expected SABOTEUR_MIN=0.5 after post-clamp", s[7]); + + // Curiosity should be near max (stagnant_or_worse ≈ 1) + // curiosity_bound = 0.3, curiosity_dynamic ≈ 0.3, curiosity_floor = 0.06 + // pressure = max(0.3, 0.06) = 0.3 + assert!((s[6] - 0.3).abs() < 1e-3, + "curiosity_pressure at max regression = {}, expected ≈ 0.3", s[6]); } ``` @@ -711,17 +850,75 @@ In `gpu_dqn_trainer.rs`, add include + launcher method `launch_reward_subsystem_ - [ ] **Step 6: Add novelty hash buffer field to GpuDqnTrainer** -In `gpu_dqn_trainer.rs`, add: +Per `feedback_no_cpu_forwards` ("CPU is read-only"), the projection +matrix MUST be populated by a GPU kernel (Philox-driven), NOT by host +RNG writing through `host_slice_mut`. The hash buffer is allocated +zero-initialized (cuMemHostAlloc DEVICEMAP zero-fills); the projection +matrix is allocated then immediately filled by a one-shot GPU init +kernel reading a seed scalar. + +In `gpu_dqn_trainer.rs`, add the field + the init-kernel cubin: + ```rust -/// SP11 novelty visit-count hash table. 1M slots × f32 (4 MB). +/// SP11 novelty visit-count hash table. 1M slots × f32 (4 MB mapped-pinned). /// SimHash projection 42×16 → 16-bit code; bucket = (action_id, code). -/// Reset at fold boundary via state-reset registry novelty arm. -novelty_hash_buf: cudarc::driver::CudaSlice, -/// SP11 SimHash projection matrix (42×16 = 672 floats, fixed at trainer init). -novelty_simhash_proj: cudarc::driver::CudaSlice, +/// Reset at fold boundary via state-reset registry novelty arm +/// (zero-init kernel writes through device pointer; no DtoH/HtoD). +novelty_hash_buf: mapped_pinned::MappedF32Buffer, +/// SP11 SimHash projection matrix (42×16 = 672 floats). Populated ONCE at +/// trainer construct by `novelty_simhash_proj_init_kernel` (Philox-driven +/// from `config.seed`). Never re-initialized; never written from host. +novelty_simhash_proj: mapped_pinned::MappedF32Buffer, ``` -Initialize in trainer constructor: `novelty_hash_buf = device.alloc_zeros::(1 << 20)?;` and the projection matrix as 42×16 random ±1 values seeded deterministically (use a fixed seed; the trainer config already has a `seed` field). +Add a small GPU init kernel `novelty_simhash_proj_init_kernel.cu` (≤30 LOC): + +```cuda +// crates/ml/src/cuda_pipeline/novelty_simhash_proj_init_kernel.cu +// SP11 — one-shot GPU init for the SimHash projection matrix. +// Philox(seed) → uniform u32 → low bit picks +1 or -1. +// Per `feedback_no_cpu_forwards` the projection MUST be GPU-populated. + +#include +#include "philox_inline.cuh" // existing project helper + +extern "C" __global__ void novelty_simhash_proj_init_kernel( + float* __restrict__ proj, // [42 × 16 = 672] + const unsigned long long seed, + const int n_floats) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= n_floats) return; + const unsigned int u = philox_uint(seed, (unsigned int)i, 0u); + proj[i] = (u & 1u) ? 1.0f : -1.0f; +} +``` + +Initialize in trainer constructor: +```rust +let novelty_hash_buf = unsafe { mapped_pinned::MappedF32Buffer::new(1 << 20)? }; +// cuMemHostAlloc DEVICEMAP zero-initializes — matches state-reset semantics. + +let novelty_simhash_proj = unsafe { mapped_pinned::MappedF32Buffer::new(42 * 16)? }; +// One-shot GPU init using a Philox stream seeded from config.seed. +// No host write; the matrix is populated on-device. +self.launch_novelty_simhash_proj_init( + &novelty_simhash_proj, + config.seed.wrapping_add(SP11_PROJ_SEED_SALT), + &self.stream, +)?; +``` + +Where `SP11_PROJ_SEED_SALT: u64 = 0x_5511_0001` is defined in +`sp11_isv_slots.rs` alongside the other SP11 constants. The launcher +method `launch_novelty_simhash_proj_init` runs the init kernel with +grid 21 × block 32 (covers 672 elements). + +Register cubin in `build.rs`. The kernels read/write through +`device_ptr()` / `device_ptr_mut()`. No `htod_copy`, no `dtoh_sync_copy`, +no `alloc_zeros`, no host RNG. Per `feedback_wire_everything_up` the +init kernel is wired to the constructor that allocates the buffer — +no orphan addition. - [ ] **Step 7: Write SimHash kernel** @@ -782,11 +979,13 @@ extern "C" __global__ void novelty_simhash_update_kernel( float* __restrict__ hash_table, const int batch_size) { - // Each thread processes one (state, action). Linear-probe-on-conflict not - // needed because we're EMA-style updating (collisions just blend counts). - // Per feedback_no_atomicadd: serialize per bucket via deterministic - // single-thread-per-bucket. Use a simple non-atomic increment (race accepted — - // count drift from races is bounded; semantic still "approximate visit count"). + // Each thread processes one (state, action). Per `feedback_no_atomicadd`, + // we forbid atomicAdd. Concurrent writers to the same bucket race; + // last-writer-wins loses increments. This biases counts DOWNWARD, which + // biases novelty UPWARD (1/sqrt(1+count) is decreasing in count). + // Over-novelty is the SAFE direction — the policy explores MORE, never + // masks an already-visited state-action as novel. Document the safety + // argument in this header (consumers depend on it). const int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= batch_size) return; @@ -796,7 +995,11 @@ extern "C" __global__ void novelty_simhash_update_kernel( } ``` -NOTE on no-atomicAdd compliance: the update kernel uses a non-atomic write under "approximate visit count" semantics. Collisions from concurrent thread races simply lose increments, biasing novelty *upward* (the safe direction — overestimates novelty rather than masking it). Document this explicitly in the kernel header comment. +NOTE on `feedback_no_atomicadd` compliance: the update kernel uses a +non-atomic write. Race-driven count under-counts bias novelty UPWARD, +which is the safe direction (model explores more, never less). Spec +§3.5.2 explicitly accepts this trade-off. The header comment must +preserve this safety argument. - [ ] **Step 8: Register cubin + launcher methods** @@ -808,7 +1011,7 @@ In `training_loop.rs`, in the per-step section (after rewards and grad EMAs are ```rust // SP11 — populate canary slots, then run controller -self.gpu.launch_reward_component_grad_ratio_compute(&self.stream)?; +self.gpu.launch_reward_component_mag_ratio_compute(&self.stream)?; self.gpu.launch_saboteur_engagement_compute(/*reward bufs*/, &self.stream)?; // val_sharpe_delta only fires at val emit boundaries (Task A1.1 wired separately) // pnl_reward_magnitude_ema is updated via existing reward_component_ema_kernel — no new launch needed @@ -933,42 +1136,107 @@ For each remaining site from `/tmp/sp11_audit.txt`, replace the hardcoded weight If the audit finds NO additional sites for popart/trail/micro/opp_cost/bonus, that means those reward components currently bypass an explicit weight — they enter the reward sum directly. In that case, modify the reward composition site (search for `reward_total = popart + cf + trail + ...`) to multiply each component by its ISV weight before summing. -- [ ] **Step 5: Migrate saboteur scale** +- [ ] **Step 5: Migrate saboteur scale (GPU-side multiplication only)** -In `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs:3328`, find: -```rust -let ps = self.saboteur_perturbation_scale; -``` -Replace with: -```rust -// SP11: saboteur intensity multiplier from controller output -let mult = self.gpu.read_isv_slot(SABOTEUR_INTENSITY_MULT_INDEX); -let ps = self.saboteur_perturbation_scale * mult; +Per `feedback_no_cpu_compute_strict` AND `feedback_no_cpu_forwards` +("CPU is read-only"), the saboteur multiplication MUST happen inside the +saboteur perturbation kernel — NOT on CPU. Do NOT introduce a host-side +`read_isv_slot` + Rust multiplication. + +In `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` at the +saboteur perturbation kernel call site (~line 3328), change the kernel +launch to pass: +1. `saboteur_perturbation_scale` as the existing scalar parameter (f32), +2. The ISV device pointer + `SABOTEUR_INTENSITY_MULT_INDEX as i32`, + +and modify the saboteur perturbation kernel itself to compute: +```cuda +// SP11 — multiplier comes from controller output (ISV slot 347), NOT host. +const float mult = isv[saboteur_intensity_mult_slot]; +const float effective_scale = base_scale * mult; +// ... apply effective_scale to fill/slippage perturbation ... ``` -The `read_isv_slot` helper must use the existing mapped-pinned ISV mirror (it already exists per SP4 patterns; if not, add a `MappedF32Buffer` mirror that's DtoH'd at epoch boundaries — never per-step). For per-bar reads, prefer pulling the value into a local `MappedF32Buffer` once per step before the saboteur loop. +The existing `saboteur_perturbation_scale` Rust f32 stays on the trainer +struct as the BASE rate; the ISV multiplier is the adaptive curriculum +modulator on top. Both arrive at the kernel as parameters; multiplication +happens on-device. -- [ ] **Step 6: Add replay-time curiosity** +If the existing saboteur perturbation kernel signature does not yet +take an ISV pointer, add it to the signature, update the launcher, and +update any other call sites. Per `feedback_no_partial_refactor` this +signature change is part of the same Layer B atomic commit. -Find the replay sampling kernel. Search for `replay_sample` or `sample_batch` in `crates/ml/src/cuda_pipeline/`. Inside the kernel that produces the per-replayed-bar reward, add curiosity: +- [ ] **Step 6: Add replay-time curiosity in `graph_utility_kernels.cu`** + +The replay-time reward gather lives in +`crates/ml/src/cuda_pipeline/graph_utility_kernels.cu:71` +(`gather_f32_scalar`, the kernel that pulls rewards/dones/IS-weights +from the replay buffer by index). Add a sibling kernel +`gather_replay_reward_with_curiosity` that: ```cuda -// SP11 — recompute curiosity at replay time (NOT collect time) to avoid -// stale-signal contamination of the policy. -const float curiosity_p = isv[CURIOSITY_PRESSURE_INDEX]; -// novelty pre-computed per-tuple via novelty_simhash_lookup kernel -const float novelty = novelty_signals[i]; -const float r_with_curiosity = r_base[i] + curiosity_p * novelty; +// crates/ml/src/cuda_pipeline/graph_utility_kernels.cu — append. +// SP11 — replay-time curiosity bonus. +// +// Recompute curiosity at REPLAY time (NOT collect time) per spec §3.5.1 +// to bind exploration pressure to current visit-count + current ISV[346], +// not stale at-collect values. Stale-curiosity replay would lock the +// policy into the ep1 fixed point we are trying to escape. + +extern "C" __global__ void gather_replay_reward_with_curiosity( + const float* __restrict__ reward_buf, // base rewards in replay buffer + const int* __restrict__ replay_indices, // [batch_size] + const float* __restrict__ novelty_signals, // [batch_size] — pre-filled by lookup + const float* __restrict__ isv, + float* __restrict__ reward_out, // [batch_size] + const int batch_size, + const int curiosity_pressure_slot) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= batch_size) return; + + const float r_base = reward_buf[replay_indices[i]]; + const float curiosity_p = isv[curiosity_pressure_slot]; + reward_out[i] = r_base + curiosity_p * novelty_signals[i]; +} ``` -In the replay batch dispatch, schedule the `novelty_simhash_lookup_kernel` immediately before reward computation and the `novelty_simhash_update_kernel` immediately after. The update happens AFTER lookup so this batch sees the pre-update count. +In the replay batch dispatch (Rust side, in `gpu_dqn_trainer.rs` near +the existing `gather_f32_scalar` launch for rewards), the new sequence is: + +```rust +// 1. Lookup current novelty for each (state, action) in the replay batch. +self.launch_novelty_simhash_lookup( + &replay_states_buf, &replay_actions_buf, + &self.novelty_simhash_proj, &self.novelty_hash_buf, + &mut self.novelty_signals_buf, batch_size, &self.stream)?; + +// 2. Gather replay reward + curiosity bonus (replaces existing scalar gather). +self.launch_gather_replay_reward_with_curiosity( + &replay_reward_buf, &replay_indices_buf, + &self.novelty_signals_buf, + &mut reward_out, batch_size, &self.stream)?; + +// 3. Update novelty counts AFTER lookup so this batch saw the pre-update count. +self.launch_novelty_simhash_update( + &replay_states_buf, &replay_actions_buf, + &self.novelty_simhash_proj, &mut self.novelty_hash_buf, + batch_size, &self.stream)?; +``` + +The pre-existing scalar reward gather call is REPLACED, not augmented — +per `feedback_no_legacy_aliases` the old call is deleted, not kept as a +fallback path. - [ ] **Step 7: Run `cargo check`** ```bash SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -20 ``` -Expected: clean compile. If `read_isv_slot` doesn't exist, add it (mapped-pinned read). +Expected: clean compile. The saboteur kernel signature change is the +most likely break-site if any other call site exists; update all +callers in the same commit per `feedback_no_partial_refactor`. - [ ] **Step 8: Run cargo test** @@ -1104,7 +1372,7 @@ type: pearl Every degree of freedom in the reward path — component weights, shaping multipliers, exploration bonuses, curriculum knobs — is an output of a unified controller whose inputs are improvement-rate canaries (val sharpe -delta + variance, per-component grad ratios, saboteur engagement, PnL +delta + variance, per-component magnitude ratios, saboteur engagement, PnL magnitude). The controller's classification of regime ("improving / stagnant / regressing") is encoded in a Z-score and sigmoid; no hardcoded threshold survives at the regime layer.