From b92dcc3dfc66dc7e8ff69b55bb941248fc0955a2 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 4 May 2026 15:14:54 +0200 Subject: [PATCH] chore(sp11): remove reward-chain diagnostic instrumentation MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Instrumentation from 774d7552a served its purpose — empirically identified the asymmetric reward cap at experience_kernels.cu:2788 as the inflater (popart min reaching -210336 by F0 ep3 pre-fix). Fixed in 35db31089; validated by smoke-test-trk72 (PASSED). Removed: - reward_chain_diag_reduce_kernel.cu - 12 per-sample diagnostic buffers in gpu_experience_collector - kernel parameter threading in experience_kernels.cu - launcher + reader + mapped-pinned output in gpu_dqn_trainer - wire-up site + HEALTH_DIAG emit in training_loop - audit doc section for the transient instrumentation This brings the worktree back to its pre-instrumentation state on the SP11 reward chain. Symmetric reward cap fix (35db31089) and plan_isv symmetric clamps (Commit A) remain. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 8 - .../src/cuda_pipeline/experience_kernels.cu | 96 +--------- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 167 ----------------- .../cuda_pipeline/gpu_experience_collector.rs | 134 -------------- .../reward_chain_diag_reduce_kernel.cu | 170 ------------------ .../src/trainers/dqn/trainer/training_loop.rs | 88 --------- docs/dqn-wire-up-audit.md | 115 ------------ 7 files changed, 1 insertion(+), 777 deletions(-) delete mode 100644 crates/ml/src/cuda_pipeline/reward_chain_diag_reduce_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index eb9c4bece..1677f1b3b 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -573,14 +573,6 @@ fn main() { // (contamination → controller over-weighed popart → 10× sharpe // drop). Spec §4 amendment "Why slot 360". "popart_component_ema_kernel.cu", - // SP11 instrumentation (2026-05-04, diagnostic-only): block- - // tree-reduce for the 12 per-sample reward-chain diagnostic - // buffers populated by `experience_env_step`. Outputs (min, - // mean, max) per buffer = 36 floats to a mapped-pinned scratch - // buffer the host reads at HEALTH_DIAG emit time. Single block, - // 256 threads, no atomicAdd. Will be removed in the follow-up - // commit once the inflater is identified and properly fixed. - "reward_chain_diag_reduce_kernel.cu", // SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller + // SimHash novelty buffer. Three new kernels in this layer: // 1. reward_subsystem_controller — main producer (5 canaries → 10 diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index dcc10f6b1..a65f3b694 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1833,32 +1833,7 @@ extern "C" __global__ void experience_env_step( * guards every write with the NULL check. Pure GPU buffer (the * trainer holds it as a `MappedF32Buffer`); reads happen only inside * `popart_component_ema_kernel` (kernel-to-kernel dataflow). */ - float* __restrict__ popart_component_per_sample, - /* SP11 instrumentation (2026-05-04, diagnostic-only): 12 per-sample - * reward-chain checkpoint buffers, each [N*L] f32. Smoke - * `smoke-test-gwfn8` on `fd24b5383` showed mean(|reward|) hitting - * 5054 at F0 ep2 despite all known multiplicative modifiers being - * structurally bounded. The inflater is somewhere in the pre- - * composition or modifier chain that isn't currently visible in - * HEALTH_DIAG. These buffers capture the value of the running - * `reward` cascade variable (and the per-component locals) at every - * checkpoint so the consumer reduction kernel can produce - * (min, mean, max) per stage. NULL-tolerant — pre-wire-up early - * init paths skip the writes via guards. Will be removed in the - * follow-up commit once the inflater is identified and properly - * fixed per pearl_bounded_modifier_outputs_require_structural_activation. */ - float* __restrict__ diag_r_popart_per_sample, /* [N*L] r_popart pre-composition */ - float* __restrict__ diag_r_trail_per_sample, /* [N*L] r_trail pre-composition */ - float* __restrict__ diag_r_micro_per_sample, /* [N*L] r_micro pre-composition */ - float* __restrict__ diag_r_opp_cost_per_sample, /* [N*L] r_opp_cost pre-composition */ - float* __restrict__ diag_r_bonus_per_sample, /* [N*L] r_bonus pre-composition */ - float* __restrict__ diag_r_weighted_per_sample, /* [N*L] r_weighted post-composition, pre-modifier */ - float* __restrict__ diag_post_dd_per_sample, /* [N*L] reward after drawdown penalty */ - float* __restrict__ diag_post_inv_per_sample, /* [N*L] reward after inventory penalty */ - float* __restrict__ diag_post_churn_per_sample, /* [N*L] reward after churn penalty */ - float* __restrict__ diag_post_conv_per_sample, /* [N*L] reward after conviction scale */ - float* __restrict__ diag_position_abs_per_sample, /* [N*L] |position| sanity */ - float* __restrict__ diag_conviction_per_sample /* [N*L] conviction value sanity (1.0 when modifier inert) */ + float* __restrict__ popart_component_per_sample ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -1955,25 +1930,6 @@ extern "C" __global__ void experience_env_step( if (popart_component_per_sample != NULL) { popart_component_per_sample[out_off] = 0.0f; } - /* SP11 instrumentation (diagnostic-only): default per-sample reward- - * chain checkpoint values. Real values are overwritten at each - * checkpoint below. The default 0 carries through to the reduction - * kernel as "stage didn't fire on this (i,t)" — non-firing slots - * dilute the mean toward 0 but don't poison min/max (most stages can - * legitimately produce 0). NULL-tolerant for early init before the - * collector wires the buffers. */ - if (diag_r_popart_per_sample != NULL) diag_r_popart_per_sample[out_off] = 0.0f; - if (diag_r_trail_per_sample != NULL) diag_r_trail_per_sample[out_off] = 0.0f; - if (diag_r_micro_per_sample != NULL) diag_r_micro_per_sample[out_off] = 0.0f; - if (diag_r_opp_cost_per_sample != NULL) diag_r_opp_cost_per_sample[out_off] = 0.0f; - if (diag_r_bonus_per_sample != NULL) diag_r_bonus_per_sample[out_off] = 0.0f; - if (diag_r_weighted_per_sample != NULL) diag_r_weighted_per_sample[out_off] = 0.0f; - if (diag_post_dd_per_sample != NULL) diag_post_dd_per_sample[out_off] = 0.0f; - if (diag_post_inv_per_sample != NULL) diag_post_inv_per_sample[out_off] = 0.0f; - if (diag_post_churn_per_sample != NULL) diag_post_churn_per_sample[out_off] = 0.0f; - if (diag_post_conv_per_sample != NULL) diag_post_conv_per_sample[out_off] = 0.0f; - if (diag_position_abs_per_sample != NULL) diag_position_abs_per_sample[out_off] = 0.0f; - if (diag_conviction_per_sample != NULL) diag_conviction_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; @@ -3359,21 +3315,6 @@ extern "C" __global__ void experience_env_step( * value sourced from the spec — `feedback_no_quickfixes` carve-out * for Invariant-1 anchors per spec §3.4.2. * ================================================================ */ - /* SP11 instrumentation (diagnostic-only): capture each per-component - * local before the controller-weighted composition. Five mutually- - * exclusive trade-state branches set exactly one of {r_popart, - * r_trail, r_micro, r_opp_cost} per (i,t); r_bonus accumulates - * additively across the entry/exit bonus sites. Sanity check: - * |position| (before the capital-floor branch potentially zeros it - * below). The conviction sanity write happens inline at the - * conviction modifier site (~line 3520) where its value is in scope. */ - if (diag_r_popart_per_sample != NULL) diag_r_popart_per_sample[out_off] = r_popart; - if (diag_r_trail_per_sample != NULL) diag_r_trail_per_sample[out_off] = r_trail; - if (diag_r_micro_per_sample != NULL) diag_r_micro_per_sample[out_off] = r_micro; - if (diag_r_opp_cost_per_sample != NULL) diag_r_opp_cost_per_sample[out_off] = r_opp_cost; - if (diag_r_bonus_per_sample != NULL) diag_r_bonus_per_sample[out_off] = r_bonus; - if (diag_position_abs_per_sample != NULL) diag_position_abs_per_sample[out_off] = fabsf(position); - { /* On-policy composition: 5 components (popart / trail / micro / * opp_cost / bonus). The cf-reward path is composed separately in @@ -3405,14 +3346,6 @@ extern "C" __global__ void experience_env_step( + w_oc * r_opp_cost + w_bn * r_bonus; reward = r_weighted; /* the post-composition modifiers below operate on r_weighted */ - /* SP11 instrumentation: post-composition, pre-modifier checkpoint. - * If r_weighted is already large, the inflater is upstream - * (mis-bounded component or runaway weight). If r_weighted is - * small but post_* checkpoints below balloon, the inflater is - * one of the modifier sites. */ - if (diag_r_weighted_per_sample != NULL) { - diag_r_weighted_per_sample[out_off] = r_weighted; - } } /* ---- Drawdown penalty (from trade_physics.cuh) ────────────────────── @@ -3422,10 +3355,6 @@ extern "C" __global__ void experience_env_step( * (validation mode) we measure pure equity-change reward only. */ float equity_now = cash + position * raw_close; reward += shaping_scale * compute_drawdown_penalty(equity_now, peak_equity, dd_threshold, w_dd); - /* SP11 instrumentation: reward after drawdown penalty. */ - if (diag_post_dd_per_sample != NULL) { - diag_post_dd_per_sample[out_off] = reward; - } /* ---- Portfolio value for floor check ---- */ float new_portfolio_value_pre_floor = cash + position * raw_close; @@ -3535,10 +3464,6 @@ extern "C" __global__ void experience_env_step( * Phase 3: gated by shaping_scale (behavioral; not in validation reward). */ float inventory_penalty = holding_cost_rate * fabsf(position); reward -= shaping_scale * inventory_penalty; - /* SP11 instrumentation: reward after inventory penalty. */ - if (diag_post_inv_per_sample != NULL) { - diag_post_inv_per_sample[out_off] = reward; - } /* Churn penalty: graduated cost for rapid position flips. * Not a hard gate — the model CAN exit early, but it costs extra. @@ -3548,36 +3473,17 @@ extern "C" __global__ void experience_env_step( float churn_penalty = churn_penalty_scale * (churn_threshold - hold_time) / churn_threshold; reward -= shaping_scale * churn_penalty; } - /* SP11 instrumentation: reward after churn penalty (or unchanged if - * the gate didn't fire — same write-the-running-reward semantics as - * the inventory checkpoint above). */ - if (diag_post_churn_per_sample != NULL) { - diag_post_churn_per_sample[out_off] = reward; - } /* Plan conviction as reward scaling — ONLY when plan head is mature. * Before readiness >= plan_thr_env, conviction defaults to 1.0 (identity) so the * reward signal flows undampened. Without this gate, random init conviction * (~0.1) kills the gradient signal and prevents the model from learning. */ - /* SP11 instrumentation: track the effective conviction used by the - * modifier. Defaults to 1.0 (identity) when the gate is inert — that - * matches the kernel semantics ("reward unchanged"). */ - float diag_conviction_value = 1.0f; if (plan_params_ptr != NULL && fabsf(reward) > 1e-8f) { float readiness_val = (readiness_ptr != NULL) ? readiness_ptr[0] : 0.0f; float conviction = (readiness_val >= plan_thr_env) ? plan_params_ptr[i * 6 + PLAN_PARAM_CONVICTION] /* [0, 1] from plan head */ : 1.0f; /* identity until plan matures */ reward *= conviction; - diag_conviction_value = conviction; - } - /* SP11 instrumentation: reward after conviction scale, plus the - * conviction value itself for sanity (should be ∈ [0, 1] in scope). */ - if (diag_post_conv_per_sample != NULL) { - diag_post_conv_per_sample[out_off] = reward; - } - if (diag_conviction_per_sample != NULL) { - diag_conviction_per_sample[out_off] = diag_conviction_value; } /* G7: Counterfactual reward flip — negate reward when features were flipped. diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 2ca709323..12aec6739 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -577,16 +577,6 @@ static SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN: &[u8] = static SP11_POPART_COMPONENT_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/popart_component_ema_kernel.cubin")); -/// SP11 instrumentation (2026-05-04, diagnostic-only): block-tree-reduce -/// for the 12 per-sample reward-chain diagnostic buffers populated by -/// `experience_env_step`. Outputs (min, mean, max) per buffer = 36 -/// floats to a mapped-pinned scratch buffer the host reads at -/// HEALTH_DIAG emit time. Will be removed in the follow-up commit once -/// the inflater is identified and properly fixed per -/// pearl_bounded_modifier_outputs_require_structural_activation. -static SP11_REWARD_CHAIN_DIAG_REDUCE_CUBIN: &[u8] = - include_bytes!(concat!(env!("OUT_DIR"), "/reward_chain_diag_reduce_kernel.cubin")); - /// SP11 Fix 39 (2026-05-04, Task A2): one-shot GPU init for the SimHash /// projection matrix. Philox-driven (deterministic from a host-passed seed /// derived from the SP11 salt) writes ±1 values into a 42×16 = 672-element @@ -4936,28 +4926,6 @@ pub struct GpuDqnTrainer { /// unwired). pub(crate) sp11_popart_component_dev_ptr: u64, pub(crate) sp11_popart_component_len: usize, - /// SP11 instrumentation (2026-05-04, diagnostic-only): block-tree- - /// reduce kernel for the 12 per-sample reward-chain checkpoint - /// buffers populated by `experience_env_step`. Outputs (min, mean, - /// max) per buffer = 36 floats to - /// `sp11_reward_chain_diag_out_pinned`. Loaded from - /// `reward_chain_diag_reduce_kernel.cubin`. Will be removed in the - /// follow-up commit once the inflater is identified and properly - /// fixed per pearl_bounded_modifier_outputs_require_structural_activation. - sp11_reward_chain_diag_reduce_kernel: CudaFunction, - /// SP11 instrumentation (2026-05-04, diagnostic-only): mapped-pinned - /// 36 f32 output buffer for the reward-chain diagnostic reduction. - /// Layout = 3 stats × 12 buffers (see kernel docstring). The host - /// reads via `read_sp11_reward_chain_diag` after stream sync. - pub(crate) sp11_reward_chain_diag_out_pinned: super::mapped_pinned::MappedF32Buffer, - /// SP11 instrumentation (2026-05-04, diagnostic-only): device - /// pointers for the 12 per-sample reward-chain checkpoint buffers, - /// owned by the experience collector. Wired in the same commit by - /// `set_sp11_reward_chain_diag_bufs` mirroring the popart-component - /// pattern. Initial 0s / 0 (kernel guarded — - /// `launch_sp11_reward_chain_diag_reduce` no-ops when unwired). - pub(crate) sp11_reward_chain_diag_dev_ptrs: [u64; 12], - pub(crate) sp11_reward_chain_diag_len: usize, /// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller /// kernel. Single-block, 10-thread producer reading 5 canary ISV slots /// [350..360) and writing 10 floats to scratch[SCRATCH_SP11_CONTROLLER_BASE..+10). @@ -12687,43 +12655,6 @@ impl GpuDqnTrainer { self.sp11_popart_component_len = n_bars; } - /// SP11 instrumentation (2026-05-04, diagnostic-only): wire the - /// experience collector's 12 per-sample reward-chain checkpoint - /// buffers into the trainer so `launch_sp11_reward_chain_diag_reduce` - /// reads the right buffers. Mirror of `set_sp11_popart_component_buf` - /// — the collector owns the per-bar buffers (alloc_episodes × - /// alloc_timesteps each); the trainer caches the pointers + length - /// for the per-step launcher path. Per `feedback_wire_everything_up`: - /// this is wired in the same commit that adds the buffers + kernel. - /// Will be removed in the follow-up commit once the inflater is - /// identified and properly fixed. - pub(crate) fn set_sp11_reward_chain_diag_bufs( - &mut self, - dev_ptrs: [u64; 12], - n_bars: usize, - ) { - self.sp11_reward_chain_diag_dev_ptrs = dev_ptrs; - self.sp11_reward_chain_diag_len = n_bars; - } - - /// SP11 instrumentation (2026-05-04, diagnostic-only): host read of - /// the 36-element reward-chain diagnostic output buffer. Caller MUST - /// have synchronised the producer stream first (the reduction - /// kernel writes via mapped-pinned with `__threadfence_system()`, - /// host visibility requires the post-launch sync). Returns - /// `[buf_idx*3 + 0/1/2] = (min, mean, max)` for buf_idx ∈ 0..12 in - /// the order documented in - /// `reward_chain_diag_reduce_kernel.cu` (r_popart, r_trail, r_micro, - /// r_opp_cost, r_bonus, r_weighted, post_dd, post_inv, post_churn, - /// post_conv, position_abs, conviction). - pub(crate) fn read_sp11_reward_chain_diag(&self) -> [f32; 36] { - let raw = self.sp11_reward_chain_diag_out_pinned.read_all(); - let mut out = [0.0_f32; 36]; - let n = raw.len().min(36); - out[..n].copy_from_slice(&raw[..n]); - out - } - /// SP11 Fix 39 (2026-05-04, A1.1): launch the val-sharpe Δ + variance /// canary producer. /// @@ -13113,73 +13044,6 @@ impl GpuDqnTrainer { Ok(()) } - /// SP11 instrumentation (2026-05-04, diagnostic-only): launch the - /// reward-chain diagnostic block-tree-reduce. - /// - /// Reads the 12 per-sample reward-chain checkpoint buffers (wired in - /// via `set_sp11_reward_chain_diag_bufs`) and writes (min, mean, - /// max) per buffer = 36 floats to `sp11_reward_chain_diag_out_pinned`. - /// Single block, 256 threads, shared_mem = 3 × 256 f32 = 3 KiB - /// (sum + min + max scratch arrays). NO atomicAdd per - /// `feedback_no_atomicadd.md`. Pure GPU compute per - /// `feedback_no_cpu_compute_strict.md`. - /// - /// No-op when the buffers haven't been wired yet (early init before - /// the collector exists). Mirrors the popart_component launcher's - /// NULL-guard pattern. - /// - /// Will be removed in the follow-up commit once the inflater is - /// identified and properly fixed. - pub(crate) fn launch_sp11_reward_chain_diag_reduce(&self) -> Result<(), MLError> { - // Precondition guard — buffers wired by - // `set_sp11_reward_chain_diag_bufs` after collector construction. - // Until then, no-op (NOT a stub return value per - // `feedback_no_stubs.md` — a precondition guard that defers the - // launch). Once wired, every subsequent step launches. - if self.sp11_reward_chain_diag_len == 0 { - return Ok(()); - } - - let n_bars_i32: i32 = i32::try_from(self.sp11_reward_chain_diag_len) - .unwrap_or(i32::MAX); - let out_dev = self.sp11_reward_chain_diag_out_pinned.dev_ptr; - let p = &self.sp11_reward_chain_diag_dev_ptrs; - - // The 12 per-sample buffer device pointers must be passed by - // value as kernel args. cudarc's launch_builder accepts u64 args - // for raw device pointers. Order MUST match the kernel parameter - // declaration in `reward_chain_diag_reduce_kernel.cu`. - unsafe { - self.stream - .launch_builder(&self.sp11_reward_chain_diag_reduce_kernel) - .arg(&p[0]) // r_popart - .arg(&p[1]) // r_trail - .arg(&p[2]) // r_micro - .arg(&p[3]) // r_opp_cost - .arg(&p[4]) // r_bonus - .arg(&p[5]) // r_weighted - .arg(&p[6]) // post_dd - .arg(&p[7]) // post_inv - .arg(&p[8]) // post_churn - .arg(&p[9]) // post_conv - .arg(&p[10]) // position_abs - .arg(&p[11]) // conviction - .arg(&out_dev) - .arg(&n_bars_i32) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (256, 1, 1), - /* sum + min + max shared scratch, BLOCK_SIZE=256 each. */ - shared_mem_bytes: 3 * 256 * std::mem::size_of::() as u32, - }) - .map_err(|e| MLError::ModelError( - format!("sp11 reward_chain_diag_reduce: {e}") - ))?; - } - - Ok(()) - } - /// SP11 Fix 39 (2026-05-04, Task A2): launch the reward-subsystem /// controller producer + chained Pearls A+D output smoothing. /// @@ -16640,25 +16504,6 @@ impl GpuDqnTrainer { format!("SP11 popart_component_per_sample alloc: {e}") ))?; - // SP11 instrumentation (2026-05-04, diagnostic-only): block-tree- - // reduce kernel + 36 f32 mapped-pinned output for the reward- - // chain diagnostic chain. Will be removed in the follow-up - // commit once the inflater is identified and properly fixed. - let sp11_reward_chain_diag_reduce_kernel = { - let module = stream.context() - .load_cubin(SP11_REWARD_CHAIN_DIAG_REDUCE_CUBIN.to_vec()) - .map_err(|e| MLError::ModelError(format!("sp11 reward_chain_diag_reduce cubin load: {e}")))?; - module.load_function("reward_chain_diag_reduce_kernel") - .map_err(|e| MLError::ModelError(format!("sp11 reward_chain_diag_reduce load: {e}")))? - }; - // 36 f32 = 3 stats (min, mean, max) × 12 buffers. Layout matches - // the kernel's `out_min_mean_max[buf_idx*3 + stat]`. - let sp11_reward_chain_diag_out_pinned = unsafe { - super::mapped_pinned::MappedF32Buffer::new(36) - }.map_err(|e| MLError::ModelError( - format!("SP11 reward_chain_diag_out_pinned alloc: {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 @@ -20068,18 +19913,6 @@ impl GpuDqnTrainer { sp11_popart_component_ema_kernel, sp11_popart_component_dev_ptr: 0, sp11_popart_component_len: 0, - // SP11 instrumentation (2026-05-04, diagnostic-only): reward- - // chain diagnostic kernel + mapped-pinned output buffer. The - // 12 per-sample buffer device pointers are owned by the - // experience collector; cached here as 0 / 0 and wired in - // via `set_sp11_reward_chain_diag_bufs` after collector - // construction (training_loop.rs init path). Will be - // removed in the follow-up commit once the inflater is - // identified and properly fixed. - sp11_reward_chain_diag_reduce_kernel, - sp11_reward_chain_diag_out_pinned, - sp11_reward_chain_diag_dev_ptrs: [0u64; 12], - sp11_reward_chain_diag_len: 0, // SP11 Fix 39 (Task A2): controller + SimHash novelty kernels // and their backing buffers. Projection matrix was populated // on-device above by `launch_novelty_simhash_proj_init`; hash diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 67b1acfb5..69e31c793 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -628,29 +628,6 @@ pub struct GpuExperienceCollector { /// "Why slot 360". Pure GPU buffer (kernel-to-kernel only); the SP11 /// popart-component EMA kernel reads via raw device pointer. popart_component_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - /// SP11 instrumentation (2026-05-04, diagnostic-only): 12 per-sample - /// reward-chain checkpoint buffers. Smoke `smoke-test-gwfn8` showed - /// mean(|reward|) hitting 5054 at F0 ep2 despite all known - /// multiplicative modifiers being structurally bounded. These - /// buffers capture the running `reward` cascade variable (and per- - /// component locals) at every checkpoint so the consumer reduction - /// kernel can produce (min, mean, max) per stage. Pure GPU buffers - /// (kernel-to-kernel only); the trainer reads through the reduction - /// kernel which writes a 36-slot mapped-pinned scratch buffer. - /// Will be removed in the follow-up commit once the inflater is - /// identified and properly fixed. - diag_r_popart_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_r_trail_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_r_micro_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_r_opp_cost_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_r_bonus_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_r_weighted_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_post_dd_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_post_inv_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_post_churn_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_post_conv_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_position_abs_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] - diag_conviction_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. @@ -1248,50 +1225,6 @@ impl GpuExperienceCollector { let popart_component_per_sample = stream .alloc_zeros::(total_output) .map_err(|e| MLError::ModelError(format!("alloc popart_component_per_sample: {e}")))?; - // SP11 instrumentation (2026-05-04, diagnostic-only): 12 per-sample - // reward-chain checkpoint buffers, each [total_output] f32. The - // experience kernel writes the running `reward` cascade variable - // (and per-component locals) at every checkpoint; the SP11 - // reduction kernel computes (min, mean, max) per buffer for the - // HEALTH_DIAG `reward_chain_diag` line. Will be removed in the - // follow-up commit once the inflater is identified and properly - // fixed per pearl_bounded_modifier_outputs_require_structural_activation. - let diag_r_popart_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_r_popart_per_sample: {e}")))?; - let diag_r_trail_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_r_trail_per_sample: {e}")))?; - let diag_r_micro_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_r_micro_per_sample: {e}")))?; - let diag_r_opp_cost_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_r_opp_cost_per_sample: {e}")))?; - let diag_r_bonus_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_r_bonus_per_sample: {e}")))?; - let diag_r_weighted_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_r_weighted_per_sample: {e}")))?; - let diag_post_dd_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_post_dd_per_sample: {e}")))?; - let diag_post_inv_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_post_inv_per_sample: {e}")))?; - let diag_post_churn_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_post_churn_per_sample: {e}")))?; - let diag_post_conv_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_post_conv_per_sample: {e}")))?; - let diag_position_abs_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_position_abs_per_sample: {e}")))?; - let diag_conviction_per_sample = stream - .alloc_zeros::(total_output) - .map_err(|e| MLError::ModelError(format!("alloc diag_conviction_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). @@ -1792,21 +1725,6 @@ impl GpuExperienceCollector { total_reward_per_sample, saboteur_delta_reward_per_sample, popart_component_per_sample, - // SP11 instrumentation (2026-05-04, diagnostic-only): 12 per- - // sample reward-chain checkpoint buffers. Will be removed in - // the follow-up commit once the inflater is identified. - diag_r_popart_per_sample, - diag_r_trail_per_sample, - diag_r_micro_per_sample, - diag_r_opp_cost_per_sample, - diag_r_bonus_per_sample, - diag_r_weighted_per_sample, - diag_post_dd_per_sample, - diag_post_inv_per_sample, - diag_post_churn_per_sample, - diag_post_conv_per_sample, - diag_position_abs_per_sample, - diag_conviction_per_sample, reward_components_per_sample, flat_to_pos_per_sample, readiness_per_sample, @@ -2252,38 +2170,6 @@ impl GpuExperienceCollector { (dev_ptr, n_bars) } - /// SP11 instrumentation (2026-05-04, diagnostic-only): expose the 12 - /// per-sample reward-chain checkpoint buffer device pointers + the - /// shared length to the trainer. Returns - /// `(ptrs[12], n_bars)` where `n_bars = alloc_episodes × - /// alloc_timesteps`. The trainer caches the array via - /// `GpuDqnTrainer::set_sp11_reward_chain_diag_bufs` after the - /// collector is constructed (training_loop.rs init path) so - /// `launch_sp11_reward_chain_diag_reduce` has the consumer-side - /// pointers ready every step. Order MUST match the kernel parameter - /// declaration order (r_popart, r_trail, r_micro, r_opp_cost, - /// r_bonus, r_weighted, post_dd, post_inv, post_churn, post_conv, - /// position_abs, conviction). Will be removed in the follow-up - /// commit once the inflater is identified. - pub fn sp11_reward_chain_diag_dev_ptrs(&self) -> ([u64; 12], usize) { - let ptrs = [ - self.diag_r_popart_per_sample.device_ptr(&self.stream).0, - self.diag_r_trail_per_sample.device_ptr(&self.stream).0, - self.diag_r_micro_per_sample.device_ptr(&self.stream).0, - self.diag_r_opp_cost_per_sample.device_ptr(&self.stream).0, - self.diag_r_bonus_per_sample.device_ptr(&self.stream).0, - self.diag_r_weighted_per_sample.device_ptr(&self.stream).0, - self.diag_post_dd_per_sample.device_ptr(&self.stream).0, - self.diag_post_inv_per_sample.device_ptr(&self.stream).0, - self.diag_post_churn_per_sample.device_ptr(&self.stream).0, - self.diag_post_conv_per_sample.device_ptr(&self.stream).0, - self.diag_position_abs_per_sample.device_ptr(&self.stream).0, - self.diag_conviction_per_sample.device_ptr(&self.stream).0, - ]; - let n_bars = self.alloc_episodes * self.alloc_timesteps; - (ptrs, n_bars) - } - /// Upload pre-computed expert demonstration actions to GPU. /// /// `expert_actions[bar_index]` = expert exposure action index (-1 = no opinion). @@ -4222,26 +4108,6 @@ impl GpuExperienceCollector { // in training_loop.rs init). Spec §4 amendment "Why // slot 360". .arg(&mut self.popart_component_per_sample) - // SP11 instrumentation (2026-05-04, diagnostic-only): - // 12 per-sample reward-chain checkpoint buffers. Order - // MUST match the kernel parameter declaration order in - // `experience_kernels.cu` (r_popart, r_trail, r_micro, - // r_opp_cost, r_bonus, r_weighted, post_dd, post_inv, - // post_churn, post_conv, position_abs, conviction). - // Will be removed in the follow-up commit once the - // inflater is identified and properly fixed. - .arg(&mut self.diag_r_popart_per_sample) - .arg(&mut self.diag_r_trail_per_sample) - .arg(&mut self.diag_r_micro_per_sample) - .arg(&mut self.diag_r_opp_cost_per_sample) - .arg(&mut self.diag_r_bonus_per_sample) - .arg(&mut self.diag_r_weighted_per_sample) - .arg(&mut self.diag_post_dd_per_sample) - .arg(&mut self.diag_post_inv_per_sample) - .arg(&mut self.diag_post_churn_per_sample) - .arg(&mut self.diag_post_conv_per_sample) - .arg(&mut self.diag_position_abs_per_sample) - .arg(&mut self.diag_conviction_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_chain_diag_reduce_kernel.cu b/crates/ml/src/cuda_pipeline/reward_chain_diag_reduce_kernel.cu deleted file mode 100644 index 909a4cac0..000000000 --- a/crates/ml/src/cuda_pipeline/reward_chain_diag_reduce_kernel.cu +++ /dev/null @@ -1,170 +0,0 @@ -// crates/ml/src/cuda_pipeline/reward_chain_diag_reduce_kernel.cu -// -// SP11 instrumentation (2026-05-04, diagnostic-only): block-tree-reduce -// for the 12 per-sample reward-chain diagnostic buffers populated by -// `experience_env_step`. -// -// Smoke `smoke-test-gwfn8` on `fd24b5383` showed mean(|reward|) hitting -// 5054 at F0 ep2 despite all known multiplicative modifiers being -// structurally bounded (capped_pnl ≤ 10, conviction ∈ (0,1) via sigmoid, -// cf_flip ∈ ±1, drawdown ∈ [-5*w_dd, 0], shaping_scale ∈ [0,1]). The -// inflater is somewhere in the pre-composition or modifier chain that -// isn't currently visible in HEALTH_DIAG. This kernel block-reduces 12 -// per-sample buffers to (min, mean, max) per buffer = 36 scalars and -// writes them into a single mapped-pinned scratch buffer the host reads -// at HEALTH_DIAG emit time. -// -// Layout of `out_min_mean_max`: 36 floats, 3 per buffer in the order -// [buf_idx*3 + 0] = min -// [buf_idx*3 + 1] = mean -// [buf_idx*3 + 2] = max -// where `buf_idx` matches the order of the 12 const-pointer inputs: -// 0 = r_popart, 1 = r_trail, 2 = r_micro, 3 = r_opp_cost, 4 = r_bonus, -// 5 = r_weighted, 6 = post_dd, 7 = post_inv, 8 = post_churn, -// 9 = post_conv, 10 = position_abs, 11 = conviction. -// -// Single block (BLOCK_SIZE=256). Per-buffer the kernel does THREE -// block tree-reduces (sum / min / max) in shared memory, all initialised -// from a strided thread-local fold over the buffer. NO atomicAdd per -// `feedback_no_atomicadd.md`. Pure GPU compute per -// `feedback_no_cpu_compute_strict.md`. -// -// This is INSTRUMENTATION (transient). Removed in the follow-up commit -// once the inflater is identified and properly fixed. No ISV slot -// allocations — output goes to a dedicated mapped-pinned diagnostic -// buffer the host reads directly. - -#include -#include - -#define BLOCK_SIZE 256 -#define NUM_DIAG_BUFFERS 12 -#define STATS_PER_BUFFER 3 /* min, mean, max */ - -/* Per-buffer reduction helper. Reads `n_samples` floats from `src`, - * tree-reduces sum, min, max into shared memory, writes to - * `out_min_mean_max[buf_idx*3..+3)`. NULL-tolerant — if `src` is null - * (buffer not wired yet) the slot stays 0. */ -__device__ __forceinline__ void reduce_one_buffer( - const float* __restrict__ src, - int n_samples, - int buf_idx, - float* __restrict__ out_min_mean_max, - float* __restrict__ shared_sum, - float* __restrict__ shared_min, - float* __restrict__ shared_max) -{ - /* NULL guard — nothing wired yet, output stays at 0 default. */ - if (src == NULL) { - if (threadIdx.x == 0) { - out_min_mean_max[buf_idx * STATS_PER_BUFFER + 0] = 0.0f; - out_min_mean_max[buf_idx * STATS_PER_BUFFER + 1] = 0.0f; - out_min_mean_max[buf_idx * STATS_PER_BUFFER + 2] = 0.0f; - } - return; - } - - /* Strided fold: each thread reduces a slice into private accumulators. */ - float t_sum = 0.0f; - float t_min = FLT_MAX; - float t_max = -FLT_MAX; - for (int i = threadIdx.x; i < n_samples; i += blockDim.x) { - float v = src[i]; - t_sum += v; - t_min = fminf(t_min, v); - t_max = fmaxf(t_max, v); - } - shared_sum[threadIdx.x] = t_sum; - shared_min[threadIdx.x] = t_min; - shared_max[threadIdx.x] = t_max; - __syncthreads(); - - /* Standard log2(BLOCK_SIZE) tree reduction — no atomicAdd. The three - * reductions run in lockstep (same `s` divisor, same `__syncthreads`) - * so we don't pay extra barriers per reduction. */ - for (int s = blockDim.x / 2; s > 0; s >>= 1) { - if (threadIdx.x < s) { - shared_sum[threadIdx.x] += shared_sum[threadIdx.x + s]; - shared_min[threadIdx.x] = fminf(shared_min[threadIdx.x], - shared_min[threadIdx.x + s]); - shared_max[threadIdx.x] = fmaxf(shared_max[threadIdx.x], - shared_max[threadIdx.x + s]); - } - __syncthreads(); - } - - if (threadIdx.x == 0) { - float n_f = fmaxf(1.0f, (float)n_samples); - out_min_mean_max[buf_idx * STATS_PER_BUFFER + 0] = shared_min[0]; - out_min_mean_max[buf_idx * STATS_PER_BUFFER + 1] = shared_sum[0] / n_f; - out_min_mean_max[buf_idx * STATS_PER_BUFFER + 2] = shared_max[0]; - } - /* Pre-next-buffer barrier: subsequent calls reuse shared_sum/min/max, - * so all threads must finish observing this buffer's writes before the - * next buffer's strided fold begins. */ - __syncthreads(); -} - -/// Block-tree-reduce min/mean/max for the 12 reward-chain diagnostic -/// buffers. Single block; the launcher dispatches grid_dim=(1,1,1), -/// block_dim=(BLOCK_SIZE=256,1,1), shared_mem_bytes=3*BLOCK_SIZE*sizeof(float). -/// -/// `out_min_mean_max` must point to 36 mapped-pinned f32 slots -/// (3 stats × 12 buffers). The kernel writes via `__threadfence_system()` -/// at the end so the host sees the updates after stream sync. -extern "C" __global__ void reward_chain_diag_reduce_kernel( - const float* __restrict__ buf_r_popart, /* [n_samples] */ - const float* __restrict__ buf_r_trail, /* [n_samples] */ - const float* __restrict__ buf_r_micro, /* [n_samples] */ - const float* __restrict__ buf_r_opp_cost, /* [n_samples] */ - const float* __restrict__ buf_r_bonus, /* [n_samples] */ - const float* __restrict__ buf_r_weighted, /* [n_samples] */ - const float* __restrict__ buf_post_dd, /* [n_samples] */ - const float* __restrict__ buf_post_inv, /* [n_samples] */ - const float* __restrict__ buf_post_churn, /* [n_samples] */ - const float* __restrict__ buf_post_conv, /* [n_samples] */ - const float* __restrict__ buf_position_abs, /* [n_samples] */ - const float* __restrict__ buf_conviction, /* [n_samples] */ - float* __restrict__ out_min_mean_max, /* [36] mapped-pinned */ - int n_samples) -{ - if (blockIdx.x != 0) return; - - /* Three shared-memory buffers laid out contiguously in dynamic shmem. - * Total = 3 × BLOCK_SIZE × sizeof(float) = 3 KiB at BLOCK_SIZE=256. */ - extern __shared__ float sdata[]; - float* shared_sum = sdata; - float* shared_min = sdata + BLOCK_SIZE; - float* shared_max = sdata + 2 * BLOCK_SIZE; - - reduce_one_buffer(buf_r_popart, n_samples, 0, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_r_trail, n_samples, 1, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_r_micro, n_samples, 2, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_r_opp_cost, n_samples, 3, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_r_bonus, n_samples, 4, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_r_weighted, n_samples, 5, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_post_dd, n_samples, 6, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_post_inv, n_samples, 7, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_post_churn, n_samples, 8, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_post_conv, n_samples, 9, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_position_abs, n_samples, 10, out_min_mean_max, - shared_sum, shared_min, shared_max); - reduce_one_buffer(buf_conviction, n_samples, 11, out_min_mean_max, - shared_sum, shared_min, shared_max); - - /* Cross-host fence: ensure mapped-pinned writes are visible to the host - * after stream sync. Matches the popart_component_ema_kernel pattern. */ - if (threadIdx.x == 0) { - __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 4e60d3dd9..a2c3cd174 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -837,27 +837,6 @@ impl DQNTrainer { } } - // SP11 instrumentation (2026-05-04, diagnostic-only): wire - // the experience collector's 12 per-sample reward-chain - // checkpoint buffers into the trainer. Mirror of the popart- - // component wire-up above. The collector's - // `experience_env_step` writes the running reward cascade - // variable (and per-component locals) at every checkpoint; - // the trainer's `launch_sp11_reward_chain_diag_reduce` - // block-tree-reduces (min, mean, max) into the trainer's - // mapped-pinned output, which the host reads at - // HEALTH_DIAG emit time. Per `feedback_wire_everything_up`: - // producer (collector) and consumer (trainer) wire in the - // same commit. Will be removed in the follow-up commit - // once the inflater is identified and properly fixed. - if let Some(ref collector) = self.gpu_experience_collector { - let (diag_dev_ptrs, n_bars) = collector.sp11_reward_chain_diag_dev_ptrs(); - if let Some(ref mut fused_mut) = self.fused_ctx { - fused_mut.trainer_mut() - .set_sp11_reward_chain_diag_bufs(diag_dev_ptrs, 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); @@ -4634,73 +4613,6 @@ impl DQNTrainer { ); } - // SP11 instrumentation (2026-05-04, diagnostic-only): - // reward_chain_diag HEALTH_DIAG line. Smoke `smoke-test-gwfn8` - // on `fd24b5383` showed mean(|reward|) hitting 5054 at F0 ep2 - // despite all known multiplicative modifiers being - // structurally bounded. This diag captures (min, mean, max) - // at every checkpoint of the reward chain so the inflater - // can be localised to a specific stage: - // per-component pre-composition (popart/trail/micro/ - // opp_cost/bonus) — bounds expectation - // r_weighted post-composition, pre-modifier — controller - // weight × component sanity - // post_dd / post_inv / post_churn / post_conv — modifier- - // site running-reward - // position_abs / conviction — sanity multiplicand checks - // - // Launches the GPU reduction kernel, syncs the producer - // stream, then reads the 36 floats from the mapped-pinned - // output. Will be removed in the follow-up commit once the - // inflater is identified and properly fixed per - // pearl_bounded_modifier_outputs_require_structural_activation. - if let Some(ref fused) = self.fused_ctx { - let trainer = fused.trainer(); - if let Err(e) = trainer.launch_sp11_reward_chain_diag_reduce() { - tracing::warn!(error = %e, - "SP11 reward_chain_diag_reduce launch failed"); - } else { - // Sync the producer stream so the mapped-pinned - // output is host-visible. Same pattern as - // `read_isv_signal_at` callers that depend on a - // just-launched producer. - if let Err(e) = trainer.stream().synchronize() { - tracing::warn!(error = %e, - "SP11 reward_chain_diag_reduce stream sync failed"); - } else { - let v = trainer.read_sp11_reward_chain_diag(); - tracing::info!( - "HEALTH_DIAG[{}]: reward_chain_diag \ - popart [min={:.3} mean={:.3} max={:.3}] \ - trail [min={:.3} mean={:.3} max={:.3}] \ - micro [min={:.3} mean={:.3} max={:.3}] \ - opp_cost [min={:.3} mean={:.3} max={:.3}] \ - bonus [min={:.3} mean={:.3} max={:.3}] \ - r_weighted [min={:.3} mean={:.3} max={:.3}] \ - post_dd [min={:.3} mean={:.3} max={:.3}] \ - post_inv [min={:.3} mean={:.3} max={:.3}] \ - post_churn [min={:.3} mean={:.3} max={:.3}] \ - post_conv [min={:.3} mean={:.3} max={:.3}] \ - position_abs [min={:.3} mean={:.3} max={:.3}] \ - conviction [min={:.3} mean={:.3} max={:.3}]", - epoch, - v[0], v[1], v[2], - v[3], v[4], v[5], - v[6], v[7], v[8], - v[9], v[10], v[11], - v[12], v[13], v[14], - v[15], v[16], v[17], - v[18], v[19], v[20], - v[21], v[22], v[23], - v[24], v[25], v[26], - v[27], v[28], v[29], - v[30], v[31], v[32], - v[33], v[34], v[35], - ); - } - } - } - // Plan 4 Task 6 Commit B: aux-heads diagnostic line. Reads // ISV[AUX_NEXT_BAR_MSE_EMA_INDEX=113] + // ISV[AUX_REGIME_CE_EMA_INDEX=114] (populated by diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index bdca78332..01a8a1007 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -5573,121 +5573,6 @@ validation, scoped against fold-internal exploration/exploitation dynamics (Kelly engagement at fold mid-epoch, NoisyNet σ schedule, saboteur intensity drift), not against the controller. -## SP11 reward-chain diagnostic instrumentation (2026-05-04) - -**Status:** transient diagnostic, scheduled for removal in the -follow-up commit that fixes the inflater root-cause. - -Smoke `smoke-test-gwfn8` on `fd24b5383` showed `reward_split popart` -(= EMA of `|reward_components_per_sample[+0]|` which is post-modifier -total reward) jumping 101 → 5054 → 4174 → 2152 across F0 epochs 1-4. -Sharpe degrades within fold despite the SP11 controller working -correctly. - -All known multiplicative modifiers are structurally bounded: -- `r_popart = capped_pnl ∈ [-10, +10]` (experience_kernels.cu:2799, - bilateral as of the SP11 symmetric-cap fix — see "Resolution - (2026-05-04)" below). Pre-fix: `fminf(base_reward, 10.0f)` at - line 2788 only capped the upper tail; losses were unbounded — - that WAS the inflater this diagnostic was searching for. -- `× conviction ∈ (0, 1)` via sigmoid at experience_kernels.cu:7579 -- `× cf_flip ∈ {-1, +1}` -- `× shaping_scale ∈ [0, 1]` -- `compute_drawdown_penalty ∈ [-5*w_dd, 0]` - -Yet `|reward| ≈ 5054` — something in the chain produces values ~500× -larger than the bounded inputs would predict. To localise the -inflater, this commit adds GPU-side instrumentation (data-gathering -only, no fix) capturing the running `reward` cascade variable + per- -component locals at every checkpoint. - -### Resolution (2026-05-04) - -The diagnostic instrumentation in `smoke-test-k9drh @ 774d7552a` -captured the asymmetry directly: `ep1 r_popart min=-9186, max=+10` -and `ep2 min=-79089, max=+10` (max stuck, min growing each epoch). -Root cause: line 2788 `float capped_pnl = fminf(base_reward, 10.0f)` -— unilateral upper cap; `base_reward = 2.0f * vol_normalized_return` -where `vol_normalized_return = segment_return / vol_norm` has no -structural lower bound (signed P&L). A single large adverse -`segment_return` produced an arbitrarily negative `capped_pnl` → -`r_popart` → `r_weighted` → `reward_components[+0]` → slot 63 -(PopArt input EMA), inflating C51/IQN/Bellman normalization scale -and breaking Q-target consistency across epochs. - -Fix: -```c -float capped_pnl = fmaxf(-10.0f, fminf(base_reward, 10.0f)); -``` -Spec semantic was always bilateral `[-10, +10]`; the implementation -was unilateral. Per -`pearl_bounded_modifier_outputs_require_structural_activation`: -spec-bounded values require BILATERAL structural enforcement. - -Audit follow-through (per `feedback_no_partial_refactor`): all -other `fminf(.,K)` and `fmaxf(.,K)` clamps in -`experience_kernels.cu` reviewed for the same asymmetric-clamp -pattern. Findings: -- Reward modifier chain (3260-3500): all bounded modifiers operate - on the post-cap `r_weighted`; once `capped_pnl` is bilateral, the - whole downstream chain is bounded. -- Bilateral clamps verified correct at: 515-517 (cost bands), - 2174/2191/2227 (conviction/Kelly), 3300 (stability), - 3729/3763 (cf_reward), 3452 (capital-floor), 5076/5210/5213/ - 5461/5535 ([0,1] gates), 6353 (d_h0), 6693/6694 (ADX/CUSUM). -- Intentional asymmetries verified at: 1078 (cosine_progress ≥ 0 - structurally), 1082-1085 (eps×mult ≥ 0 structurally + EPS_FLOOR - at 1090-1094), 2564-2565 (max_dd lower-unbounded by drawdown - semantic), 2873/2949 (fabsf-based unit ≥ 0 structurally), 4574 - (guarded by `if best_pnl > 0.0f`), 5896 (input-gate underflow - toward 0 is correct sigmoid-like semantic). -- Latent finding flagged separately (NOT fixed in this commit — - feature-side, requires consumer audit): `plan_isv[PNL_VS_TARGET]` - at 850 and `plan_isv[PNL_VS_STOP]` at 853 are upper-clamped at - 2.0 but lower-unbounded; mirrored in `backtest_plan_kernel.cu`. - These feed `assemble_state` (line 956) as policy features, not - reward components, so out of scope for this reward-chain fix. - -Diagnostic instrumentation (commit 774d7552a) NOT removed in this -commit — kept for the symmetric-cap validation smoke on L40S. -Removal lands in the follow-up commit after smoke validates the -fix end-to-end. - -### What it adds - -| Component | Detail | -|-----------|--------| -| Per-sample buffers (12) | `diag_r_popart`, `diag_r_trail`, `diag_r_micro`, `diag_r_opp_cost`, `diag_r_bonus`, `diag_r_weighted`, `diag_post_dd`, `diag_post_inv`, `diag_post_churn`, `diag_post_conv`, `diag_position_abs`, `diag_conviction` — each `[alloc_episodes × alloc_timesteps]` f32, owned by `gpu_experience_collector` | -| Producer | `experience_env_step` writes each buffer at the matching checkpoint with NULL-tolerant guards; entry-block default 0 at every (i,t) before any early-return | -| Reduction kernel | `reward_chain_diag_reduce_kernel.cu` — single block, 256 threads, three lockstep block-tree-reduces (sum/min/max) per buffer, no atomicAdd; outputs `[36]` f32 (3 stats × 12 buffers) to `out_min_mean_max` | -| Trainer wiring | `set_sp11_reward_chain_diag_bufs` + `launch_sp11_reward_chain_diag_reduce` + `read_sp11_reward_chain_diag` (mapped-pinned 36-f32 output) | -| Host emit | `HEALTH_DIAG[{epoch}]: reward_chain_diag …` immediately after the existing `reward_split` line (training_loop.rs ~4614) — launches reduction kernel, syncs stream, reads mapped-pinned, emits all 12 (min, mean, max) triples | - -### What it does NOT add - -- No new ISV slots (output goes to a dedicated mapped-pinned buffer) -- No state-reset registry entries (buffers reset to 0 every step via - the kernel entry-block default writes — same pattern as the other - per-sample diagnostic buffers like `trail_triggered_per_sample`) -- No reward-composition logic changes — the cascade is *captured*, - not modified. This is INSTRUMENTATION, not a fix. - -### Removal plan - -After the inflater is identified from a fresh smoke run reading the -new HEALTH_DIAG line, the follow-up commit will: -1. Apply the proper structural fix (most likely a missing bound on - one of the bonus/penalty multiplicands per - `pearl_one_unbounded_signal_per_reward.md` + - `pearl_bounded_modifier_outputs_require_structural_activation.md`). -2. Delete the 12 buffers, the reduction kernel, the launcher/setter/ - reader trio, the kernel parameter additions, the per-sample - writes, the build.rs entry, the HEALTH_DIAG line, and this audit - section. - -Per `feedback_no_partial_refactor`: instrumentation lands as one -atomic commit; removal lands as one atomic commit (with the fix). - ## SP11 plan_isv symmetric clamp — 4 mirror sites (2026-05-04) The implementer of `35db31089` (symmetric reward cap) flagged 4