diff --git a/crates/ml/build.rs b/crates/ml/build.rs index fb26c76c2..c81abd405 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -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 diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 5709ed7c1..f0bd39841 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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 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. diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 059219c29..df8951e33 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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) + /// — 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::() 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::() 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::() 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::() 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, diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index dbd4c3ba8..74e9bba3f 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -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, // [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, // [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::(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::(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}" diff --git a/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu b/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu new file mode 100644 index 000000000..c14168a9e --- /dev/null +++ b/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu @@ -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 + +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(); + } +} diff --git a/crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu b/crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu new file mode 100644 index 000000000..a66a311d3 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/saboteur_engagement_compute_kernel.cu @@ -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 + +#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(); + } +} diff --git a/crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu b/crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu new file mode 100644 index 000000000..444b5b4b6 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu @@ -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 + +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(); +} diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index f91b3af27..48be6d748 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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 diff --git a/crates/ml/tests/sp11_producer_unit_tests.rs b/crates/ml/tests/sp11_producer_unit_tests.rs new file mode 100644 index 000000000..4e0ad2f5f --- /dev/null +++ b/crates/ml/tests/sp11_producer_unit_tests.rs @@ -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 { + 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, 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::() 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; +} diff --git a/docs/isv-slots.md b/docs/isv-slots.md index 5e38bd4a6..b179cd934 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -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 | |---|---|---|---|