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

feat(sp11): A1 — three canary producer kernels (no behavior change)

Browse Source 
Adds val_sharpe_delta + saboteur_engagement + reward_component_mag_ratio
GPU producers for the SP11 reward-as-controlled-subsystem chain. Each is
a single-block producer chained with apply_pearls_ad_kernel for Pearls
A+D smoothing per pearl_first_observation_bootstrap.md +
pearl_wiener_optimal_adaptive_alpha.md. All three write to slots in
[350..360) which no consumer reads yet — Layer A is additive; consumer
migration lands atomically in Layer B.

A1.1 — val_sharpe_delta_compute_kernel.cu
  Two-pass: writes raw delta + (delta - prev_delta_ema)^2 to scratch.
  Chained Pearls A+D (n_slots=2) → ISV[VAL_SHARPE_DELTA_EMA_INDEX=350,
  VAL_SHARPE_VAR_EMA_INDEX=351]. Host writes val_sharpe to mapped-pinned
  history[1]; rotation handled in training_loop.rs at val emit boundary
  (a literal already-computed value — no host-side compute, no htod_copy).

A1.2 — saboteur_engagement_compute_kernel.cu
  Per-bar |Δreward| > 0.01 × ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX] check
  with block tree-reduce (no atomicAdd per feedback_no_atomicadd). The
  per-bar Δreward signal is produced by experience_env_step's saboteur
  perturbation site as `traded × |reward| × max(|eff_spread − 1|,
  |eff_slip − 1|)` — a structural proxy for the cost-differential the
  saboteur imposed on bars where the model traded. Single kernel-side
  emit (no parallel reward computation), per spec §3.3.1.
  Chained Pearls A+D → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].

A1.3 — reward_component_mag_ratio_compute_kernel.cu
  Reads ISV[REWARD_POPART_EMA_INDEX..+6) (the SP4 reward-component
  magnitude EMAs), normalises to ratios, and mirrors popart magnitude
  into scratch[6] as a side-output. ONE non-pointer parameter
  (popart_ema_base_slot) — no _unused param per feedback_no_stubs.
  Two chained Pearls A+D launches:
    n_slots=6 → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6)
    n_slots=1 → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
  (slots non-contiguous: 352..358 then 359.)

Wire-up (per feedback_wire_everything_up):
- 3 cubin entries appended to crates/ml/build.rs
- 3 kernel handles + val_sharpe_history_pinned (MappedF32Buffer[2]) +
  saboteur_delta_reward dev-ptr cache fields on GpuDqnTrainer
- 3 launchers (launch_sp11_*) + 1 setter (set_sp11_saboteur_delta_reward_buf)
- saboteur_delta_reward_per_sample buffer field on GpuExperienceCollector
- experience_env_step kernel signature extended with the new buffer arg;
  every call site in the same commit per feedback_no_partial_refactor
- training_loop.rs init wires collector→trainer setter; val emit boundary
  invokes launch_sp11_val_sharpe_delta_compute; per-epoch metrics block
  invokes launch_sp11_mag_ratio_compute then
  launch_sp11_saboteur_engagement_compute (mag_ratio first so the
  signal-relative threshold base is populated before the saboteur reader)
- SP5_SCRATCH_TOTAL grown 266 → 276 (10 new scratch slots: 2+1+7)
- docs/isv-slots.md SP11 section updated to reflect A1 producers

3 GPU oracle tests in crates/ml/tests/sp11_producer_unit_tests.rs
pass on RTX 3050 Ti via MappedF32Buffer fixtures (zero htod_copy /
dtoh_sync_copy / alloc_zeros — feedback_no_htod_htoh_only_mapped_pinned
compliant).

Note on Step 8a path: the plan offered two routes for the saboteur
Δreward producer — in-kernel diff emission OR a small dedicated
reader-of-existing-buffers. The existing reward path emits ONE reward
(not both with/without), so the dedicated-reader alternative was
infeasible. The in-kernel emission landed as a small write site at the
END of experience_env_step (after total_reward_per_sample is finalised),
threading saboteur_eff_spread/saboteur_eff_slip from the perturbation
site forward to the END via stack vars. Single new kernel parameter,
single new GPU-only buffer, single existing call site updated.

Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md §5
Plan: docs/superpowers/plans/2026-05-04-sp11-reward-as-controlled-subsystem.md (Task A1)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-04 01:40:38 +02:00
parent 201b59dfbc
commit 91b48bc7a5
10 changed files with 1218 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
crates/ml/build.rs
View File

