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

plan(sp18 v2): Phase 0 Task 0.1 — D-leg per-action reward decomposition diagnostic

Browse Source 
Adds the SP18 v2 Phase 0 D-leg observability scaffold per the plan's
"Phase 0 — diagnostic emit (NO functional change)" task:

- New kernel `reward_decomp_diag_kernel.cu`: block tree-reduce
  (4 blocks × 256 threads, one block per direction-axis bin) reading
  `reward_components_per_sample [N×6]` + `actions_out [N]` and emitting
  5 per-bin stats (mean r_micro / mean r_opp_cost / mean r_popart /
  mean |reward| / fire_rate) into a 20-float row-major output. Bin
  order: Hold(0)→Long(1)→Short(2)→Flat(3); col order: micro→opp→
  popart→abs→fire. Empty-bin guard emits 0.0 (NOT NaN) per the
  consumer-side KILL CRITERION arithmetic contract.

- Cubin manifest entry in `crates/ml/build.rs` + `REWARD_DECOMP_DIAG_
  CUBIN` re-export in `gpu_dqn_trainer.rs`.

- 20-float `MappedF32Buffer sp18_reward_decomp_diag_buf` field on
  `GpuDqnTrainer` + accessor pair (`sp18_reward_decomp_diag_dev_ptr`
  for the writer-side launcher; `read_sp18_reward_decomp_diag` for the
  HEALTH_DIAG reader). Buffer is constructor-zeroed so cold-start
  HEALTH_DIAG emits a deterministic zero block.

