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

diag(sp11): instrument reward chain to find the 5000x inflater

Browse Source 
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<f32> 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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-04 13:55:03 +02:00
parent 7a19c51522
commit 774d7552a0
7 changed files with 721 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
crates/ml/build.rs
View File

@@ -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

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

@@ -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.

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

@@ -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::<f32>() 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

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

@@ -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<f32>, // [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<f32>, // [alloc_episodes * alloc_timesteps]
diag_r_trail_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_r_micro_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_r_opp_cost_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_r_bonus_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_r_weighted_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_post_dd_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_post_inv_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_post_churn_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_post_conv_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_position_abs_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
diag_conviction_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
/// C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2).
/// [alloc_episodes * alloc_timesteps * 6] row-major.
/// Components: [0]=popart(total), [1]=cf, [2]=trail(0), [3]=micro, [4]=opp_cost, [5]=bonus.
@@ -1225,6 +1248,50 @@ impl GpuExperienceCollector {
let popart_component_per_sample = stream
.alloc_zeros::<f32>(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::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_r_popart_per_sample: {e}")))?;
let diag_r_trail_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_r_trail_per_sample: {e}")))?;
let diag_r_micro_per_sample = stream
.alloc_zeros::<f32>(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::<f32>(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::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_r_bonus_per_sample: {e}")))?;
let diag_r_weighted_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_r_weighted_per_sample: {e}")))?;
let diag_post_dd_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_post_dd_per_sample: {e}")))?;
let diag_post_inv_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_post_inv_per_sample: {e}")))?;
let diag_post_churn_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_post_churn_per_sample: {e}")))?;
let diag_post_conv_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_post_conv_per_sample: {e}")))?;
let diag_position_abs_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc diag_position_abs_per_sample: {e}")))?;
let diag_conviction_per_sample = stream
.alloc_zeros::<f32>(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}"

170
crates/ml/src/cuda_pipeline/reward_chain_diag_reduce_kernel.cu Normal file
View File

@@ -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 <cuda_runtime.h>
#include <cfloat>
#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();
}
}

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

@@ -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

59
docs/dqn-wire-up-audit.md
View File

@@ -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).