From 774d7552a08e72f1415a5ee869f799290ee4c821 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 4 May 2026 13:55:03 +0200 Subject: [PATCH] diag(sp11): instrument reward chain to find the 5000x inflater MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Smoke smoke-test-gwfn8 on fd24b5383 showed mean(|reward|) hitting 5054 at F0 ep2 despite all known multiplicative modifiers being structurally bounded (conviction in (0,1) via sigmoid at line 7579, cf_flip in +/-1, drawdown in [-5*w_dd,0], shaping_scale in [0,1]). The inflater is somewhere in the pre-composition or modifier chain that isn't currently visible in HEALTH_DIAG. Adds 12 per-sample diagnostic buffers + reduction kernel + HEALTH_DIAG emit for min/mean/max at every checkpoint in the reward chain: - per-component (r_popart, r_trail, r_micro, r_opp_cost, r_bonus) - r_weighted (post-composition, pre-modifier) - post-modifier sequential (post_dd, post_inv, post_churn, post_conv) - sanity checks (position_abs, conviction) Implementation: - 12 new per-sample CudaSlice buffers in gpu_experience_collector (alloc_episodes * alloc_timesteps each); zero-init at every (i,t) in the kernel entry block before any early-return; written at each checkpoint with NULL-tolerant guards. - new reward_chain_diag_reduce_kernel.cu: single-block 256-thread block-tree-reduce over the 12 buffers, three lockstep reductions (sum/min/max) per buffer in shared memory; outputs 36 floats (3 stats x 12 buffers) to a 36-slot mapped-pinned scratch buffer on the trainer; no atomicAdd per feedback_no_atomicadd, pure GPU compute per feedback_no_cpu_compute_strict, mapped-pinned host visibility per feedback_no_htod_htoh_only_mapped_pinned. - trainer: cubin load, mapped-pinned 36-f32 output, set_sp11_reward_ chain_diag_bufs setter, launch_sp11_reward_chain_diag_reduce launcher, read_sp11_reward_chain_diag host accessor. - training_loop.rs: wires the 12 collector buffers post-construction (mirror of the popart-component wire-up); HEALTH_DIAG `reward_ chain_diag` emit added immediately after `reward_split` — launches reduction kernel, syncs stream, reads the 36 floats. - build.rs: adds reward_chain_diag_reduce_kernel.cu to the cubin manifest. - docs/dqn-wire-up-audit.md: new section documenting the instrumentation scope, additions, exclusions (no ISV slots, no state-reset registry entries), and removal plan. No state-reset registry entry: this is a transient diagnostic, not persistent state — 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 ISV slots are added: the host reads the mapped-pinned scratch directly to keep this lightweight and avoid permanent ISV growth. Will be removed in a follow-up commit once the inflater is identified and properly fixed per pearl_bounded_modifier_outputs_require_ structural_activation. 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 | 59 ++++++ 7 files changed, 721 insertions(+), 1 deletion(-) create 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 1677f1b3b..eb9c4bece 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -573,6 +573,14 @@ 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 36023b6c0..97a645b99 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1830,7 +1830,32 @@ 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 + 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) */ ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -1927,6 +1952,25 @@ 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; @@ -3301,6 +3345,21 @@ 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 @@ -3332,6 +3391,14 @@ 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) ────────────────────── @@ -3341,6 +3408,10 @@ 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; @@ -3450,6 +3521,10 @@ 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. @@ -3459,17 +3534,36 @@ 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 12aec6739..2ca709323 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -577,6 +577,16 @@ 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 @@ -4926,6 +4936,28 @@ 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). @@ -12655,6 +12687,43 @@ 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. /// @@ -13044,6 +13113,73 @@ 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. /// @@ -16504,6 +16640,25 @@ 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 @@ -19913,6 +20068,18 @@ 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 69e31c793..67b1acfb5 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -628,6 +628,29 @@ 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. @@ -1225,6 +1248,50 @@ 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). @@ -1725,6 +1792,21 @@ 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, @@ -2170,6 +2252,38 @@ 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). @@ -4108,6 +4222,26 @@ 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 new file mode 100644 index 000000000..909a4cac0 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/reward_chain_diag_reduce_kernel.cu @@ -0,0 +1,170 @@ +// 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 a2c3cd174..4e60d3dd9 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -837,6 +837,27 @@ 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); @@ -4613,6 +4634,73 @@ 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 d5278ac35..3085a6c36 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -5572,3 +5572,62 @@ degradation will be triaged as a follow-up task post-T10 50-epoch 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` (experience_kernels.cu:2744) +- `× 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. + +### 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).