- `sp18_reward_decomp_diag_kernel` field on `GpuExperienceCollector` +
  cubin load on the collector's stream + `launch_sp18_reward_decomp_
  diag(n, b1, b2, b3, out_dev_ptr)` launcher. Wired in
  `training_loop.rs` at the per-step boundary, BEFORE
  `launch_reward_component_ema_inplace` (which `memset_zeros` the
  source buffer after consuming it) per `pearl_canary_input_freshness_
  launch_order`.

- New per-epoch HEALTH_DIAG line emit at the existing per-epoch
  boundary (after the SP17 dueling line):

    HEALTH_DIAG[N]: reward_decomp [hold(micro=X opp=Y popart=Z abs=W
                    fire=F) long(...) short(...) flat(...)]

  Reads the mapped-pinned 20-float diag buffer directly via the
  collector→trainer host_ptr — no DtoH copy.

- New `crates/ml/tests/sp18_hold_reward_oracle_tests.rs`:
  * `reward_decomp_per_action_cpu_oracle` (CPU oracle pinning the
    per-bin reduction math against a 4-sample synthetic batch).
  * `reward_decomp_per_action_gpu_oracle` (GPU oracle, ignored unless
    `--ignored`; asserts kernel matches CPU oracle bit-for-bit within
    1e-6 f32 budget).
  * `reward_decomp_empty_bin_emits_zero_not_nan` (empty-bin contract
    guard).

Pure observability — no production-path consumer in this commit. No
reward changes, no Bellman target changes, no kernel modifications to
the action-selection or training paths. Per
`feedback_no_partial_refactor` the kernel + cubin manifest + buffer +
launcher + production wire-up + HEALTH_DIAG emit + GPU oracle test all
land atomically.

Verification:
  SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean.
  CPU oracle test passes; GPU oracle + empty-bin guard both pass on
  RTX 3050 Ti (2.13s).

Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
      § Phase 0 Task 0.1.
Audit: docs/dqn-wire-up-audit.md § "SP18 v2 Phase 0 Task 0.1".

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 01:06:48 +02:00
parent 52100ae7a1
commit 15b50ac38f
7 changed files with 921 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
crates/ml/build.rs
View File

@@ -879,6 +879,21 @@ fn main() {
// `isv[ISV_HOLD_COST_IDX=380]` by `isv[461]`. See
// docs/dqn-wire-up-audit.md § "SP16 Phase 2" for rationale.
"hold_cost_scale_update_kernel.cu",
// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
// decomposition diagnostic. Reads `reward_components_per_sample`
// ([n_samples × 6]) + `actions` ([n_samples], factored-action
// ints) from the experience collector, bins per direction
// (DIR_HOLD / DIR_LONG / DIR_SHORT / DIR_FLAT — 4 bins), and
// emits 5 stats per bin (mean r_micro, mean r_opp_cost, mean
// r_popart, mean |reward|, fire_rate) into a 20-float
// mapped-pinned `reward_decomp_diag_buf`. Block tree-reduce
// (4 blocks × 256 threads, one block per bin); no atomicAdd
// per `feedback_no_atomicadd.md`. Pure observability — does
// NOT modify any production-path buffer; written values are
// consumed exclusively by the per-epoch HEALTH_DIAG line emit.
// Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
// § Phase 0 Task 0.1.
"reward_decomp_diag_kernel.cu",
// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
// Single-block 256-thread kernel computing RMS = sqrt(mean(x²))
// over `h_s2_aux [B, SH2]` (aux trunk final output, no activation)

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

@@ -114,6 +114,15 @@ static ATOMS_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/ato
static Q_QUANTILE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/q_quantile_kernel.cubin"));
/// Plan 3 Task 1 C.2: per-component |reward| EMA into ISV[63..69).
pub(crate) static REWARD_COMPONENT_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_component_ema_kernel.cubin"));
/// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
/// decomposition diagnostic. Reads `reward_components_per_sample`
/// + `actions_out` from the collector, bins per-direction
/// (DIR_HOLD / DIR_LONG / DIR_SHORT / DIR_FLAT), emits 5 stats per bin
/// (mean r_micro / r_opp_cost / r_popart / |reward| / fire_rate) into
/// a 20-float mapped-pinned `reward_decomp_diag_buf`. Block tree-
/// reduce — 4 blocks × 256 threads. Pure observability — no
/// production-path consumer in this commit. Plan: SP18 Phase 0 Task 0.1.
pub(crate) static REWARD_DECOMP_DIAG_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_decomp_diag_kernel.cubin"));
/// Plan 3 Task 3 B.2: Flat→Positioned transition rate EMA into ISV[71].
pub(crate) static TRADE_RATE_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/trade_rate_ema_kernel.cubin"));
/// Plan 3 Task 4 B.4: per-batch readiness EMA + derived plan_threshold.
@@ -6803,6 +6812,25 @@ pub struct GpuDqnTrainer {
/// host writes — kernel-owned). Lifetime matches the trainer; reset
/// once at construct, written-then-read inside the kernel each launch.
sp17_advantage_clip_scratch: super::mapped_pinned::MappedF32Buffer,
// ── SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
// decomposition diagnostic buffer ──────────────────────────────────
/// 20-float mapped-pinned diagnostic buffer for the SP18 D-leg
/// per-action reward decomposition. Layout (row-major, 4×5):
///
/// row 0 = Hold, row 1 = Long, row 2 = Short, row 3 = Flat
/// col 0 = mean(r_micro), col 1 = mean(r_opp_cost),
/// col 2 = mean(r_popart), col 3 = mean(|reward|),
/// col 4 = fire_rate
///
/// Written by the collector's
/// `launch_sp18_reward_decomp_diag(...)` kernel each per-step
/// boundary, BEFORE `launch_reward_component_ema_inplace` zeros the
/// source buffer. Consumed exclusively by the per-epoch HEALTH_DIAG
/// emit at `training_loop.rs`. Pure observability — no production-
/// path consumer in this commit. Plan: SP18 Phase 0 Task 0.1.
pub(crate) sp18_reward_decomp_diag_buf: super::mapped_pinned::MappedF32Buffer,
// SP14 Layer C Phase C.1 (2026-05-08): α-machinery struct fields deleted
// atomically with the kernel files (alpha_grad_compute_kernel,
// gradient_hack_detect_kernel, sp14_scale_wire_col_kernel) per
@@ -9407,6 +9435,48 @@ impl GpuDqnTrainer {
Ok(self.read_isv_signal_at(ADVANTAGE_CLIP_BOUND_INDEX))
}
/// SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
/// decomposition diagnostic device pointer.
///
/// Returns the mapped-pinned device pointer of the 20-float
/// `sp18_reward_decomp_diag_buf` so the collector's
/// `launch_sp18_reward_decomp_diag` can write into it. The host-side
/// HEALTH_DIAG emit reads via `read_sp18_reward_decomp_diag()` after
/// stream sync. Pure observability — caller is expected to be the
/// per-step training loop wiring point.
pub fn sp18_reward_decomp_diag_dev_ptr(&self) -> u64 {
self.sp18_reward_decomp_diag_buf.dev_ptr
}
/// SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
/// decomposition diagnostic readback.
///
/// Returns the 20-float diagnostic block as a `[f32; 20]`, row-major
/// 4 bins × 5 stats:
///
/// [Hold_micro, Hold_opp, Hold_popart, Hold_abs, Hold_fire,
/// Long_micro, Long_opp, Long_popart, Long_abs, Long_fire,
/// Short_micro, Short_opp, Short_popart, Short_abs, Short_fire,
/// Flat_micro, Flat_opp, Flat_popart, Flat_abs, Flat_fire]
///
/// Reads the mapped-pinned host_ptr directly — caller MUST have
/// stream-synced after the collector's launch (e.g. via the natural
/// per-step boundary that already runs `cuStreamSynchronize` between
/// collection and training). No HtoD/DtoH copy.
///
/// On cold-start (before any kernel launch) returns the
/// constructor-initialised zero block.
pub fn read_sp18_reward_decomp_diag(&self) -> [f32; 20] {
let v = self.sp18_reward_decomp_diag_buf.read_all();
debug_assert_eq!(
v.len(), 20,
"sp18_reward_decomp_diag_buf size invariant: 4 bins × 5 stats"
);
let mut out = [0.0_f32; 20];
for (i, x) in v.iter().take(20).enumerate() { out[i] = *x; }
out
}
/// Update per-branch liquid tau modulation — GPU kernel (RK4/Euler adaptive ODE).
///
/// Reads per_branch_q_gaps (pinned device-mapped) and qlstm_context from device.
@@ -20568,6 +20638,27 @@ impl GpuDqnTrainer {
"sp17_advantage_clip_scratch alloc ({sp17_clip_scratch_len} f32): {e}"
)))?;
// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
// decomposition diagnostic buffer. 4 bins × 5 stats = 20 floats
// mapped-pinned. The collector's
// `launch_sp18_reward_decomp_diag` writes per-step; the host-side
// HEALTH_DIAG emit reads via the mapped host_ptr after stream
// sync. Pure observability — no production-path consumer.
const SP18_REWARD_DECOMP_DIAG_LEN: usize = 4 * 5;
let sp18_reward_decomp_diag_buf = unsafe {
super::mapped_pinned::MappedF32Buffer::new(SP18_REWARD_DECOMP_DIAG_LEN)
}
.map_err(|e| MLError::ModelError(format!(
"sp18_reward_decomp_diag_buf alloc ({SP18_REWARD_DECOMP_DIAG_LEN} f32): {e}"
)))?;
// Initialise to 0.0 so the first HEALTH_DIAG read on cold-start
// (before the collector has fired) is a deterministic zero block
// rather than indeterminate page contents. Pearl-A first-
// observation bootstrap fires inside the kernel itself when
// n_samples > 0; this just zero-fills before the kernel ever runs.
sp18_reward_decomp_diag_buf
.write_from_slice(&vec![0.0_f32; SP18_REWARD_DECOMP_DIAG_LEN]);
// SP14 Layer C Phase C.1 (2026-05-08): α-machinery cubin loads
// (alpha_grad_compute, gradient_hack_detect, sp14_scale_wire_col)
// deleted atomically with the kernel files. Coupling A
@@ -24810,6 +24901,11 @@ impl GpuDqnTrainer {
sp17_v_share_kernel,
sp17_advantage_clip_bound_kernel,
sp17_advantage_clip_scratch,
// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action
// reward decomposition diagnostic buffer. Mapped-pinned 20
// floats (4 bins × 5 stats). Written by collector kernel,
// read host-side at HEALTH_DIAG cadence.
sp18_reward_decomp_diag_buf,
sp14_dir_concat_qaux_kernel,
sp14_dir_qaux_concat_scratch,
// SP14 backward dX scratch — column-SH2 wire is dropped at

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

@@ -888,6 +888,19 @@ pub struct GpuExperienceCollector {
/// Launched after each collect_experiences_gpu epoch from launch_reward_component_ema_inplace().
reward_component_ema_kernel: CudaFunction,
/// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
/// decomposition diagnostic kernel. Reads
/// `reward_components_per_sample [n_samples × 6]` + `actions_out
/// [n_samples]` (factored-action ints) and bins per-direction
/// (DIR_HOLD / DIR_LONG / DIR_SHORT / DIR_FLAT). Writes a 20-float
/// per-action-bin diagnostic block (4 bins × 5 stats per bin) into
/// a host-supplied mapped-pinned `out_buf`. Block tree-reduce —
/// 4 blocks × 256 threads. Launched BEFORE
/// `launch_reward_component_ema_inplace` (which zeroes the source
/// buffer). Pure observability — no production-path consumer.
/// Plan: SP18 Phase 0 Task 0.1.
sp18_reward_decomp_diag_kernel: CudaFunction,
/// SP4 GPU-only Pearls A+D applicator (2026-05-01). Loaded from the
/// shared `APPLY_PEARLS_AD_CUBIN`; consumed by
/// `launch_reward_component_ema_inplace` after the
@@ -1964,6 +1977,24 @@ impl GpuExperienceCollector {
.map_err(|e| MLError::ModelError(format!("reward_component_ema load: {e}")))?
};
// SP18 v2 Phase 0 Task 0.1 (2026-05-08): per-action reward
// decomposition diagnostic kernel. Loaded on the collector's
// stream so the launcher chains directly after
// `experience_action_select` (which writes `actions_out` packed
// factored action_idx) and BEFORE
// `launch_reward_component_ema_inplace` (which zeros the
// reward-component buffer).
let sp18_reward_decomp_diag_kernel = {
use super::gpu_dqn_trainer::REWARD_DECOMP_DIAG_CUBIN;
let m = stream.context()
.load_cubin(REWARD_DECOMP_DIAG_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"sp18 reward_decomp_diag cubin load: {e}")))?;
m.load_function("reward_decomp_diag")
.map_err(|e| MLError::ModelError(format!(
"sp18 reward_decomp_diag load: {e}")))?
};
// SP4 GPU-only Pearls A+D applicator (2026-05-01). Loaded
// separately on this collector's stream so the cross-boundary
// `launch_reward_component_ema_inplace` can apply Pearls A+D
@@ -2583,6 +2614,7 @@ impl GpuExperienceCollector {
market_dim_bn: market_dim_cfg,
bn_tanh_concat_fn,
reward_component_ema_kernel,
sp18_reward_decomp_diag_kernel,
apply_pearls_ad_kernel,
// SP13 v3 P0a.T4 (2026-05-04): Hold-rate observer + fixed-α
// EMA applicator + 1-elem mapped-pinned scratch — chained
@@ -3322,6 +3354,83 @@ impl GpuExperienceCollector {
Ok([0.0, cf_rate, trail_r, micro_r])
}
/// SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
/// decomposition diagnostic launcher.
///
/// Launches `reward_decomp_diag` (4 blocks × 256 threads, one block
/// per direction-axis bin) on this collector's stream. Reads
/// `reward_components_per_sample [n_samples × 6]` + `actions_out
/// [n_samples]` written by the last `collect_experiences_gpu`
/// epoch and bins per-direction (DIR_HOLD / DIR_LONG / DIR_SHORT /
/// DIR_FLAT). Writes a 20-float diagnostic block (4 bins × 5 stats
/// per bin) into the host-supplied mapped-pinned `out_dev_ptr`.
///
/// Output layout (row-major, 4×5):
///
/// row 0 = Hold, row 1 = Long, row 2 = Short, row 3 = Flat
/// col 0 = mean(r_micro) (component idx 3)
/// col 1 = mean(r_opp_cost) (component idx 4)
/// col 2 = mean(r_popart) (component idx 0)
/// col 3 = mean(|reward|) (sum of |components|)
/// col 4 = fire_rate (frac bars with non-zero comp)
///
/// **Launch ordering** (CRITICAL): MUST run BEFORE
/// `launch_reward_component_ema_inplace` because the latter
/// `memset_zeros` the source `reward_components_per_sample` after
/// consuming it. Per `pearl_canary_input_freshness_launch_order`.
///
/// Pure observability — does NOT modify any production-path buffer.
/// `n_samples` should be `alloc_episodes * alloc_timesteps` (base
/// batch, not CF-doubled). Pass `out_dev_ptr=0` to noop the launch
/// (e.g. before the trainer wires the diagnostic buffer).
///
/// Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
/// § Phase 0 Task 0.1.
pub fn launch_sp18_reward_decomp_diag(
&mut self,
n_samples: usize,
b1: usize,
b2: usize,
b3: usize,
out_dev_ptr: u64,
) -> Result<(), crate::MLError> {
if n_samples == 0 || out_dev_ptr == 0 {
return Ok(());
}
let n_i32 = n_samples as i32;
let b1_i32 = b1 as i32;
let b2_i32 = b2 as i32;
let b3_i32 = b3 as i32;
let rc_ptr = self.reward_components_per_sample.raw_ptr();
let act_ptr = self.actions_out.raw_ptr();
// Block tree-reduce: 4 blocks × 256 threads. Dynamic shmem =
// 5 × BLOCK_SIZE × sizeof(f32) for the 5 reduction tiles
// (micro / opp_cost / popart / abs_sum / count). Matches the
// SP18_DECOMP_BLOCK_SIZE=256 kernel-side constant.
const BLOCK_DIM: u32 = 256;
const NUM_BINS: u32 = 4;
let shmem_bytes: u32 = 5 * BLOCK_DIM * std::mem::size_of::<f32>() as u32;
unsafe {
self.stream.launch_builder(&self.sp18_reward_decomp_diag_kernel)
.arg(&rc_ptr)
.arg(&act_ptr)
.arg(&n_i32)
.arg(&b1_i32)
.arg(&b2_i32)
.arg(&b3_i32)
.arg(&out_dev_ptr)
.launch(LaunchConfig {
grid_dim: (NUM_BINS, 1, 1),
block_dim: (BLOCK_DIM, 1, 1),
shared_mem_bytes: shmem_bytes,
})
.map_err(|e| crate::MLError::ModelError(format!(
"sp18 reward_decomp_diag launch: {e}"
)))?;
}
Ok(())
}
/// C.2 Plan 3 Task 1 (spec §4.C.2); SP4 Layer A Task A13.5 retrofit
/// (2026-05-01); GPU-Pearls refactor (2026-05-01): GPU-driven
/// reward-component step-observation producer + GPU Pearls A+D update.

229
crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu Normal file
View File

@@ -0,0 +1,229 @@
/* ══════════════════════════════════════════════════════════════════════════
* SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
* decomposition diagnostic.
*
* Reads the per-sample reward-component buffer + actions buffer produced
* by `experience_kernels.cu`, bins per direction-axis (Hold / Long /
* Short / Flat), and emits 5 per-bin statistics into a 20-float
* mapped-pinned diagnostic buffer:
*
* col 0: mean(r_micro) per action bin (component idx 3)
* col 1: mean(r_opp_cost) per action bin (component idx 4)
* col 2: mean(r_popart) per action bin (component idx 0)
* col 3: mean(|reward|) per action bin (sum of absolute values
* across all 6 components)
* col 4: fire_rate per action bin (frac of bars with any
* non-zero component)
*
* Buffer layout (row-major, 4 rows × 5 cols):
*
* out[bin * 5 + col]
*
* with bin order:
*
* bin 0 = Hold (DIR_HOLD = 1)
* bin 1 = Long (DIR_LONG = 2)
* bin 2 = Short (DIR_SHORT = 0)
* bin 3 = Flat (DIR_FLAT = 3)
*
* The bin order is logical (Hold first — primary observable for SP18 KILL
* CRITERION) and explicitly does NOT match the DIR_* enum order. Each
* thread reads its sample's `actions[i]`, decodes the direction via
* `action / (b1 * b2 * b3)`, and increments counters for the matching
* bin.
*
* The reward-component layout is the SP15 Wave 2 reshape per
* `reward_component_ema_kernel.cu`:
*
* [c=0] popart — final on-policy reward (PopArt input)
* [c=1] cf — counterfactual replay
* [c=2] trail — 0.0 placeholder
* [c=3] micro — dense OFI-alignment micro-reward
* [c=4] opp_cost — Flat opportunity cost
* [c=5] bonus — trade-attempt / timing bonus
*
* Pearls + invariants:
*
* - `feedback_no_atomicadd.md` — block tree-reduce; one block per bin,
* 5 shmem accumulators each, no atomicAdd.
*
* - `pearl_no_host_branches_in_captured_graph.md` — pure kernel launch,
* no host-side dispatch within the capture region (the call site is
* outside the captured graph anyway: per-step diagnostic at the same
* cadence as `launch_reward_component_ema_inplace`).
*
* - `feedback_no_stubs.md` — full bin-aware reduction, no zero-output
* placeholder.
*
* - `feedback_no_htod_htoh_only_mapped_pinned.md` — `out` is a mapped-
* pinned `MappedF32Buffer.dev_ptr`; host reads after stream sync.
*
* - `pearl_first_observation_bootstrap.md` — the kernel writes raw
* means each launch (no Wiener-α blend at this Phase 0 cadence —
* diagnostic only). Host-side EMA blend is a separate concern; the
* Phase 0 KILL CRITERION reads the raw per-epoch means.
*
* Args:
*
* rc — `reward_components_per_sample[n_samples * 6]` device
* pointer. Row-major; row i is `[popart, cf, trail, micro,
* opp_cost, bonus]`.
*
* actions — `actions_out[n_samples]` device pointer (i32). Each
* int decodes via `dir = action / (b1 * b2 * b3)`.
*
* n_samples — total sample count (production: `alloc_episodes ×
* alloc_timesteps`; the CF-doubled buffer half is NOT
* consumed by this kernel — pass the base count).
*
* b1, b2, b3 — magnitude / order / urgency branch sizes (used to
* decode direction from the factored-action int).
*
* out — `out[20]` mapped-pinned f32, row-major 4×5 (4 bins × 5
* stats per bin). See top-of-file layout comment.
*
* Launch: grid=(4, 1, 1), block=(BLOCK_SIZE, 1, 1).
* Shared mem: 5 × BLOCK_SIZE × sizeof(float) for the 5 reduction tiles.
* ══════════════════════════════════════════════════════════════════════════ */
#include <cuda_runtime.h>
#include "state_layout.cuh"
#define SP18_DECOMP_BLOCK_SIZE 256
/* Component indices in `reward_components_per_sample` row. */
#define RCP_POPART_IDX 0
#define RCP_MICRO_IDX 3
#define RCP_OPP_COST_IDX 4
#define RCP_NUM_COMPS 6
/* Output layout: 4 bins × 5 cols = 20 floats. */
#define SP18_DECOMP_NUM_BINS 4
#define SP18_DECOMP_NUM_COLS 5
#define SP18_DECOMP_OUT_SIZE (SP18_DECOMP_NUM_BINS * SP18_DECOMP_NUM_COLS)
/* Logical bin → DIR_* enum lookup. Order: Hold, Long, Short, Flat. */
__device__ __forceinline__ int sp18_decomp_bin_to_dir(int bin) {
switch (bin) {
case 0: return DIR_HOLD;
case 1: return DIR_LONG;
case 2: return DIR_SHORT;
case 3: return DIR_FLAT;
default: return -1; /* unreachable */
}
}
extern "C" __global__
void reward_decomp_diag(
const float* __restrict__ rc, /* [n_samples × 6] */
const int* __restrict__ actions, /* [n_samples] */
int n_samples,
int b1,
int b2,
int b3,
float* __restrict__ out /* [20] mapped-pinned */
) {
extern __shared__ float s_buf[];
float* s_micro = s_buf; /* [BLOCK_SIZE] */
float* s_opp = s_buf + SP18_DECOMP_BLOCK_SIZE; /* [BLOCK_SIZE] */
float* s_popart = s_buf + 2 * SP18_DECOMP_BLOCK_SIZE; /* [BLOCK_SIZE] */
float* s_abs_sum = s_buf + 3 * SP18_DECOMP_BLOCK_SIZE; /* [BLOCK_SIZE] */
float* s_count = s_buf + 4 * SP18_DECOMP_BLOCK_SIZE; /* [BLOCK_SIZE] */
const int tid = threadIdx.x;
const int bin = blockIdx.x;
if (bin >= SP18_DECOMP_NUM_BINS) return;
const int target_dir = sp18_decomp_bin_to_dir(bin);
const int dir_div = b1 * b2 * b3;
/* Per-thread accumulators — strided over n_samples. */
float l_micro = 0.0f;
float l_opp = 0.0f;
float l_popart = 0.0f;
float l_abs_sum = 0.0f;
float l_count = 0.0f; /* count of bars matching this bin */
float l_fire = 0.0f; /* count of bars with non-zero component */
for (int i = tid; i < n_samples; i += SP18_DECOMP_BLOCK_SIZE) {
const int dir_i = (dir_div > 0) ? (actions[i] / dir_div) : -1;
if (dir_i != target_dir) continue;
const int row = i * RCP_NUM_COMPS;
const float r_popart = rc[row + RCP_POPART_IDX];
const float r_micro = rc[row + RCP_MICRO_IDX];
const float r_opp = rc[row + RCP_OPP_COST_IDX];
/* sum of |components| → mean of |reward| across all 6 components.
* Read all 6 explicitly so any future component addition fails
* loudly via the row-stride contract instead of silently skipping. */
float abs_total = 0.0f;
#pragma unroll
for (int c = 0; c < RCP_NUM_COMPS; ++c) {
abs_total += fabsf(rc[row + c]);
}
l_micro += r_micro;
l_opp += r_opp;
l_popart += r_popart;
l_abs_sum += abs_total;
l_count += 1.0f;
if (abs_total > 0.0f) {
l_fire += 1.0f;
}
}
/* Track count for division. We reuse s_count to reduce both the
* matching-bin count AND the firing-bar count with a two-phase pattern:
* first phase reduces matching count (denominator for means); second
* phase reduces firing count (numerator for fire_rate). */
s_micro[tid] = l_micro;
s_opp[tid] = l_opp;
s_popart[tid] = l_popart;
s_abs_sum[tid] = l_abs_sum;
s_count[tid] = l_count;
__syncthreads();
/* Block tree-reduce. */
for (int s = SP18_DECOMP_BLOCK_SIZE / 2; s > 0; s >>= 1) {
if (tid < s) {
s_micro[tid] += s_micro[tid + s];
s_opp[tid] += s_opp[tid + s];
s_popart[tid] += s_popart[tid + s];
s_abs_sum[tid] += s_abs_sum[tid + s];
s_count[tid] += s_count[tid + s];
}
__syncthreads();
}
const float total_micro = s_micro[0];
const float total_opp = s_opp[0];
const float total_popart = s_popart[0];
const float total_abs_sum = s_abs_sum[0];
const float total_count = s_count[0];
__syncthreads();
/* Phase 2: reduce firing count using the same s_count tile. */
s_count[tid] = l_fire;
__syncthreads();
for (int s = SP18_DECOMP_BLOCK_SIZE / 2; s > 0; s >>= 1) {
if (tid < s) {
s_count[tid] += s_count[tid + s];
}
__syncthreads();
}
const float total_fire = s_count[0];
if (tid == 0) {
/* Means default to 0.0 when no samples landed in the bin (rather
* than NaN — the consumer reads these as means and would mis-
* interpret NaN as a kernel failure). fire_rate is bounded
* [0, 1] by construction. */
const float denom = (total_count > 0.0f) ? total_count : 1.0f;
const int o = bin * SP18_DECOMP_NUM_COLS;
out[o + 0] = (total_count > 0.0f) ? (total_micro / denom) : 0.0f;
out[o + 1] = (total_count > 0.0f) ? (total_opp / denom) : 0.0f;
out[o + 2] = (total_count > 0.0f) ? (total_popart / denom) : 0.0f;
out[o + 3] = (total_count > 0.0f) ? (total_abs_sum / denom) : 0.0f;
out[o + 4] = (total_count > 0.0f) ? (total_fire / denom) : 0.0f;
__threadfence_system(); /* PCIe-visible to mapped-pinned host_ptr */
}
}

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

@@ -4073,6 +4073,44 @@ impl DQNTrainer {
.as_ref()
.map(|c| c.alloc_episodes() * c.alloc_timesteps())
.unwrap_or(0);
// SP18 v2 Phase 0 Task 0.1 (2026-05-08):
// D-leg per-action reward decomposition
// diagnostic. Runs BEFORE
// `launch_reward_component_ema_inplace`
// (which zeros the source `reward_
// components_per_sample` buffer after
// consuming it). Per
// `pearl_canary_input_freshness_launch_order`.
// Pure observability — writes a 20-float
// mapped-pinned buffer on `fused.trainer
// ()` consumed only by the per-epoch
// HEALTH_DIAG emit; no production-path
// consumer in this commit. ISV/dev_ptr
// is 0 unless the trainer's diag buffer
// was constructed (always true post-PP.2;
// the launcher's `out_dev_ptr=0` guard
// noops if the wire-up regresses).
if let Some(ref fused) = self.fused_ctx {
if let Some(ref mut c) =
self.gpu_experience_collector
{
let cfg = fused.trainer().config();
let b1 = cfg.branch_1_size;
let b2 = cfg.branch_2_size;
let b3 = cfg.branch_3_size;
let out_ptr = fused.trainer()
.sp18_reward_decomp_diag_dev_ptr();
if let Err(e) = c.launch_sp18_reward_decomp_diag(
n_main, b1, b2, b3, out_ptr,
) {
tracing::warn!(
"SP18 Phase 0 reward_decomp_diag \
launch failed (non-fatal, \
diagnostic-only): {e}"
);
}
}
}
if let Some(ref mut c) = self.gpu_experience_collector {
if let Err(e) =
c.launch_reward_component_ema_inplace(n_main)
@@ -5343,6 +5381,51 @@ impl DQNTrainer {
);
}
// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action
// reward decomposition diagnostic emit. Reads the 20-float
// mapped-pinned `sp18_reward_decomp_diag_buf` (4 bins × 5
// stats: micro / opp / popart / abs / fire). Bin order:
// Hold / Long / Short / Flat. The buffer is written by the
// collector kernel `launch_sp18_reward_decomp_diag` (per-step
// boundary, BEFORE `launch_reward_component_ema_inplace`
// zeroes the source buffer) and is host-readable directly via
// its mapped-pinned host_ptr — no DtoH copy. The KILL
// CRITERION at Phase 0 Task 0.4 reads:
//
// ABORT D-leg if |micro_hold| > 2 × |micro_long| AND Hold%
// trajectory is FALLING.
// PROCEED if |micro_hold| < 0.5 × max(|popart_long|,
// |popart_short|)
// AND Hold% trajectory is RISING.
//
// Host-side cold-start (pre-first-launch) the buffer contains
// the constructor-zeroed block; HEALTH_DIAG just prints
// zeroes — distinguishable from a real launch via the
// associated `episodes_collected` / step counter that other
// diags surface this epoch.
{
let decomp = if let Some(ref fused) = self.fused_ctx {
fused.trainer().read_sp18_reward_decomp_diag()
} else {
[0.0_f32; 20]
};
// Layout: bin*5 + col, with bin order Hold(0), Long(1),
// Short(2), Flat(3) and col order
// micro(0), opp(1), popart(2), abs(3), fire(4).
tracing::info!(
"HEALTH_DIAG[{}]: reward_decomp \
[hold(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4}) \
long(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4}) \
short(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4}) \
flat(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4})]",
epoch,
decomp[0], decomp[1], decomp[2], decomp[3], decomp[4],
decomp[5], decomp[6], decomp[7], decomp[8], decomp[9],
decomp[10], decomp[11], decomp[12], decomp[13], decomp[14],
decomp[15], decomp[16], decomp[17], decomp[18], decomp[19],
);
}
// SP7 observability (grad_decomp_pinned): surface the exact contents of
// grad_decomp_result_pinned[0..3, 12..15, 36..48] (iqn/cql_raw/c51 ×
// [mag, dir, trunk]) that the SP7 loss-balance controller read in the

342
crates/ml/tests/sp18_hold_reward_oracle_tests.rs Normal file
View File

@@ -0,0 +1,342 @@
//! SP18 v2 Phase 0 — D-leg + B-leg diagnostic-emit oracle tests.
//!
//! Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
//!
//! Phase 0 is observability-only: kernels write to new ISV slots
//! [493..505) and to a 20-float mapped-pinned diagnostic buffer; no
//! production-path consumer is wired. Tests therefore lock the
//! per-action reward decomposition kernel's bit-exact output against a
//! synthetic CPU oracle and confirm the host-side V_SHARE-trend +
//! TD-error-magnitude diagnostics emit the expected HEALTH_DIAG line
//! formats.
//!
//! Run on a GPU host:
//!
//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//! cargo test -p ml --test sp18_hold_reward_oracle_tests \
//! --features cuda -- --ignored --nocapture
//!
//! All B-tests (`#[ignore = "requires GPU"]`) load cubins via the
//! workspace `OUT_DIR` and use `MappedF32Buffer` per
//! `feedback_no_htod_htoh_only_mapped_pinned`.
// ── Layer A (CPU oracle): per-action reward decomposition math ───────────
/// CPU oracle for the SP18 D-leg per-action reward decomposition.
///
/// Construction (4 samples, 4 directions):
///
/// sample 0: dir=Hold, components=[popart=0.10, cf=0.0, trail=0.0,
/// micro=0.20, opp=0.05, bonus=0.0]
/// sample 1: dir=Long, components=[popart=0.50, cf=0.0, trail=0.0,
/// micro=0.30, opp=0.0, bonus=0.0]
/// sample 2: dir=Short, components=[popart=-0.40, cf=0.0, trail=0.0,
/// micro=-0.20, opp=0.0, bonus=0.0]
/// sample 3: dir=Flat, components=[popart=0.0, cf=0.0, trail=0.0,
/// micro=0.0, opp=0.10, bonus=0.0]
///
/// Expected per-bin output (bin order Hold, Long, Short, Flat):
///
/// Hold: micro=0.20, opp=0.05, popart= 0.10, |abs|=0.35, fire=1.0
/// Long: micro=0.30, opp=0.0, popart= 0.50, |abs|=0.80, fire=1.0
/// Short: micro=-0.20, opp=0.0, popart=-0.40, |abs|=0.60, fire=1.0
/// Flat: micro=0.0, opp=0.10, popart= 0.0, |abs|=0.10, fire=1.0
///
/// Pins the per-bin reduction math independent of GPU implementation.
#[test]
fn reward_decomp_per_action_cpu_oracle() {
// Components: [popart, cf, trail, micro, opp_cost, bonus]
let rc: Vec<[f32; 6]> = vec![
[ 0.10, 0.0, 0.0, 0.20, 0.05, 0.0], // Hold
[ 0.50, 0.0, 0.0, 0.30, 0.0, 0.0], // Long
[-0.40, 0.0, 0.0, -0.20, 0.0, 0.0], // Short
[ 0.0, 0.0, 0.0, 0.0, 0.10, 0.0], // Flat
];
// Logical bin → DIR_*: 0=Hold(1), 1=Long(2), 2=Short(0), 3=Flat(3)
const BIN_DIRS: [i32; 4] = [1, 2, 0, 3];
const N: usize = 4;
// Per-sample dir assignment matches the construction above.
let dirs: [i32; N] = [1, 2, 0, 3];
// CPU oracle: row-major 4×5 (bin × col).
let mut out = [0.0_f32; 20];
for (bin, target_dir) in BIN_DIRS.iter().enumerate() {
let mut s_micro = 0.0f32;
let mut s_opp = 0.0f32;
let mut s_popart = 0.0f32;
let mut s_abs = 0.0f32;
let mut s_count = 0.0f32;
let mut s_fire = 0.0f32;
for i in 0..N {
if dirs[i] != *target_dir { continue; }
let row = rc[i];
let abs_total: f32 = row.iter().map(|x| x.abs()).sum();
s_popart += row[0];
s_micro += row[3];
s_opp += row[4];
s_abs += abs_total;
s_count += 1.0;
if abs_total > 0.0 { s_fire += 1.0; }
}
let denom = if s_count > 0.0 { s_count } else { 1.0 };
let o = bin * 5;
out[o + 0] = if s_count > 0.0 { s_micro / denom } else { 0.0 };
out[o + 1] = if s_count > 0.0 { s_opp / denom } else { 0.0 };
out[o + 2] = if s_count > 0.0 { s_popart / denom } else { 0.0 };
out[o + 3] = if s_count > 0.0 { s_abs / denom } else { 0.0 };
out[o + 4] = if s_count > 0.0 { s_fire / denom } else { 0.0 };
}
// Hold (bin 0).
assert!((out[0] - 0.20).abs() < 1e-6, "hold micro: got {}", out[0]);
assert!((out[1] - 0.05).abs() < 1e-6, "hold opp: got {}", out[1]);
assert!((out[2] - 0.10).abs() < 1e-6, "hold popart: got {}", out[2]);
assert!((out[3] - 0.35).abs() < 1e-6, "hold abs: got {}", out[3]);
assert!((out[4] - 1.00).abs() < 1e-6, "hold fire: got {}", out[4]);
// Long (bin 1).
assert!((out[5] - 0.30).abs() < 1e-6, "long micro: got {}", out[5]);
assert!((out[6] - 0.00).abs() < 1e-6, "long opp: got {}", out[6]);
assert!((out[7] - 0.50).abs() < 1e-6, "long popart: got {}", out[7]);
assert!((out[8] - 0.80).abs() < 1e-6, "long abs: got {}", out[8]);
assert!((out[9] - 1.00).abs() < 1e-6, "long fire: got {}", out[9]);
// Short (bin 2).
assert!((out[10] - (-0.20)).abs() < 1e-6, "short micro: got {}", out[10]);
assert!((out[11] - 0.00 ).abs() < 1e-6, "short opp: got {}", out[11]);
assert!((out[12] - (-0.40)).abs() < 1e-6, "short popart: got {}", out[12]);
assert!((out[13] - 0.60 ).abs() < 1e-6, "short abs: got {}", out[13]);
assert!((out[14] - 1.00 ).abs() < 1e-6, "short fire: got {}", out[14]);
// Flat (bin 3).
assert!((out[15] - 0.00).abs() < 1e-6, "flat micro: got {}", out[15]);
assert!((out[16] - 0.10).abs() < 1e-6, "flat opp: got {}", out[16]);
assert!((out[17] - 0.00).abs() < 1e-6, "flat popart: got {}", out[17]);
assert!((out[18] - 0.10).abs() < 1e-6, "flat abs: got {}", out[18]);
assert!((out[19] - 1.00).abs() < 1e-6, "flat fire: got {}", out[19]);
}
// ── Layer B (GPU oracle): kernel matches CPU oracle bit-for-bit ──────────
#[cfg(feature = "cuda")]
#[allow(unsafe_code)]
mod gpu {
use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
/// Cubin handle for the SP18 D-leg per-action reward decomposition
/// diagnostic kernel.
const SP18_REWARD_DECOMP_DIAG_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/reward_decomp_diag_kernel.cubin"));
fn make_test_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_reward_decomp_diag(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP18_REWARD_DECOMP_DIAG_CUBIN.to_vec())
.expect("load reward_decomp_diag_kernel cubin");
module
.load_function("reward_decomp_diag")
.expect("load reward_decomp_diag function")
}
/// SP18 v2 Phase 0 Task 0.1 — GPU oracle for the per-action reward
/// decomposition diagnostic kernel.
///
/// Builds a 4-sample synthetic batch matching the CPU oracle above
/// (one sample per direction-axis bin), launches the kernel, and
/// asserts the 20-float mapped-pinned output matches the closed-form
/// expected values bit-for-bit (within f32 rounding budget).
///
/// Branch sizes (b1=b2=b3=3) match production; the factored-action
/// encoding is `dir * (b1*b2*b3) + ...` so DIR_HOLD=1 → action=27,
/// DIR_LONG=2 → action=54, DIR_SHORT=0 → action=0, DIR_FLAT=3 →
/// action=81.
#[test]
#[ignore = "requires GPU"]
fn reward_decomp_per_action_gpu_oracle() {
let stream = make_test_stream();
let kernel = load_reward_decomp_diag(&stream);
const N: usize = 4;
const NUM_COMPS: usize = 6;
const NUM_BINS: usize = 4;
const NUM_COLS: usize = 5;
const OUT_LEN: usize = NUM_BINS * NUM_COLS;
// ── Inputs ───────────────────────────────────────────────────
// reward_components_per_sample[N × 6], row-major.
let rc_host: Vec<f32> = vec![
// sample 0 — Hold
0.10, 0.0, 0.0, 0.20, 0.05, 0.0,
// sample 1 — Long
0.50, 0.0, 0.0, 0.30, 0.0, 0.0,
// sample 2 — Short
-0.40, 0.0, 0.0, -0.20, 0.0, 0.0,
// sample 3 — Flat
0.0, 0.0, 0.0, 0.0, 0.10, 0.0,
];
assert_eq!(rc_host.len(), N * NUM_COMPS);
let rc_buf = unsafe { MappedF32Buffer::new(N * NUM_COMPS) }
.expect("alloc rc buf");
rc_buf.write_from_slice(&rc_host);
// actions_out[N], factored-action ints. dir = action / (b1*b2*b3).
const B1: i32 = 3;
const B2: i32 = 3;
const B3: i32 = 3;
const DIR_DIV: i32 = B1 * B2 * B3; // 27
// DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3.
let actions_host: [i32; N] = [
1 * DIR_DIV, // Hold
2 * DIR_DIV, // Long
0 * DIR_DIV, // Short
3 * DIR_DIV, // Flat
];
let actions_buf = unsafe { MappedI32Buffer::new(N) }
.expect("alloc actions buf");
actions_buf.write_from_slice(&actions_host);
// Output buffer — 20 floats, mapped-pinned. Pre-fill with NaN
// so any unwritten lane shows up as a test failure.
let out_buf = unsafe { MappedF32Buffer::new(OUT_LEN) }
.expect("alloc out buf");
out_buf.write_from_slice(&vec![f32::NAN; OUT_LEN]);
// ── Launch ───────────────────────────────────────────────────
let n_i32 = N as i32;
const BLOCK_DIM: u32 = 256;
const NUM_BIN_BLOCKS: u32 = NUM_BINS as u32;
// Dynamic shmem = 5 × BLOCK_DIM × sizeof(f32) — matches the
// kernel's `s_buf` carve-up (5 reduction tiles).
let shmem_bytes: u32 = 5 * BLOCK_DIM * std::mem::size_of::<f32>() as u32;
unsafe {
stream
.launch_builder(&kernel)
.arg(&rc_buf.dev_ptr)
.arg(&actions_buf.dev_ptr)
.arg(&n_i32)
.arg(&B1)
.arg(&B2)
.arg(&B3)
.arg(&out_buf.dev_ptr)
.launch(LaunchConfig {
grid_dim: (NUM_BIN_BLOCKS, 1, 1),
block_dim: (BLOCK_DIM, 1, 1),
shared_mem_bytes: shmem_bytes,
})
.expect("launch reward_decomp_diag");
}
stream.synchronize().expect("sync after reward_decomp_diag");
// ── Verify ───────────────────────────────────────────────────
let out = out_buf.read_all();
for (i, x) in out.iter().enumerate() {
assert!(x.is_finite(), "out[{i}] is non-finite ({x}) — \
kernel did not write this lane");
}
// Closed-form expected values (matches the CPU oracle above).
let expected: [f32; OUT_LEN] = [
// Hold
0.20, 0.05, 0.10, 0.35, 1.00,
// Long
0.30, 0.00, 0.50, 0.80, 1.00,
// Short
-0.20, 0.00, -0.40, 0.60, 1.00,
// Flat
0.00, 0.10, 0.00, 0.10, 1.00,
];
let eps = 1e-6_f32;
for i in 0..OUT_LEN {
let bin = i / NUM_COLS;
let col = i % NUM_COLS;
let bin_name = match bin { 0 => "Hold", 1 => "Long", 2 => "Short", 3 => "Flat", _ => "?" };
let col_name = match col { 0 => "micro", 1 => "opp", 2 => "popart", 3 => "abs", 4 => "fire", _ => "?" };
assert!(
(out[i] - expected[i]).abs() < eps,
"out[{bin_name}.{col_name}] expected {:.6}, got {:.6} (diff {:.3e})",
expected[i], out[i], (out[i] - expected[i]).abs(),
);
}
}
/// Empty-bin guard test — when no samples land in a given bin,
/// the kernel must emit 0.0 for every column (NOT NaN, which would
/// be the natural division-by-zero default). The HEALTH_DIAG
/// consumer relies on this contract: a zero block is interpreted as
/// "bin saw no traffic this epoch", whereas NaN would cascade
/// downstream into the KILL CRITERION arithmetic.
#[test]
#[ignore = "requires GPU"]
fn reward_decomp_empty_bin_emits_zero_not_nan() {
let stream = make_test_stream();
let kernel = load_reward_decomp_diag(&stream);
const N: usize = 1;
const NUM_BINS: usize = 4;
const NUM_COLS: usize = 5;
const OUT_LEN: usize = NUM_BINS * NUM_COLS;
// Single Long-only sample → bins Hold/Short/Flat are all empty.
let rc_host: Vec<f32> = vec![0.5, 0.0, 0.0, 0.3, 0.0, 0.0];
let rc_buf = unsafe { MappedF32Buffer::new(rc_host.len()) }
.expect("alloc rc buf");
rc_buf.write_from_slice(&rc_host);
const B1: i32 = 3; const B2: i32 = 3; const B3: i32 = 3;
const DIR_DIV: i32 = B1 * B2 * B3;
let actions_host: [i32; N] = [2 * DIR_DIV]; // Long
let actions_buf = unsafe { MappedI32Buffer::new(N) }
.expect("alloc actions buf");
actions_buf.write_from_slice(&actions_host);
let out_buf = unsafe { MappedF32Buffer::new(OUT_LEN) }
.expect("alloc out buf");
out_buf.write_from_slice(&vec![f32::NAN; OUT_LEN]);
let n_i32 = N as i32;
const BLOCK_DIM: u32 = 256;
let shmem_bytes: u32 = 5 * BLOCK_DIM * std::mem::size_of::<f32>() as u32;
unsafe {
stream
.launch_builder(&kernel)
.arg(&rc_buf.dev_ptr)
.arg(&actions_buf.dev_ptr)
.arg(&n_i32)
.arg(&B1)
.arg(&B2)
.arg(&B3)
.arg(&out_buf.dev_ptr)
.launch(LaunchConfig {
grid_dim: (NUM_BINS as u32, 1, 1),
block_dim: (BLOCK_DIM, 1, 1),
shared_mem_bytes: shmem_bytes,
})
.expect("launch reward_decomp_diag");
}
stream.synchronize().expect("sync");
let out = out_buf.read_all();
// Hold (bin 0), Short (bin 2), Flat (bin 3) — all 5 cols must be 0.0.
for &bin in &[0_usize, 2, 3] {
for col in 0..NUM_COLS {
let v = out[bin * NUM_COLS + col];
assert!(
v == 0.0,
"empty-bin guard: bin={bin} col={col} expected 0.0, got {v}"
);
}
}
// Long (bin 1) — non-zero (sanity).
assert!(out[1 * NUM_COLS + 0] > 0.0, "long micro should be > 0");
assert!(out[1 * NUM_COLS + 4] > 0.0, "long fire should be > 0");
}
}

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

@@ -2,6 +2,53 @@
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
## 2026-05-08 — SP18 v2 Phase 0 Task 0.1: D-leg per-action reward decomposition diagnostic
Phase 0 observability: kernel + Rust launcher + per-epoch HEALTH_DIAG emit + GPU oracle test, all in the same atomic commit per `feedback_no_partial_refactor`.
**Files added/modified:**
- NEW `crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu` — block tree-reduce kernel (4 blocks × 256 threads, one block per direction-axis bin) reading `reward_components_per_sample [n_samples × 6]` + `actions_out [n_samples]` and emitting 5 stats per bin (mean r_micro / mean r_opp_cost / mean r_popart / mean |reward| / fire_rate) into a 20-float row-major output. Bin order: Hold(0)→Long(1)→Short(2)→Flat(3); col order: micro→opp→popart→abs→fire. Empty-bin guard emits 0.0 (NOT NaN) per the consumer-side KILL CRITERION arithmetic contract.
- `crates/ml/build.rs` — cubin manifest entry for `reward_decomp_diag_kernel.cu`.
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — `REWARD_DECOMP_DIAG_CUBIN` re-export + 20-float `MappedF32Buffer sp18_reward_decomp_diag_buf` field on `GpuDqnTrainer` + accessor pair (`sp18_reward_decomp_diag_dev_ptr` writer-side, `read_sp18_reward_decomp_diag` reader-side). Buffer is constructor-initialised to 0.0 so cold-start HEALTH_DIAG emits a deterministic zero block.
- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — `sp18_reward_decomp_diag_kernel` field + cubin load on the collector's stream + `launch_sp18_reward_decomp_diag(n, b1, b2, b3, out_dev_ptr)` launcher. Launches AT THE PER-STEP BOUNDARY, BEFORE `launch_reward_component_ema_inplace` (which `memset_zeros` the source buffer after consuming it) per `pearl_canary_input_freshness_launch_order`.
- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — wire the launch right before `launch_reward_component_ema_inplace`; emit a new HEALTH_DIAG line at the per-epoch boundary (right after the SP17 dueling line):
```
HEALTH_DIAG[N]: reward_decomp [hold(micro=X opp=Y popart=Z abs=W fire=F) long(...) short(...) flat(...)]
```
- NEW `crates/ml/tests/sp18_hold_reward_oracle_tests.rs`:
- `reward_decomp_per_action_cpu_oracle` — CPU oracle pinning the per-bin reduction math against a 4-sample synthetic batch (one sample per direction-axis).
- `reward_decomp_per_action_gpu_oracle` (`#[ignore = "requires GPU"]`) — GPU oracle: same synthetic batch, asserts kernel output matches the CPU oracle bit-for-bit (1e-6 f32 budget).
- `reward_decomp_empty_bin_emits_zero_not_nan` (`#[ignore]`) — empty-bin guard test.
**Atomicity envelope** (per `feedback_no_partial_refactor`): kernel + cubin manifest + Rust launcher + production wire-up + HEALTH_DIAG emit + GPU oracle test all land in a single commit. The per-epoch HEALTH_DIAG line is consumer; the kernel + launcher + buffer are producer; no intermediate state where one half exists without the other.
**Wire-up summary** (per Invariant 7):
| Layer | Site | Direction |
|---|---|---|
| Kernel | `reward_decomp_diag_kernel.cu::reward_decomp_diag` | producer (block tree-reduce, no atomicAdd) |
| Cubin | `build.rs` manifest entry | infra |
| Cubin re-export | `gpu_dqn_trainer.rs::REWARD_DECOMP_DIAG_CUBIN` | infra |
| Buffer | `gpu_dqn_trainer.rs::sp18_reward_decomp_diag_buf [20]` mapped-pinned | output sink |
| Kernel handle | `gpu_experience_collector.rs::sp18_reward_decomp_diag_kernel` | launcher |
| Launcher | `gpu_experience_collector.rs::launch_sp18_reward_decomp_diag` | producer call |
| Wire site | `training_loop.rs` per-step `collect_experiences_gpu` boundary, BEFORE `launch_reward_component_ema_inplace` | production wire |
| HEALTH_DIAG emit | `training_loop.rs` per-epoch boundary, after the SP17 `dueling` line | observability |
| GPU oracle test | `tests/sp18_hold_reward_oracle_tests.rs` | regression guard |
**Pearls + invariants:**
- `feedback_no_atomicadd` — block tree-reduce (one block per bin, 5 shmem accumulators each).
- `feedback_no_htod_htoh_only_mapped_pinned` — output is a `MappedF32Buffer.dev_ptr`; host reads via the mapped host_ptr after stream sync.
- `pearl_no_host_branches_in_captured_graph` — pure kernel launch, no host-side dispatch within the capture region. The wire site is outside `exp_fwd_graph` capture (per-step boundary, alongside `launch_reward_component_ema_inplace`).
- `pearl_canary_input_freshness_launch_order` — launches BEFORE `launch_reward_component_ema_inplace` so the per-action means are computed against the live `reward_components_per_sample` buffer (the latter zeroes it after consuming).
- `feedback_no_stubs` — full bin-aware reduction body, no zero-output placeholder; empty-bin guard emits 0.0 explicitly per the test contract.
**Plan**: `docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md` § Phase 0 Task 0.1.
## 2026-05-08 — SP18 v2 Pre-Phase Task PP.5: H100 collection-time profiling baseline (DEFERRED-TO-PHASE-4 MARKER)
PP.5 is the B-DD6 headroom-check task: profile current `gpu_experience_collector::collect_experience` end-to-end ms/cycle on H100 to confirm the new target-net forward pass (B-leg Phase 4) has the budget for an estimated ~30% slowdown. The plan explicitly marks this as **reviewer-only** ("Reviewer-run profiling … `./scripts/argo-train.sh --model dqn --epochs 1 --label sp18-pp5-profile --profile nsys`"). It does NOT modify production code — it just produces a baseline measurement for Phase 4 to validate against.