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

plan(sp11): patch all 5 violations + 2 gaps from review

Browse Source 
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<f32> + 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.
This commit is contained in:
jgrusewski
2026-05-04 00:40:36 +02:00
parent d8b44e0829
commit e0d3abd9d2
1 changed files with 375 additions and 107 deletions
Download Patch File Download Diff File Expand all files Collapse all files

482
docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md
View File

@@ -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<f32> = vec![10.0, 15.0]; // delta = 5.0
let history_dev = dev.htod_copy(history).unwrap();
let mut scratch_dev = dev.alloc_zeros::<f32>(2).unwrap();
let mut isv_dev = dev.alloc_zeros::<f32>(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 <cuda_runtime.h>
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<f32> = vec![1.0, 1.005, 1.05, 1.5];
let r_no: Vec<f32> = 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::<f32>(1).unwrap();
let mut isv_dev = dev.alloc_zeros::<f32>(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 <cuda_runtime.h>
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<f32>,
/// SP11 SimHash projection matrix (42×16 = 672 floats, fixed at trainer init).
novelty_simhash_proj: cudarc::driver::CudaSlice<f32>,
/// 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::<f32>(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 <cuda_runtime.h>
#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.