@@ -548,6 +548,22 @@ fn main() {
"intent_eval_divergence_compute_kernel.cu",
"q_var_mag_ema_compute_kernel.cu",
"kelly_warmup_floor_compute_kernel.cu",
// SP11 Fix 39 (2026-05-04, Task A1): three canary producer kernels
// for the reward-as-controlled-subsystem chain. Each is a single-
// block producer chained through `apply_pearls_ad_kernel` for
// Pearl A sentinel-bootstrap + Pearl D Wiener-α steady-state
// smoothing. All three write to slots in [350..360) which no
// consumer reads yet (Layer A is additive; consumer migration
// lands atomically in Layer B). See spec §3.3 (canary slots) and
// §5 (producer kernel listing).
// 1. val_sharpe_delta_compute → ISV[VAL_SHARPE_DELTA_EMA_INDEX=350,
// VAL_SHARPE_VAR_EMA_INDEX=351]
// 2. saboteur_engagement_compute → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358]
// 3. reward_component_mag_ratio → ISV[REWARD_COMPONENT_MAG_RATIO_BASE=352..358)
// ∪ ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
"val_sharpe_delta_compute_kernel.cu",
"saboteur_engagement_compute_kernel.cu",
"reward_component_mag_ratio_compute_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

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

@@ -1761,7 +1761,29 @@ extern "C" __global__ void experience_env_step(
* already produced by q_stats_kernel.cu — no new producer needed.
* Per pearl_one_unbounded_signal_per_reward.md +
* feedback_isv_for_adaptive_bounds.md. */
int q_dir_abs_ref_idx
int q_dir_abs_ref_idx,
/* SP11 Fix 39 (2026-05-04, A1.2): per-bar saboteur engagement signal
* for the reward-as-controlled-subsystem chain. Layout [N*L] f32, row-
* major (i, t). Producer: this kernel writes the engagement proxy at
* the END (after `total_reward_per_sample` is finalised) — a
* structurally-meaningful approximation of |Δreward_with - Δreward_without|
* per spec §3.3.1:
* delta = traded × |reward| × max(|eff_spread - 1|, |eff_slip - 1|)
* where `eff_spread`/`eff_slip` are the saboteur-perturbed cost
* multipliers and `traded` is 1 iff the bar carried a trade. The
* kernel `saboteur_engagement_compute_kernel` consumes this buffer
* and applies a `0.01 × ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX]`
* threshold to classify per-bar engagement, then block tree-reduces
* to a fraction → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358] via Pearls
* A+D. NULL = saboteur disabled this epoch (writes zero per bar
* → engagement collapses to zero, Pearl A holds the EMA at sentinel
* until the saboteur reactivates).
*
* Pure GPU buffer (CudaSlice<f32> on the host side), never read
* host-side — kernel-to-kernel dataflow only, so the no-CudaSlice
* rule from feedback_no_htod_htoh_only_mapped_pinned does not apply
* (that rule covers CPU↔GPU paths). */
float* __restrict__ saboteur_delta_reward_per_sample
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -1777,6 +1799,15 @@ extern "C" __global__ void experience_env_step(
exploration_scale = fminf(1.0f, fmaxf(0.0f, exploration_scale));
shaping_scale = fminf(1.0f, fmaxf(0.0f, shaping_scale));
/* SP11 Fix 39 (A1.2): track the saboteur perturbation magnitude per
* thread so the END-of-kernel engagement-signal write can compute
* `traded × |reward| × max(|eff_spread − 1|, |eff_slip − 1|)`. The
* non-saboteur path leaves these at 1.0 (identity perturbation) so
* the engagement signal naturally collapses to zero — Pearl A holds
* the EMA at sentinel until the saboteur activates. */
float saboteur_eff_spread = 1.0f;
float saboteur_eff_slip = 1.0f;
/* #33 Per-episode saboteur overrides — lerped by exploration_scale.
* sab_eff = 1 + exploration_scale * (sab - 1) → identity at scale=0,
* full saboteur at scale=1. Same lerp form for both spread and slippage. */
@@ -1792,6 +1823,8 @@ extern "C" __global__ void experience_env_step(
* portfolio sim via trade_physics.cuh. The exploration_scale gating
* here covers spread + slippage; fill probability stays at base. */
(void)sab_fill; /* still read for register layout consistency */
saboteur_eff_spread = eff_spread;
saboteur_eff_slip = eff_slip;
}
/* G2: Transaction cost curriculum — scale spread_cost and tx_cost_multiplier */
@@ -1829,6 +1862,14 @@ extern "C" __global__ void experience_env_step(
micro_reward_per_sample[out_off] = 0.0f;
total_reward_per_sample[out_off] = 0.0f;
var_scale_per_sample[out_off] = 0.0f;
/* SP11 Fix 39 (A1.2): default per-bar saboteur engagement signal to 0.
* Real value is overwritten near `total_reward_per_sample[out_off] = reward`
* after the trade lifecycle resolves. NULL pointer = saboteur disabled
* for this epoch's collect call (caller passes 0 from the launcher); the
* kernel guards the write with the NULL check. */
if (saboteur_delta_reward_per_sample != NULL) {
saboteur_delta_reward_per_sample[out_off] = 0.0f;
}
{
long long cf_off_default = out_off + (long long)N * L;
cf_flip_per_sample[cf_off_default] = 0;
@@ -3176,6 +3217,28 @@ extern "C" __global__ void experience_env_step(
/* C.2: component [0] = final on-policy reward (PopArt input denominator). */
reward_components_per_sample[out_off * 6 + 0] = reward;
/* SP11 Fix 39 (A1.2): per-bar saboteur engagement signal — see param
* docstring above. Computed after `total_reward_per_sample` is
* finalised. The proxy `traded × |reward| × max(|eff_spread − 1|,
* |eff_slip − 1|)` is non-zero exactly when (a) the bar carried a
* trade (so the saboteur perturbation actually altered transaction
* cost in this bar's reward) AND (b) the saboteur perturbed the
* cost multipliers (eff_* != 1.0). The downstream kernel applies a
* `0.01 × ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX]` threshold to classify
* engagement and reduces to a per-step fraction.
*
* Note: the raw-portfolio-return path (raw_returns_out) is NOT
* affected by saboteur perturbations — those go into shaped reward
* only. So this signal is a meaningful proxy specifically for
* "saboteur affected the LEARNING signal", which is what the
* controller cares about. */
if (saboteur_delta_reward_per_sample != NULL) {
float perturb_mag = fmaxf(fabsf(saboteur_eff_spread - 1.0f),
fabsf(saboteur_eff_slip - 1.0f));
float traded_f = (float)(traded_per_sample[out_off]);
saboteur_delta_reward_per_sample[out_off] = traded_f * fabsf(reward) * perturb_mag;
}
/* ---- Write RAW portfolio return (unshaped) for accurate Sharpe/MaxDD ---- */
/* True fractional return: (equity_t - equity_{t-1}) / equity_{t-1}.
* No dense/sparse weighting, no soft-clamp, no CEA bonus.

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

@@ -522,6 +522,37 @@ static SP9_Q_VAR_MAG_EMA_COMPUTE_CUBIN: &[u8] =
static SP9_KELLY_WARMUP_FLOOR_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/kelly_warmup_floor_compute_kernel.cubin"));
/// SP11 Fix 39 (2026-05-04, Task A1): val-sharpe Δ + variance canary
/// producer. Single block, single thread; reads a 2-element mapped-pinned
/// `val_sharpe_history` ([prev_epoch, curr_epoch]) plus
/// ISV[VAL_SHARPE_DELTA_EMA_INDEX]; writes `(curr-prev)` and
/// `(delta - prev_delta_ema)^2` to scratch[SCRATCH_SP11_VAL_SHARPE_DELTA_BASE..+2).
/// Downstream `apply_pearls_ad_kernel` (n_slots=2) smooths into
/// ISV[VAL_SHARPE_DELTA_EMA_INDEX=350, VAL_SHARPE_VAR_EMA_INDEX=351].
static SP11_VAL_SHARPE_DELTA_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/val_sharpe_delta_compute_kernel.cubin"));
/// SP11 Fix 39 (2026-05-04, Task A1): per-bar saboteur engagement canary
/// producer. Single block, 256 threads + grid-stride; reads
/// `saboteur_delta_reward_buf` populated by `experience_env_step` at the
/// saboteur perturbation site, computes per-bar engagement
/// (`|delta| > 0.01 × ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX]`), block tree-
/// reduces to a fraction. Downstream `apply_pearls_ad_kernel` (n_slots=1)
/// smooths into ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
static SP11_SABOTEUR_ENGAGEMENT_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/saboteur_engagement_compute_kernel.cubin"));
/// SP11 Fix 39 (2026-05-04, Task A1): per-component reward-magnitude-ratio
/// canary producer. Single block, 6 threads; reads 6 contiguous component
/// magnitude EMAs from ISV[REWARD_POPART_EMA_INDEX..+6) (populated by
/// SP4's reward_component_ema_kernel), normalises to ratios, and mirrors
/// the popart magnitude into scratch[6] as a side-output. Downstream
/// `apply_pearls_ad_kernel` smooths the 6 ratios into
/// ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6) and the singleton mirror
/// into ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359].
static SP11_MAG_RATIO_COMPUTE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/reward_component_mag_ratio_compute_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
@@ -1197,8 +1228,14 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
/// require no scratch storage. The eval_dist base slides 265 → 262.
/// SP10 (Fix 38, 2026-05-03) adds:
/// 1 float for eval Thompson temp (SCRATCH_SP10_THOMPSON_TEMP=265, at [265..266))
/// Combined: 259 + 6 + 1 = 266 scratch slots [0..266).
pub const SP5_SCRATCH_TOTAL: usize = 266;
/// SP11 (Fix 39, 2026-05-04, A1) adds:
/// 2 floats for val_sharpe_delta_compute (SCRATCH_SP11_VAL_SHARPE_DELTA_BASE=266, [266..268))
/// 1 float for saboteur_engagement_compute (SCRATCH_SP11_SABOTEUR_ENGAGEMENT=268, [268..269))
/// 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;
/// SP5 Layer D Task D1 (rewrite, 2026-05-02): scratch index base for
/// `pnl_aggregation_update`.
@@ -1311,6 +1348,19 @@ pub const SCRATCH_SP9_EVAL_DIST_BASE: usize = 262; // → ISV[EVAL_DIS
/// ISV[EVAL_THOMPSON_TEMP_INDEX=339].
pub const SCRATCH_SP10_THOMPSON_TEMP: usize = 265; // → ISV[EVAL_THOMPSON_TEMP_INDEX=339]
/// SP11 (Fix 39, 2026-05-04, Task A1): scratch slots for the three canary
/// producers in the reward-as-controlled-subsystem chain. Each producer
/// writes to its scratch block; downstream `apply_pearls_ad_kernel`
/// smooths into the corresponding ISV slots in [350..360) via Pearl A
/// sentinel-bootstrap + Pearl D Wiener-α steady-state smoothing.
///
/// 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
/// (Task A1 Steps 1-13).
pub const SCRATCH_SP11_VAL_SHARPE_DELTA_BASE: usize = 266; // [266..268) → ISV[VAL_SHARPE_DELTA_EMA_INDEX=350, VAL_SHARPE_VAR_EMA_INDEX=351]
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]
/// 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 →
@@ -4633,6 +4683,65 @@ pub struct GpuDqnTrainer {
/// adaptive max-floor on the Kelly cap. Loaded from
/// `kelly_warmup_floor_compute_kernel.cubin`.
sp9_kelly_warmup_floor_compute_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, A1.1): val-sharpe Δ + variance canary
/// producer. Single-block, single-thread kernel reads
/// `val_sharpe_history_pinned` (mapped-pinned 2-element [prev, curr])
/// + `ISV[VAL_SHARPE_DELTA_EMA_INDEX]`; writes raw delta and squared
/// deviation to `scratch[SCRATCH_SP11_VAL_SHARPE_DELTA_BASE..+2)`.
/// Followed by 1 apply_pearls_ad_kernel launch (n_slots=2) →
/// ISV[VAL_SHARPE_DELTA_EMA_INDEX=350, VAL_SHARPE_VAR_EMA_INDEX=351].
/// Loaded from `val_sharpe_delta_compute_kernel.cubin`.
sp11_val_sharpe_delta_compute_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, A1.2): per-bar saboteur engagement canary
/// producer. Single-block, 256-thread kernel with grid-stride iteration
/// over `saboteur_delta_reward_buf` (per-bar |Δreward| populated by
/// `experience_env_step` at the saboteur perturbation site). Block tree-
/// reduces engagement count, divides by n_bars; writes fraction to
/// `scratch[SCRATCH_SP11_SABOTEUR_ENGAGEMENT=268]`. Followed by 1
/// apply_pearls_ad_kernel launch → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
/// Loaded from `saboteur_engagement_compute_kernel.cubin`.
sp11_saboteur_engagement_compute_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, A1.3): per-component reward-magnitude-ratio
/// canary producer. Single-block, 6-thread kernel reads 6 contiguous
/// EMAs at `ISV[REWARD_POPART_EMA_INDEX..+6)`, normalises to ratios,
/// and mirrors the popart magnitude into scratch[6] as a side-output.
/// Writes 7 floats to `scratch[SCRATCH_SP11_MAG_RATIO_BASE..+7)`.
/// Followed by 2 apply_pearls_ad_kernel launches:
/// (n_slots=6) → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6)
/// (n_slots=1) → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
/// Loaded from `reward_component_mag_ratio_compute_kernel.cubin`.
sp11_mag_ratio_compute_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, A1.1): mapped-pinned 2-element val-sharpe
/// history buffer consumed by `sp11_val_sharpe_delta_compute_kernel`.
/// Layout: `[prev_epoch_val_sharpe, curr_epoch_val_sharpe]`. The host
/// writes the current epoch's val sharpe to slot[1] each val emit
/// boundary in `training_loop.rs` (rotating slot[1] → slot[0] before
/// the new write). Buffer is constructor-zero-initialised, which
/// Pearl A's sentinel-bootstrap path consumes naturally on the first
/// observation. The host write is a permitted CPU↔GPU path: it goes
/// through mapped-pinned memory (no `htod_copy`), and writes a literal
/// already-computed value rather than performing host-side compute,
/// so it does NOT violate `feedback_no_htod_htoh_only_mapped_pinned`
/// or `feedback_no_cpu_compute_strict`.
pub(crate) val_sharpe_history_pinned: super::mapped_pinned::MappedF32Buffer,
/// SP11 Fix 39 (2026-05-04, A1.2): per-bar saboteur Δreward GPU buffer
/// populated by `experience_env_step` at the saboteur perturbation site.
/// Length = `alloc_episodes × alloc_timesteps`; layout matches
/// `total_reward_per_sample` (row-major [N, L]). Consumed by
/// `saboteur_engagement_compute_kernel`. Pure GPU buffer (CudaSlice<f32>)
/// — never read host-side, so the no-CudaSlice rule does not apply
/// (it's a GPU-only path between two kernels). Owned by
/// `gpu_experience_collector`; the trainer holds the device pointer
/// (populated post-experience-collection) and the bar count.
/// Per `feedback_wire_everything_up`: this field is initialised in
/// the same commit that adds it; the consumer in `training_loop.rs`
/// reads `self.gpu_experience_collector.saboteur_delta_reward_dev_ptr()`
/// just before each per-step launch.
pub(crate) sp11_saboteur_delta_reward_dev_ptr: u64,
/// Length of `sp11_saboteur_delta_reward_dev_ptr`'s underlying buffer
/// (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,
/// 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
@@ -12274,6 +12383,284 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP11 Fix 39 (2026-05-04, A1.1): set the GPU device pointer + bar
/// count for the per-bar saboteur Δreward buffer populated by
/// `experience_env_step`. The trainer caches this so
/// `launch_sp11_saboteur_engagement_compute` has the consumer pointer
/// available without re-querying the collector each step. Per
/// `feedback_wire_everything_up`: the buffer is allocated by the
/// collector and wired into the trainer in the same commit that adds
/// it. Initial value is zero/zero (constructor); the training-loop
/// init path calls this setter once after collector construction.
pub(crate) fn set_sp11_saboteur_delta_reward_buf(
&mut self,
dev_ptr: u64,
n_bars: usize,
) {
self.sp11_saboteur_delta_reward_dev_ptr = dev_ptr;
self.sp11_saboteur_delta_reward_len = n_bars;
}
/// SP11 Fix 39 (2026-05-04, A1.1): launch the val-sharpe Δ + variance
/// canary producer.
///
/// Reads `val_sharpe_history_pinned` (mapped-pinned 2-element
/// `[prev_epoch, curr_epoch]` written from `training_loop.rs` at the
/// val emit boundary) plus `ISV[VAL_SHARPE_DELTA_EMA_INDEX]`; writes
/// raw delta + squared deviation to
/// `scratch[SCRATCH_SP11_VAL_SHARPE_DELTA_BASE..+2)`. Chained
/// `apply_pearls_ad_kernel` (n_slots=2) smooths into
/// ISV[VAL_SHARPE_DELTA_EMA_INDEX=350, VAL_SHARPE_VAR_EMA_INDEX=351].
///
/// Layer A is additive — no consumer reads either canary slot in this
/// commit. The controller producer (Layer A2) and reward-shaping
/// consumers (Layer B) wire downstream.
pub(crate) fn launch_sp11_val_sharpe_delta_compute(&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::{
VAL_SHARPE_DELTA_EMA_INDEX, VAL_SHARPE_VAR_EMA_INDEX,
};
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_sp11_val_sharpe_delta_compute: isv_signals_dev_ptr must be allocated");
debug_assert_eq!(self.val_sharpe_history_pinned.len, 2,
"launch_sp11_val_sharpe_delta_compute: val_sharpe_history_pinned must be 2 f32");
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;
let history_dev = self.val_sharpe_history_pinned.dev_ptr;
let scratch_base_i32: i32 = SCRATCH_SP11_VAL_SHARPE_DELTA_BASE as i32;
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;
// Step 1: producer kernel writes scratch[base..base+2).
// The kernel's scratch_out arg is offset by the base scratch index
// via pointer arithmetic — the kernel writes to `scratch_out[0..2)`
// relative to whatever pointer it receives. Pass the offset device
// pointer rather than a base index parameter (keeps the kernel
// signature minimal and matches the SP9 producer-launcher
// pattern at `launch_intent_eval_divergence_compute`).
let scratch_offset_bytes = (SCRATCH_SP11_VAL_SHARPE_DELTA_BASE as u64)
* std::mem::size_of::<f32>() as u64;
let scratch_out_dev: u64 = scratch_dev + scratch_offset_bytes;
unsafe {
self.stream
.launch_builder(&self.sp11_val_sharpe_delta_compute_kernel)
.arg(&history_dev)
.arg(&scratch_out_dev)
.arg(&isv_dev)
.arg(&delta_ema_slot_i32)
.arg(&var_ema_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("sp11 val_sharpe_delta_compute: {e}")))?;
}
// Step 2: chained apply_pearls_ad_kernel (n_slots=2) — both
// VAL_SHARPE_DELTA_EMA_INDEX and VAL_SHARPE_VAR_EMA_INDEX are
// contiguous (350, 351) so a single n_slots=2 launch covers them.
let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3;
let wiener_off = base_wiener_offset
+ (delta_ema_slot_i32 - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, scratch_base_i32,
isv_dev, delta_ema_slot_i32,
wiener_dev, wiener_off,
2,
ALPHA_META,
)?;
}
Ok(())
}
/// SP11 Fix 39 (2026-05-04, A1.2): launch the per-bar saboteur
/// engagement canary producer.
///
/// Reads `sp11_saboteur_delta_reward_dev_ptr` (a per-bar |Δreward|
/// GPU buffer populated by `experience_env_step` at the saboteur
/// perturbation site). Block tree-reduces engagement count over all
/// `sp11_saboteur_delta_reward_len` bars, writes fraction to
/// `scratch[SCRATCH_SP11_SABOTEUR_ENGAGEMENT]`. Chained
/// `apply_pearls_ad_kernel` (n_slots=1) smooths into
/// ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
///
/// Pre-condition: `set_sp11_saboteur_delta_reward_buf` has been called
/// with the collector's saboteur Δreward dev ptr. If the saboteur is
/// inactive on a given epoch (delta_reward stays zero), engagement
/// reads zero and Pearl A's sentinel-bootstrap path holds the EMA
/// at zero — no spurious activation.
pub(crate) fn launch_sp11_saboteur_engagement_compute(&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::{
SABOTEUR_ENGAGEMENT_RATE_INDEX, PNL_REWARD_MAGNITUDE_EMA_INDEX,
};
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_sp11_saboteur_engagement_compute: isv_signals_dev_ptr must be allocated");
// No-op when the producer hasn't been wired yet (e.g. early init
// before the collector exists). This is NOT a stub return value
// — it's a precondition guard that defers the launch until the
// wire-up is complete. Once `set_sp11_saboteur_delta_reward_buf`
// is called with a valid pointer, every subsequent step launches
// the kernel for real.
if self.sp11_saboteur_delta_reward_dev_ptr == 0
|| self.sp11_saboteur_delta_reward_len == 0
{
return Ok(());
}
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;
let delta_dev = self.sp11_saboteur_delta_reward_dev_ptr;
let scratch_offset_bytes = (SCRATCH_SP11_SABOTEUR_ENGAGEMENT as u64)
* std::mem::size_of::<f32>() as u64;
let scratch_out_dev: u64 = scratch_dev + scratch_offset_bytes;
let scratch_idx_i32: i32 = SCRATCH_SP11_SABOTEUR_ENGAGEMENT as i32;
let n_bars_i32: i32 = i32::try_from(self.sp11_saboteur_delta_reward_len)
.unwrap_or(i32::MAX);
let pnl_mag_slot_i32: i32 = PNL_REWARD_MAGNITUDE_EMA_INDEX as i32;
// Step 1: producer kernel — block tree-reduce, BLOCK_SIZE=256.
// Shared memory = 256 ints = 1024 bytes.
unsafe {
self.stream
.launch_builder(&self.sp11_saboteur_engagement_compute_kernel)
.arg(&delta_dev)
.arg(&scratch_out_dev)
.arg(&isv_dev)
.arg(&n_bars_i32)
.arg(&pnl_mag_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: 256 * std::mem::size_of::<i32>() as u32,
})
.map_err(|e| MLError::ModelError(format!("sp11 saboteur_engagement_compute: {e}")))?;
}
// Step 2: chained apply_pearls_ad_kernel singleton →
// ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3;
let isv_idx_i32 = SABOTEUR_ENGAGEMENT_RATE_INDEX as i32;
let wiener_off = base_wiener_offset + (isv_idx_i32 - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, scratch_idx_i32,
isv_dev, isv_idx_i32,
wiener_dev, wiener_off,
1,
ALPHA_META,
)?;
}
Ok(())
}
/// SP11 Fix 39 (2026-05-04, A1.3): launch the per-component reward-
/// magnitude-ratio canary producer.
///
/// Reads 6 contiguous EMAs at `ISV[REWARD_POPART_EMA_INDEX..+6)`,
/// normalises to ratios in `scratch[SCRATCH_SP11_MAG_RATIO_BASE..+6)`,
/// and mirrors the popart magnitude into `scratch[base+6]`. Followed
/// by two chained `apply_pearls_ad_kernel` launches:
/// 1. n_slots=6 → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6)
/// 2. n_slots=1 → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
/// (The ISV slots are non-contiguous: 352..358 plus 359, so two
/// launches rather than one n_slots=7. The kernel still emits 7
/// scratch values in one shot.)
pub(crate) fn launch_sp11_mag_ratio_compute(&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_COMPONENT_MAG_RATIO_BASE, REWARD_COMPONENT_MAG_RATIO_COUNT,
PNL_REWARD_MAGNITUDE_EMA_INDEX,
};
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_sp11_mag_ratio_compute: 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;
let scratch_offset_bytes = (SCRATCH_SP11_MAG_RATIO_BASE as u64)
* std::mem::size_of::<f32>() as u64;
let scratch_out_dev: u64 = scratch_dev + scratch_offset_bytes;
let popart_ema_base_slot_i32: i32 = REWARD_POPART_EMA_INDEX as i32;
unsafe {
self.stream
.launch_builder(&self.sp11_mag_ratio_compute_kernel)
.arg(&isv_dev)
.arg(&scratch_out_dev)
.arg(&popart_ema_base_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (6, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!("sp11 mag_ratio_compute: {e}")))?;
}
let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3;
// Pearls A+D #1: 6 contiguous slots → ISV[352..358).
let ratios_isv_base_i32 = REWARD_COMPONENT_MAG_RATIO_BASE as i32;
let ratios_scratch_idx_i32 = SCRATCH_SP11_MAG_RATIO_BASE as i32;
let ratios_wiener_off = base_wiener_offset
+ (ratios_isv_base_i32 - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, ratios_scratch_idx_i32,
isv_dev, ratios_isv_base_i32,
wiener_dev, ratios_wiener_off,
REWARD_COMPONENT_MAG_RATIO_COUNT as i32,
ALPHA_META,
)?;
}
// Pearls A+D #2: singleton mirror → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359].
let mirror_isv_idx_i32 = PNL_REWARD_MAGNITUDE_EMA_INDEX as i32;
let mirror_scratch_idx_i32 = (SCRATCH_SP11_MAG_RATIO_BASE
+ REWARD_COMPONENT_MAG_RATIO_COUNT) as i32;
let mirror_wiener_off = base_wiener_offset
+ (mirror_isv_idx_i32 - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, mirror_scratch_idx_i32,
isv_dev, mirror_isv_idx_i32,
wiener_dev, mirror_wiener_off,
1,
ALPHA_META,
)?;
}
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.
@@ -15572,6 +15959,46 @@ impl GpuDqnTrainer {
.map_err(|e| MLError::ModelError(format!("sp9 kelly_warmup_floor_compute load: {e}")))?
};
// SP11 Fix 39 (2026-05-04, Task A1): load three canary producer
// kernels for the reward-as-controlled-subsystem chain. Each is a
// single-block producer chained through `apply_pearls_ad_kernel`
// for Pearl A sentinel-bootstrap + Pearl D Wiener-α smoothing.
// No consumer reads any of the canary slots [350..360) yet
// (Layer A is additive); Layer B atomically migrates consumers.
let sp11_val_sharpe_delta_compute_kernel = {
let module = stream.context()
.load_cubin(SP11_VAL_SHARPE_DELTA_COMPUTE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp11 val_sharpe_delta_compute cubin load: {e}")))?;
module.load_function("val_sharpe_delta_compute_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 val_sharpe_delta_compute load: {e}")))?
};
let sp11_saboteur_engagement_compute_kernel = {
let module = stream.context()
.load_cubin(SP11_SABOTEUR_ENGAGEMENT_COMPUTE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp11 saboteur_engagement_compute cubin load: {e}")))?;
module.load_function("saboteur_engagement_compute_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 saboteur_engagement_compute load: {e}")))?
};
let sp11_mag_ratio_compute_kernel = {
let module = stream.context()
.load_cubin(SP11_MAG_RATIO_COMPUTE_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp11 mag_ratio_compute cubin load: {e}")))?;
module.load_function("reward_component_mag_ratio_compute_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 mag_ratio_compute load: {e}")))?
};
// SP11 Fix 39 (2026-05-04, A1.1): allocate the val-sharpe history
// mapped-pinned buffer. 2-element [prev_epoch, curr_epoch] — the
// host writes the current val sharpe at every val emit boundary
// in `training_loop.rs`, then rotates slot[1] into slot[0] before
// the next observation. Constructor-zero-initialised, which Pearl
// A's first-observation replacement consumes naturally.
let val_sharpe_history_pinned = unsafe {
super::mapped_pinned::MappedF32Buffer::new(2)
}.map_err(|e| MLError::ModelError(
format!("SP11 val_sharpe_history_pinned alloc: {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
@@ -17200,7 +17627,13 @@ impl GpuDqnTrainer {
// [242..250) loss_balance_active (SP7 activation-flag fix — CQL/C51 active flags)
// [250..251) train_active_frac (SP8 Fix 36 — canary scratch slot)
// [251..259) lb_max_budget_* (SP8 Fix 36 — per-(head, branch) cap)
// Total: SP5_SCRATCH_TOTAL = 259. Audit doc records every commit's growth.
// [259..262) sp9 kelly + q_var_mag + intent_eval (SP9 Fix 37)
// [262..265) sp9 eval_dist Q/H/F (SP9 Fix 37)
// [265..266) sp10 thompson temp (SP10 Fix 38)
// [266..268) sp11 val_sharpe_delta + var (SP11 Fix 39 A1.1)
// [268..269) sp11 saboteur_engagement (SP11 Fix 39 A1.2)
// [269..276) sp11 mag_ratio (6 ratios + popart mirror) (SP11 Fix 39 A1.3)
// Total: SP5_SCRATCH_TOTAL = 276. Audit doc records every commit's growth.
let producer_step_scratch_buf = unsafe { MappedF32Buffer::new(SP5_SCRATCH_TOTAL) }
.map_err(|e| MLError::ModelError(format!("SP4/SP5 producer_step_scratch_buf alloc: {e}")))?;
@@ -18833,6 +19266,17 @@ impl GpuDqnTrainer {
sp9_intent_eval_divergence_compute_kernel,
sp9_q_var_mag_ema_compute_kernel,
sp9_kelly_warmup_floor_compute_kernel,
sp11_val_sharpe_delta_compute_kernel,
sp11_saboteur_engagement_compute_kernel,
sp11_mag_ratio_compute_kernel,
val_sharpe_history_pinned,
// SP11 Fix 39: saboteur Δreward GPU buffer is owned by
// `gpu_experience_collector`. The trainer caches the device
// pointer + length; both default to 0 here and are wired in
// via `set_sp11_saboteur_delta_reward_buf` after the collector
// is constructed (see training_loop.rs init path).
sp11_saboteur_delta_reward_dev_ptr: 0,
sp11_saboteur_delta_reward_len: 0,
lb_budget_dispatch_kernel,
saxpy_f32_dev_ptr_aux,
scale_f32_dev_ptr_aux,

40
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -612,6 +612,14 @@ pub struct GpuExperienceCollector {
// Task 2.4: loss_aversion_per_sample REMOVED — R6 clamp relocated to
// c51_loss_kernel.cu::block_bellman_project_f (Q-target smoothing).
total_reward_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
/// SP11 Fix 39 (2026-05-04, A1.2): per-bar saboteur engagement signal.
/// [alloc_episodes * alloc_timesteps] — written by `experience_env_step`
/// at the END (after `total_reward_per_sample` is finalised) as
/// `traded × |reward| × max(|eff_spread − 1|, |eff_slip − 1|)`. Pure
/// GPU buffer (kernel-to-kernel only); the SP11 saboteur engagement
/// kernel reads via raw device pointer. See spec §3.3.1 + the param
/// docstring on `experience_env_step::saboteur_delta_reward_per_sample`.
saboteur_delta_reward_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
/// C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2).
/// [alloc_episodes * alloc_timesteps * 6] row-major.
/// Components: [0]=popart(total), [1]=cf, [2]=trail(0), [3]=micro, [4]=opp_cost, [5]=bonus.
@@ -1192,6 +1200,14 @@ impl GpuExperienceCollector {
let total_reward_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc total_reward_per_sample: {e}")))?;
// SP11 Fix 39 (2026-05-04, A1.2): per-bar saboteur engagement signal
// buffer. Same shape as `total_reward_per_sample` (one f32 per (i,t)).
// Written by `experience_env_step` after the trade lifecycle resolves
// (line ~3220, post-`total_reward_per_sample[out_off] = reward`).
// Read by the SP11 saboteur engagement kernel via raw device pointer.
let saboteur_delta_reward_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc saboteur_delta_reward_per_sample: {e}")))?;
// C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2).
// [total_output * 6] row-major: [popart, cf, trail, micro, opp_cost, bonus] per (i,t).
// Producer: experience_env_step (writes rc[0], rc[1], rc[3], rc[4], rc[5]; rc[2]=0 placeholder).
@@ -1690,6 +1706,7 @@ impl GpuExperienceCollector {
cf_flip_per_sample,
micro_reward_per_sample,
total_reward_per_sample,
saboteur_delta_reward_per_sample,
reward_components_per_sample,
flat_to_pos_per_sample,
readiness_per_sample,
@@ -2104,6 +2121,21 @@ impl GpuExperienceCollector {
self.saboteur_perturbation_scale = scale;
}
/// SP11 Fix 39 (2026-05-04, A1.2): expose the per-bar saboteur engagement
/// signal buffer to the trainer. Returns `(dev_ptr, n_bars)` where
/// `n_bars = alloc_episodes × alloc_timesteps`. The trainer caches this
/// via `GpuDqnTrainer::set_sp11_saboteur_delta_reward_buf` after the
/// collector is constructed (training_loop.rs init path) so
/// `launch_sp11_saboteur_engagement_compute` has the consumer-side
/// pointer ready every step. Per `feedback_wire_everything_up`: this
/// getter + the trainer setter + the env_step kernel write site land
/// together in this commit.
pub fn sp11_saboteur_delta_reward_dev_ptr(&self) -> (u64, usize) {
let dev_ptr = self.saboteur_delta_reward_per_sample.device_ptr(&self.stream).0;
let n_bars = self.alloc_episodes * self.alloc_timesteps;
(dev_ptr, n_bars)
}
/// Upload pre-computed expert demonstration actions to GPU.
///
/// `expert_actions[bar_index]` = expert exposure action index (-1 = no opinion).
@@ -3973,6 +4005,14 @@ impl GpuExperienceCollector {
use crate::cuda_pipeline::gpu_dqn_trainer::Q_DIR_ABS_REF_INDEX;
Q_DIR_ABS_REF_INDEX as i32
})
// SP11 Fix 39 (2026-05-04, A1.2): per-bar saboteur
// engagement signal output. The kernel writes
// `traded × |reward| × max(|eff_spread − 1|, |eff_slip − 1|)`
// here at the END of the bar processing. The SP11
// saboteur engagement kernel consumes this buffer via
// raw device pointer (see `set_sp11_saboteur_delta_reward_buf`
// wire-up in training_loop.rs init).
.arg(&mut self.saboteur_delta_reward_per_sample)
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_env_step t={t}: {e}"

86
crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu Normal file
View File

@@ -0,0 +1,86 @@
// crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
//
// SP11 (Fix 39, 2026-05-04, A1.3): per-component reward-magnitude-ratio
// canary producer.
//
// Reads 6 existing per-component reward-magnitude EMAs from ISV at the
// contiguous block `[popart_ema_base_slot..+6)` (= REWARD_POPART_EMA_INDEX
// = 63, populated by the SP4 reward-component EMA kernel). Components in
// fixed order (per `gpu_dqn_trainer.rs::REWARD_COMPONENT_COUNT`):
// 0 popart 1 cf 2 trail 3 micro 4 opp_cost 5 bonus
//
// Writes 7 floats to `scratch_out`:
// scratch_out[0..6) = ratios (each / sum), Σ=1 modulo EPS_DIV floor
// scratch_out[6] = isv[popart_ema_base_slot] — popart magnitude
// mirrored to PNL_REWARD_MAGNITUDE_EMA_INDEX = 359
// via the chained Pearls-A+D launch (n_slots=7).
//
// The mirror at scratch[6] avoids allocating a separate copy kernel:
// the popart-component EMA IS the PnL signal scale used to bound
// curiosity (spec §3.4.2) and to threshold saboteur engagement
// (spec §3.3.1). Promoting it to its own ISV slot keeps the controller's
// signal-relative bounds first-class.
//
// Source EMAs are already populated by SP4's `reward_component_ema_kernel`
// — this kernel only reads, normalises, and forwards (no compute beyond
// sum + divide). Single block, 6 threads (one per component); thread 0
// also computes the sum and writes the popart mirror.
//
// Pearls A+D applied chain-style by `apply_pearls_ad_kernel` (n_slots=7)
// targeting ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6) ∪ {PNL_REWARD_
// MAGNITUDE_EMA_INDEX}. The ISV slots are NOT contiguous (six at 352..358
// then a gap to 359), so the launcher chains two `launch_apply_pearls`
// calls — a 6-slot block then the singleton — rather than one n_slots=7
// call. The kernel signature is unchanged.
#include <cuda_runtime.h>
extern "C" __global__ void reward_component_mag_ratio_compute_kernel(
/* ISV signal bus (read-only). */
const float* __restrict__ isv,
/* Producer scratch buffer. Caller indexes the base into this array
* via the `REWARD_COMPONENT_MAG_RATIO_BASE` constant. The kernel
* writes the 7 values relative to the base provided by the launcher
* (the launcher passes a sub-slice via pointer arithmetic at the
* `arg(&scratch_dev_with_offset)` site). */
float* __restrict__ scratch_out,
/* ISV slot index of the first per-component magnitude EMA.
* = REWARD_POPART_EMA_INDEX = 63. */
int popart_ema_base_slot
) {
if (blockIdx.x != 0) return;
if (threadIdx.x >= 6) return;
/* Block-load all 6 component magnitudes into shared mem. */
__shared__ float values[6];
__shared__ float sum;
values[threadIdx.x] = isv[popart_ema_base_slot + (int)threadIdx.x];
__syncthreads();
if (threadIdx.x == 0) {
float s = 0.0f;
#pragma unroll
for (int c = 0; c < 6; c++) s += values[c];
/* EPS_DIV (1e-6) floor — Invariant-1 numerical-stability anchor
* matching SP4/SP5/SP9/SP10 conventions. Guards the impossible
* "all-zero magnitudes" case during the first observations after
* a fold reset (Pearl A bootstraps each component EMA from its
* first observation, but timing of Pearls A+D fires across the 6
* sources is per-slot, so a literal 0/0 is conceivable in the
* cold-start window). */
sum = fmaxf(s, 1e-6f);
/* Side-output mirror: popart magnitude → ISV[359] via chained
* Pearls A+D singleton launch. Read directly from values[0]
* which thread 0 just loaded above. */
scratch_out[6] = values[0];
}
__syncthreads();
/* Each of threads 0..5 writes its own ratio. */
scratch_out[threadIdx.x] = values[threadIdx.x] / sum;
if (threadIdx.x == 0) {
__threadfence_system();
}
}

88
crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu Normal file
View File

@@ -0,0 +1,88 @@
// crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu
//
// SP11 (Fix 39, 2026-05-04, A1.2): per-bar saboteur engagement canary
// producer.
//
// Reads `saboteur_delta_reward_buf` populated by `experience_env_step` at
// the saboteur perturbation site (per spec §3.3.1, single reward
// computation: the env_step kernel writes the cost-delta the saboteur
// imposed on each bar where a trade occurred). For each bar:
//
// engaged_i = (|saboteur_delta_reward[i]| > EPS_ENGAGEMENT) ? 1 : 0
// EPS_ENGAGEMENT = 0.01 × max(isv[PNL_REWARD_MAGNITUDE_EMA_INDEX], 1e-6)
//
// Reduction: per-block tree-reduce in shared memory (no atomicAdd per
// `feedback_no_atomicadd.md`). The bar count `n_bars` is the per-epoch
// `alloc_episodes × alloc_timesteps` populated by env_step. With a
// single-block launch and 256 threads + grid-stride iteration, the
// kernel handles arbitrary batch sizes.
//
// Output: scratch_out[0] = (engaged_count / n_bars) ∈ [0, 1].
// Downstream `apply_pearls_ad_kernel` (n_slots=1) smooths into
// ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358] via Pearls A+D.
//
// Pearl A: sentinel 0 → first observation replaces directly.
// Pearl D: Wiener-α smoothing chained on stream.
//
// Per spec §3.3.1: the EPS threshold is signal-relative — reading
// PNL_REWARD_MAGNITUDE_EMA_INDEX (slot 359, populated by the SP11 mag-
// ratio kernel as a side-output mirroring REWARD_POPART_EMA_INDEX) keeps
// engagement classification scale-invariant across PnL regimes.
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
extern "C" __global__ void saboteur_engagement_compute_kernel(
/* Per-bar |reward_with_saboteur - reward_without_saboteur|.
* Length = `n_bars` = alloc_episodes × alloc_timesteps. */
const float* __restrict__ saboteur_delta_reward,
/* Scratch buffer; thread 0 writes scratch_out[0] = engagement fraction. */
float* __restrict__ scratch_out,
/* ISV signal bus (read-only). */
const float* __restrict__ isv,
/* Number of bars to scan. */
int n_bars,
/* ISV slot index for the PnL-magnitude EMA (signal-relative threshold). */
int pnl_mag_ema_slot // = PNL_REWARD_MAGNITUDE_EMA_INDEX = 359
) {
__shared__ int sdata[BLOCK_SIZE];
/* Read PnL-magnitude EMA once per block (broadcast via shared) — the
* threshold is uniform across threads. The 1e-6f floor is an
* Invariant-1 numerical-stability anchor (matches SP9/SP10 EPS_DIV)
* and prevents threshold collapse to 0 before Pearl A's first
* observation populates the slot. */
const float pnl_mag = fmaxf(isv[pnl_mag_ema_slot], 1e-6f);
const float eps_engagement = 0.01f * pnl_mag;
/* Per-thread grid-stride accumulator. */
int local_engaged = 0;
if (n_bars > 0) {
for (int i = (int)threadIdx.x; i < n_bars; i += BLOCK_SIZE) {
float dr = saboteur_delta_reward[i];
if (fabsf(dr) > eps_engagement) {
local_engaged += 1;
}
}
}
sdata[threadIdx.x] = local_engaged;
__syncthreads();
/* Block tree-reduce — power-of-two BLOCK_SIZE. */
for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
if ((int)threadIdx.x < s) {
sdata[threadIdx.x] += sdata[threadIdx.x + s];
}
__syncthreads();
}
if (threadIdx.x == 0) {
const int total = sdata[0];
const float fraction = (n_bars > 0)
? (float)total / (float)n_bars
: 0.0f;
scratch_out[0] = fraction;
__threadfence_system();
}
}

72
crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu Normal file
View File

@@ -0,0 +1,72 @@
// crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu
//
// SP11 (Fix 39, 2026-05-04, A1.1): val-sharpe delta + variance canary producer.
//
// Reads a 2-element `val_sharpe_history` mapped-pinned buffer (the host
// writes [prev_epoch, curr_epoch] just before launch) and the prior
// VAL_SHARPE_DELTA_EMA_INDEX from ISV (set by Pearls A+D on the prior
// observation). Writes:
// scratch_out[0] = (curr - prev) — raw delta
// scratch_out[1] = (delta - prev_delta_ema)^2 — squared deviation
//
// Downstream `apply_pearls_ad_kernel` (n_slots=2) chained on the same
// stream smooths these into:
// ISV[VAL_SHARPE_DELTA_EMA_INDEX=350] (Δ EMA)
// ISV[VAL_SHARPE_VAR_EMA_INDEX=351] (variance EMA, true variance)
//
// Pearl A: sentinel 0 → first observation replaces directly. The squared-
// deviation against an as-yet-unwritten ISV slot is `delta * delta` on
// the first call (since `prev_delta_ema = 0` from the FoldReset sentinel),
// which Pearl A then bootstraps into `VAL_SHARPE_VAR_EMA_INDEX`.
//
// Pearl D: Wiener-α smoothing applied per slot inside apply_pearls_ad_kernel
// chained on the same stream — α adapts to each slot's own noise level.
//
// Single block, single thread; cheap arithmetic. No atomicAdd
// (`feedback_no_atomicadd.md`); `__threadfence_system()` after the writes
// to make the device-resident scratch visible to the chained Pearls launch.
#include <cuda_runtime.h>
extern "C" __global__ void val_sharpe_delta_compute_kernel(
/* val_sharpe_history: 2-element mapped-pinned buffer.
* [0] = previous-epoch val sharpe (-Inf at first observation; the
* host writes the current value into [1] then rotates).
* [1] = current-epoch val sharpe.
* Read-only.
*/
const float* __restrict__ val_sharpe_history,
/* Producer scratch buffer (mapped pinned, indexed by scratch_idx_base). */
float* __restrict__ scratch_out,
/* ISV signal bus (read-only). */
const float* __restrict__ isv,
/* ISV slot indices for Pearls A+D consumption. */
int delta_ema_slot, // = VAL_SHARPE_DELTA_EMA_INDEX = 350
int var_ema_slot // = VAL_SHARPE_VAR_EMA_INDEX = 351
) {
if (blockIdx.x != 0 || threadIdx.x != 0) return;
/* val_sharpe_history layout (host writes from training_loop.rs val
* boundary): [t-1, t]. The first observation has prev = 0.0 (the
* mapped-pinned buffer is constructor-zero-initialised), so the
* first emitted delta carries the absolute current val sharpe,
* which Pearl A then sentinel-bootstraps into the EMA. */
const float prev = val_sharpe_history[0];
const float curr = val_sharpe_history[1];
const float delta = curr - prev;
/* Squared deviation against the PREVIOUS delta_ema. After Pearl A's
* first-observation replacement this becomes the per-call deviation
* from the running mean, suitable as a true-variance-EMA input. The
* (void)var_ema_slot read keeps the kernel signature self-contained:
* the slot index is consumed by the chained apply_pearls_ad launcher,
* not by this kernel directly. */
const float prev_delta_ema = isv[delta_ema_slot];
const float dev = delta - prev_delta_ema;
scratch_out[0] = delta; // → ISV[delta_ema_slot] via Pearls A+D
scratch_out[1] = dev * dev; // → ISV[var_ema_slot] via Pearls A+D
(void)var_ema_slot; // see comment above
__threadfence_system();
}

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

@@ -802,6 +802,23 @@ impl DQNTrainer {
}
}
// SP11 Fix 39 (2026-05-04, A1.2): wire the per-bar saboteur
// engagement signal buffer from the experience collector into
// the trainer. The collector's `experience_env_step` writes
// the per-bar signal at the END of each bar; the trainer's
// `launch_sp11_saboteur_engagement_compute` reads via raw
// device pointer and reduces to a per-step engagement
// fraction → Pearls A+D → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
// Per `feedback_wire_everything_up`: producer (collector)
// and consumer (trainer) both wire in the same commit.
if let Some(ref collector) = self.gpu_experience_collector {
let (sab_dev_ptr, n_bars) = collector.sp11_saboteur_delta_reward_dev_ptr();
if let Some(ref mut fused_mut) = self.fused_ctx {
fused_mut.trainer_mut()
.set_sp11_saboteur_delta_reward_buf(sab_dev_ptr, n_bars);
}
}
// IQR exploration bonus from IQN quantile spread
if let Some(ref fused) = self.fused_ctx {
let iqr_ptr = fused.iqn_iqr_ptr().unwrap_or(0);
@@ -3387,6 +3404,49 @@ impl DQNTrainer {
tracing::warn!(error = %e,
"SP9 kelly_warmup_floor producer chain launch failed");
}
// SP11 Fix 39 (2026-05-04, A1.3): per-component
// reward-magnitude-ratio canary producer.
// Reads ISV[REWARD_POPART_EMA_INDEX..+6) (SP4
// reward-component EMAs populated by the
// experience collector's reward_component_ema
// call earlier in the epoch) and writes
// 6 normalised ratios + 1 popart mirror.
// Pearls A+D smooth into
// ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6) and
// ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359].
// Per-epoch cadence — matches the SP4 reward-
// component EMA producer firing pattern.
if let Err(e) = fused.trainer()
.launch_sp11_mag_ratio_compute()
{
tracing::warn!(error = %e,
"SP11 mag_ratio_compute launch failed");
}
// SP11 Fix 39 (2026-05-04, A1.2): per-bar
// saboteur engagement canary producer. Reads
// the per-bar saboteur Δreward signal written
// by `experience_env_step` (wired into the
// trainer at init via
// `set_sp11_saboteur_delta_reward_buf`). When
// the saboteur is inactive this epoch the
// signal stays at zero across all bars, the
// engagement fraction reads zero, and Pearl A
// holds the EMA at sentinel. Pearls A+D smooth
// into ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
//
// Pre-condition: launch_sp11_mag_ratio_compute
// ran above — so PNL_REWARD_MAGNITUDE_EMA_INDEX
// has its first observation populated by the
// time the engagement kernel reads it as the
// signal-relative threshold base.
if let Err(e) = fused.trainer()
.launch_sp11_saboteur_engagement_compute()
{
tracing::warn!(error = %e,
"SP11 saboteur_engagement_compute launch failed");
}
}
}
}
@@ -4865,6 +4925,37 @@ impl DQNTrainer {
"Epoch {}/{}: val_Sharpe={:.2} (deterministic backtest on fixed validation window)",
epoch + 1, self.hyperparams.epochs, val_sharpe,
);
// SP11 Fix 39 (2026-05-04, A1.1): launch the val-sharpe Δ + variance
// canary producer chain. Rotates `val_sharpe_history_pinned[1]`
// (current) into `[0]` (previous), writes the freshly-computed
// `val_sharpe` to `[1]`, then launches the producer kernel which
// computes `(curr - prev)` + squared-deviation against the prior
// `ISV[VAL_SHARPE_DELTA_EMA_INDEX]` and chains Pearls A+D into
// ISV[VAL_SHARPE_DELTA_EMA_INDEX=350, VAL_SHARPE_VAR_EMA_INDEX=351].
//
// The host-side write through `host_slice_mut()` is a permitted
// CPU↔GPU path: it writes to mapped-pinned memory directly (no
// `htod_copy`), and writes a literal already-computed value (no
// host-side compute), so it does NOT violate
// `feedback_no_htod_htoh_only_mapped_pinned` or
// `feedback_no_cpu_compute_strict`.
//
// Layer A is additive — no consumer reads either canary slot in
// this commit. The downstream controller wires in Layer A2.
if let Some(ref mut fused) = self.fused_ctx {
let trainer_mut = fused.trainer_mut();
{
let history = trainer_mut.val_sharpe_history_pinned.host_slice_mut();
// Rotate: prev ← curr, then curr ← new observation.
history[0] = history[1];
history[1] = val_sharpe as f32;
}
if let Err(e) = trainer_mut.launch_sp11_val_sharpe_delta_compute() {
tracing::warn!(error = %e,
"SP11 val_sharpe_delta_compute launch failed");
}
}
// Source-of-truth alignment: read the EXACT annualization factor the
// kernel applied (`sqrt(bars_per_day × trading_days_per_year)`) rather
// than recomputing from `self.hyperparams.bars_per_day` here, which

282
crates/ml/tests/sp11_producer_unit_tests.rs Normal file
View File

@@ -0,0 +1,282 @@
#![allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
//! SP11 (Fix 39) Task A1 producer kernel unit tests.
//!
//! Validates the three canary producer kernels added by Task A1:
//! 1. `val_sharpe_delta_compute_kernel.cu`
//! 2. `saboteur_engagement_compute_kernel.cu`
//! 3. `reward_component_mag_ratio_compute_kernel.cu`
//!
//! Each test loads the producer cubin directly (skipping the chained
//! Pearls A+D launch — the launchers in `gpu_dqn_trainer.rs` chain that
//! step in production) and verifies the producer's scratch outputs
//! against known synthetic inputs. The chained Pearls A+D applicator is
//! exercised separately in `sp4_producer_unit_tests.rs`.
//!
//! All tests are `#[ignore = "requires GPU"]`-gated. Run on a GPU host:
//!
//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//! cargo test -p ml --test sp11_producer_unit_tests --features cuda \
//! -- --ignored --nocapture
//!
//! Per `feedback_no_htod_htoh_only_mapped_pinned`: every CPU↔GPU buffer
//! is a `MappedF32Buffer` allocated via `cuMemHostAlloc(DEVICEMAP|PORTABLE)`.
//! Zero `htod_copy`, zero `dtoh_sync_copy`, zero `alloc_zeros`. Tests are
//! NOT exempt from this rule. The host writes inputs through `host_slice_mut()`
//! / `write_from_slice()` (mapped-pinned coherence — kernel sees the
//! values after a stream sync), launches the kernel via `dev_ptr`, then
//! reads outputs via `read_all()` (volatile reads that defeat CPU caching).
#![cfg(feature = "cuda")]
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, REWARD_COMPONENT_MAG_RATIO_BASE,
VAL_SHARPE_DELTA_EMA_INDEX, VAL_SHARPE_VAR_EMA_INDEX,
};
/// ISV total dimension (must match `gpu_dqn_trainer::ISV_TOTAL_DIM`).
/// Tests allocate a full-size mapped-pinned ISV buffer and write the
/// few slots they need; the kernels read by ISV slot index, so a
/// shorter buffer would be unsafe even with synthetic-only writes.
const ISV_TOTAL_DIM: usize = 360;
/// REWARD_POPART_EMA_INDEX from `gpu_dqn_trainer.rs` (the SP4 reward-
/// component magnitude EMA base slot, mirrored here as a const so the
/// test imports stay scoped to the SP11 producer surface).
const REWARD_POPART_EMA_INDEX: usize = 63;
/// Cubin paths for each producer kernel under test. Each cubin is built
/// by `crates/ml/build.rs` from the `.cu` source and dropped into
/// `OUT_DIR` per the standard kernel-registration pattern.
const SP11_VAL_SHARPE_DELTA_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/val_sharpe_delta_compute_kernel.cubin"));
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"));
/// Build a CUDA stream against device 0. Mirrors `make_test_stream` in
/// `sp4_producer_unit_tests.rs` and `distributional_q_tests.rs`.
fn make_test_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
/// Load a kernel function from an embedded cubin by name.
fn load_kernel(stream: &Arc<CudaStream>, cubin: &[u8], func_name: &str) -> CudaFunction {
let module = stream
.context()
.load_cubin(cubin.to_vec())
.unwrap_or_else(|e| panic!("load cubin for {func_name}: {e}"));
module
.load_function(func_name)
.unwrap_or_else(|e| panic!("load function {func_name}: {e}"))
}
/// SP11 A1.1 — val-sharpe Δ + variance producer first observation.
///
/// Inputs: history = [10.0, 15.0] (Δ = 5.0); ISV[VAL_SHARPE_DELTA_EMA_INDEX]
/// = 0 (sentinel). Expected outputs: scratch[0] = 5.0 (raw Δ), scratch[1]
/// = (5.0 − 0)² = 25.0 (squared deviation). The chained Pearls A+D
/// applicator (not exercised here) bootstraps both EMAs from these
/// values via Pearl A's first-observation replacement.
#[test]
#[ignore = "requires GPU"]
fn val_sharpe_delta_first_observation_writes_raw_delta() {
let stream = make_test_stream();
let f = load_kernel(&stream, SP11_VAL_SHARPE_DELTA_CUBIN, "val_sharpe_delta_compute_kernel");
// Safety: CUDA context active on this thread (resolved through the
// stream's context above).
let mut history = unsafe { MappedF32Buffer::new(2) }
.expect("alloc val_sharpe_history (2 f32)");
{
let h = history.host_slice_mut();
h[0] = 10.0; // prev epoch
h[1] = 15.0; // curr epoch — Δ = 5.0
}
let scratch = unsafe { MappedF32Buffer::new(2) }
.expect("alloc scratch (2 f32)");
let isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
// ISV[VAL_SHARPE_DELTA_EMA_INDEX] is constructor-zero — sentinel state
// for Pearl A's first-observation replacement. No host writes needed.
let history_dev = history.dev_ptr;
let scratch_dev = scratch.dev_ptr;
let isv_dev = isv.dev_ptr;
let delta_slot_i32: i32 = VAL_SHARPE_DELTA_EMA_INDEX as i32;
let var_slot_i32: i32 = VAL_SHARPE_VAR_EMA_INDEX as i32;
unsafe {
stream
.launch_builder(&f)
.arg(&history_dev)
.arg(&scratch_dev)
.arg(&isv_dev)
.arg(&delta_slot_i32)
.arg(&var_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch val_sharpe_delta_compute_kernel");
}
stream.synchronize().expect("sync after val_sharpe_delta launch");
let out = scratch.read_all();
let raw_delta = out[0];
let sq_dev = out[1];
assert!(
(raw_delta - 5.0_f32).abs() < 1e-6,
"expected raw delta = 5.0, got {raw_delta}"
);
// Squared deviation against zero ema = 5.0² = 25.0
assert!(
(sq_dev - 25.0_f32).abs() < 1e-4,
"expected squared deviation = 25.0 (against zero ema), got {sq_dev}"
);
}
/// SP11 A1.2 — saboteur engagement classifier counts only bars whose
/// |Δreward| exceeds the signal-relative threshold.
///
/// Inputs: 4-bar synthetic Δreward = [0.0, 0.005, 0.05, 0.5]; ISV at
/// PNL_REWARD_MAGNITUDE_EMA_INDEX = 1.0. Threshold = 0.01 × 1.0 = 0.01.
/// Bars 0 and 1 fall below the threshold; bars 2 and 3 exceed it.
/// Expected engagement = 2/4 = 0.5.
#[test]
#[ignore = "requires GPU"]
fn saboteur_engagement_counts_above_threshold_only() {
let stream = make_test_stream();
let f = load_kernel(&stream, SP11_SABOTEUR_ENGAGEMENT_CUBIN, "saboteur_engagement_compute_kernel");
let delta_reward = unsafe { MappedF32Buffer::new(4) }
.expect("alloc delta_reward (4 f32)");
delta_reward.write_from_slice(&[0.0_f32, 0.005, 0.05, 0.5]);
let scratch = unsafe { MappedF32Buffer::new(1) }
.expect("alloc scratch (1 f32)");
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
isv.host_slice_mut()[PNL_REWARD_MAGNITUDE_EMA_INDEX] = 1.0;
let delta_dev = delta_reward.dev_ptr;
let scratch_dev = scratch.dev_ptr;
let isv_dev = isv.dev_ptr;
let n_bars_i32: i32 = 4;
let pnl_slot_i32: i32 = PNL_REWARD_MAGNITUDE_EMA_INDEX as i32;
unsafe {
stream
.launch_builder(&f)
.arg(&delta_dev)
.arg(&scratch_dev)
.arg(&isv_dev)
.arg(&n_bars_i32)
.arg(&pnl_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: 256 * std::mem::size_of::<i32>() as u32,
})
.expect("launch saboteur_engagement_compute_kernel");
}
stream.synchronize().expect("sync after saboteur engagement launch");
let engagement = scratch.read_all()[0];
assert!(
(engagement - 0.5_f32).abs() < 1e-4,
"expected engagement = 0.5 (2/4 bars above threshold), got {engagement}"
);
}
/// SP11 A1.3 — per-component magnitude-ratio normaliser preserves Σ=1
/// and mirrors popart magnitude into scratch[6].
///
/// Inputs: ISV[REWARD_POPART_EMA_INDEX..+6) = [1.0, 2.0, 3.0, 4.0, 5.0,
/// 6.0]. Sum = 21. Expected ratios = [1/21, 2/21, 3/21, 4/21, 5/21,
/// 6/21]. Mirror at scratch[6] = 1.0 (popart magnitude unchanged).
/// Σ ratios = 1.0 ± 1e-5.
#[test]
#[ignore = "requires GPU"]
fn reward_component_mag_ratio_normalises_and_mirrors_popart() {
let stream = make_test_stream();
let f = load_kernel(&stream, SP11_MAG_RATIO_CUBIN, "reward_component_mag_ratio_compute_kernel");
let scratch = unsafe { MappedF32Buffer::new(7) }
.expect("alloc scratch (7 f32 — 6 ratios + 1 popart mirror)");
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
{
let s = isv.host_slice_mut();
s[REWARD_POPART_EMA_INDEX + 0] = 1.0;
s[REWARD_POPART_EMA_INDEX + 1] = 2.0;
s[REWARD_POPART_EMA_INDEX + 2] = 3.0;
s[REWARD_POPART_EMA_INDEX + 3] = 4.0;
s[REWARD_POPART_EMA_INDEX + 4] = 5.0;
s[REWARD_POPART_EMA_INDEX + 5] = 6.0;
}
let isv_dev = isv.dev_ptr;
let scratch_dev = scratch.dev_ptr;
let popart_base_i32: i32 = REWARD_POPART_EMA_INDEX as i32;
unsafe {
stream
.launch_builder(&f)
.arg(&isv_dev)
.arg(&scratch_dev)
.arg(&popart_base_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (6, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch reward_component_mag_ratio_compute_kernel");
}
stream.synchronize().expect("sync after mag_ratio launch");
let out = scratch.read_all();
let sum_in: f32 = 1.0 + 2.0 + 3.0 + 4.0 + 5.0 + 6.0;
// Per-component ratio assertions.
for c in 0..6_usize {
let expected = (c as f32 + 1.0) / sum_in;
assert!(
(out[c] - expected).abs() < 1e-5,
"component {c}: expected ratio {expected}, got {}",
out[c]
);
}
// Σ ratios == 1.0 ± 1e-5.
let sum_out: f32 = out[0..6].iter().sum();
assert!(
(sum_out - 1.0_f32).abs() < 1e-5,
"expected Σratios = 1.0, got {sum_out}"
);
// Popart mirror at scratch[6].
assert!(
(out[6] - 1.0_f32).abs() < 1e-6,
"expected popart mirror = 1.0 (= isv[REWARD_POPART_EMA_INDEX]), got {}",
out[6]
);
// Reference index touch — silence dead_code on the import even when
// a future test removes it. (REWARD_COMPONENT_MAG_RATIO_BASE = 352
// is the ISV-side destination of the first 6 scratch entries via
// the chained Pearls A+D launch.)
let _ = REWARD_COMPONENT_MAG_RATIO_BASE;
}

37
docs/isv-slots.md
View File

@@ -184,15 +184,42 @@ 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 (Task A0 reservation)
## SP11: Reward as controlled subsystem (Tasks A0 + A1)
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
controller. Constants live in `crates/ml/src/cuda_pipeline/sp11_isv_slots.rs`
and are re-exported via `cuda_pipeline::sp5_isv_slots::*`. **A0 is reservation-
only — no producer kernels or consumers are wired yet. All 20 slots remain
zero-initialized at construction and behave as no-ops until A1/A2/Layer-B
land producers and atomic consumer migration.**
and are re-exported via `cuda_pipeline::sp5_isv_slots::*`.
**Task A1 (2026-05-04): three canary producer kernels added**, each
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.
A1 producers:
- `val_sharpe_delta_compute_kernel.cu` — two-pass: raw Δ + squared
deviation against prior `VAL_SHARPE_DELTA_EMA_INDEX`. Reads
mapped-pinned 2-element val_sharpe_history; chain → ISV[350, 351].
- `saboteur_engagement_compute_kernel.cu` — block tree-reduce over
per-bar `|Δreward|` populated by `experience_env_step` at the
saboteur perturbation site (proxy: `traded × |reward| ×
max(|eff_spread − 1|, |eff_slip − 1|)`); threshold = 0.01 ×
ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX]. Chain → ISV[358].
- `reward_component_mag_ratio_compute_kernel.cu` — reads 6 EMAs at
ISV[REWARD_POPART_EMA_INDEX..+6), normalises to ratios, mirrors
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
htod_copy/dtoh_sync_copy/alloc_zeros per
`feedback_no_htod_htoh_only_mapped_pinned`).
| Range | Name constant | Family | Purpose |
|---|---|---|---|