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

feat(sp22): H6 Phase 3a — 7-component contract migration + β producer (WIP)

Browse Source 
Phase 3a builds on the Phase A foundation (464bc5f7a). Migrates the
7-component reward_components_per_sample contract atomically across
producer + readers + buffer alloc, and installs the β producer at
training-side trade-close. α kernels exist (compiled in Phase A) but
are not yet launched in captured graphs — that's Phase 3b alongside
the SP11 controller extension and A2 eval-side aux infrastructure.

Why split into 3a/3b
────────────────────
Full Phase 3 (α + β + SP11 controller + A2) is ~25-35 hr engineering
spanning ~19 files. Phase 3a is the SAFE atomic contract migration
(7-stride buffer + β producer; no α captured-graph integration yet)
— runtime-equivalent to Phase 2 (β no-op at scale_β=0 sentinel; α
kernels loaded but never launched). This commits the foundation +
contract change as a clean checkpoint per
`feedback_no_partial_refactor` (the 7-component contract spans every
consumer; partial migration would produce stride mismatches; this
commit migrates ALL consumers that read the buffer).

Files
─────
- crates/ml/src/cuda_pipeline/experience_kernels.cu:
    Preamble doc → 7-component layout.
    Buffer stride `* 6 +` → `* 7 +` (~14 sites, atomic).
    7th-slot init at the per-step zero block (`rc[6] = 0.0f`).
    β producer at the segment_complete branch:
      r_aux_align = scale_β * max(0, aux × pos_sign) * max(0, pnl)
      with NULL-safe fallbacks (aux_dir_prob_per_env NULL OR
      isv_signals_ptr NULL → β no-op).
    Two new kernel args: aux_dir_prob_per_env + aux_align_scale_idx
    (slot index for scale_β, decoupled per the loss_cap_idx pattern).

- crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu:
    Stride `idx * 6` → `idx * 7` (3 sites). Iteration stays c=0..5;
    the 7th component (aux_align) is intentionally NOT EMA'd here.
    A dedicated reward_aux_align_ema_kernel writing directly to
    ISV[REWARD_AUX_ALIGN_EMA_INDEX=536] is Phase 3b scope (avoids
    extending the apply_pearls_ad chain).
    Preamble doc updated.

- crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu:
    #define RCP_NUM_COMPS 6 → 7. The kernel's per-bin abs-sum
    (col 3) now naturally includes r_aux_align; popart/micro/
    opp_cost per-bin means unchanged.

- crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu:
    Documentation only: aux_align excluded from the 6-axis
    cf_others ratio (non-contiguous with cf_others_base_slot at
    64..68; aux_align EMA lives at ISV[536]).

- crates/ml/src/cuda_pipeline/gpu_experience_collector.rs:
    Buffer alloc `total_output * 6` → `* 7` (critical for runtime
    safety — partial migration would produce OOB writes since
    experience_env_step writes to `out_off * 7 + N`).
    experience_env_step launcher gains 2 new `.arg(...)` calls
    passing `self.prev_aux_dir_prob.raw_ptr()` and
    `SP22_AUX_ALIGN_SCALE_INDEX as i32`.

- docs/dqn-wire-up-audit.md:
    Phase 3a entry documenting the partial commit + Phase 3b
    remaining-work breakdown.

Verification
────────────
- cargo check -p ml --features cuda: 0 errors, 21 pre-existing
  warnings (Phase 2 baseline parity).
- All nvcc cubins recompile (experience_kernels, reward_component_ema,
  reward_decomp_diag, reward_component_mag_ratio_compute, plus
  Phase A's aux_to_q_dir_bias_kernel + backward).
- Runtime equivalent to Phase 2: β no-op (scale_β=0 sentinel since
  SP11 controller not yet extended), α no-op (kernels dead-code
  until Phase 3b wires them into captured graphs).

Phase 3b scope (resume in fresh session)
────────────────────────────────────────
- B6: SP11 controller extension (w_aux_align emit at ISV[537])
- B7: HEALTH_DIAG snap layout extension
- B9: α plumbing — W_aux_to_Q_dir param + Adam + captured-graph
       forward + backward integration in gpu_dqn_trainer.rs
- B10/B11: HEALTH_DIAG print-line extensions
- C1: α forward in collector's rollout-time captured graph
- D1-D7: A2 eval-side aux trunk + α + state-gather wiring
- E + F: verification gates + atomic Phase F commit + smoke + verdict

Refs
────
- docs/plans/2026-05-12-sp22-h6-phase3-alpha-beta.md (spec)
- docs/plans/2026-05-13-sp22-h6-phase3-alpha-beta-runbook.md (runbook)
- 464bc5f7a (Phase A foundation)
- pearl_no_partial_refactor (atomic 7-component contract migration)
- pearl_event_driven_reward_density_alignment (β at segment_complete)
- pearl_one_unbounded_signal_per_reward (β bounded by scale_β + alignment caps)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-13 00:56:54 +02:00
parent 464bc5f7a4
commit 2e4c7ebf60
6 changed files with 249 additions and 34 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -2051,16 +2051,27 @@ extern "C" __global__ void experience_env_step(
const float* __restrict__ magnitude_conviction_ptr,
/* C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2).
* Per-sample per-component reward magnitudes written at every (i,t).
* Layout [N*L, 6] row-major; component indices:
* [0] popart — final on-policy reward (denominator for PopArt drift)
* [1] cf — counterfactual reward written to CF slot
* [2] trail — 0.0 (placeholder; future term)
* [3] micro — dense OFI micro-reward (R5)
* [4] opp_cost — Flat opportunity-cost reward (Plan 3 Task 2 B.1)
* [5] bonus — novelty-scaled trade-attempt bonus (Plan 3 Task 3 B.2)
* Layout [N*L, 7] row-major (SP22 H6 Phase 3 contract migration);
* component indices:
* [0] popart — final on-policy reward (denominator for PopArt drift)
* [1] cf — counterfactual reward written to CF slot
* [2] trail — 0.0 (placeholder; future term)
* [3] micro — dense OFI micro-reward (R5)
* [4] opp_cost — Flat opportunity-cost reward (Plan 3 Task 2 B.1)
* [5] bonus — novelty-scaled trade-attempt bonus (Plan 3 Task 3 B.2)
* [6] aux_align — SP22 H6 Phase 3 β: event-driven aux-aligned trade-close
* bonus. Fires at `segment_complete` when policy direction
* aligns with aux conviction AND trade was profitable.
* `r_aux_align = scale_β × alignment × profit_pos` where
* `alignment = max(0, aux_at_close × position_sign)` and
* `profit_pos = max(0, realized_pnl)`. scale_β read from
* ISV[SP22_AUX_ALIGN_SCALE_INDEX] (controller-emitted).
* NULL-safe via two fallbacks: aux_dir_prob_per_env NULL
* OR isv_signals_ptr NULL → β no-op.
* Producer: this kernel (experience_env_step).
* Consumer: reward_component_ema GPU kernel → ISV[63..69). */
float* __restrict__ reward_components_per_sample, /* [N*L * 6] */
* Consumer: reward_component_ema GPU kernel → ISV[63..69) and
* ISV[REWARD_AUX_ALIGN_EMA_INDEX=536]. */
float* __restrict__ reward_components_per_sample, /* [N*L * 7] */
/* B.2 Plan 3 Task 3: Flat→Positioned transition flag per (i,t).
* Written unconditionally (0/1) at every reached slot. Consumer:
* trade_attempt_rate_ema_update → ISV[TRADE_ATTEMPT_RATE_EMA_INDEX]. */
@@ -2310,7 +2321,26 @@ extern "C" __global__ void experience_env_step(
* gather (sample_proportional → GpuBatchPtrs.aux_conf_ptr), and
* eventual Phase 5 Aux→Q gate read in the Bellman-target loss
* computation. */
float* __restrict__ aux_conf_per_sample
float* __restrict__ aux_conf_per_sample,
/* SP22 H6 Phase 3 β (2026-05-13) — per-env recentered aux directional
* probability from the PREVIOUS step's aux trunk forward. Layout [N]
* f32. Same buffer as `experience_state_gather`'s arg of the same name
* — both are populated by `aux_softmax_to_per_env_kernel`. Read here
* at the `segment_complete` branch to compute the alignment factor
* for `r_aux_align = scale_β × max(0, aux_at_close × position_sign)
* × max(0, realized_pnl)`. NULL-safe: when NULL (test scaffolds, or
* pre-aux-wireup), `aux_at_close = 0` → `alignment = 0` → β no-op.
* Identical to Phase 2 recentered encoding: 0 = neutral, +1 = up
* with full conviction, -1 = down with full conviction. */
const float* __restrict__ aux_dir_prob_per_env,
/* SP22 H6 Phase 3 β — ISV slot index for the SP11-controller-emitted
* adaptive scale_β (canonical SP22_AUX_ALIGN_SCALE_INDEX = 537).
* Passed as a kernel arg per the existing pattern (mirrors
* `loss_cap_idx`, `alpha_ema_idx`, `hold_cost_scale_idx`,
* `hold_reward_ema_idx`) so the kernel stays decoupled from SP slot
* drift. The kernel reads `isv_signals_ptr[aux_align_scale_idx]`
* NULL-safely (returns 0 if isv_signals_ptr == NULL → β no-op). */
int aux_align_scale_idx
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -2446,7 +2476,7 @@ extern "C" __global__ void experience_env_step(
* write pattern self-documenting and forward-safe. Components [4, 5] are
* overwritten later by the Flat opp-cost and Flat→Positioned bonus paths. */
{
float* rc = reward_components_per_sample + out_off * 6;
float* rc = reward_components_per_sample + out_off * 7;
rc[0] = 0.0f; /* popart (final reward) — overwritten below */
rc[1] = 0.0f; /* cf reward — overwritten below at cf_off */
rc[2] = 0.0f; /* trail penalty (future term) */
@@ -3296,12 +3326,17 @@ extern "C" __global__ void experience_env_step(
* Note: rc[+0] (PopArt input = total reward) is unconditionally
* overwritten at ~L3447, so it doesn't strictly need zero-init here, but
* including it removes ordering dependency for future readers. */
reward_components_per_sample[out_off * 6 + 0] = 0.0f;
reward_components_per_sample[out_off * 6 + 1] = 0.0f;
reward_components_per_sample[out_off * 6 + 2] = 0.0f;
reward_components_per_sample[out_off * 6 + 3] = 0.0f;
reward_components_per_sample[out_off * 6 + 4] = 0.0f;
reward_components_per_sample[out_off * 6 + 5] = 0.0f;
reward_components_per_sample[out_off * 7 + 0] = 0.0f;
reward_components_per_sample[out_off * 7 + 1] = 0.0f;
reward_components_per_sample[out_off * 7 + 2] = 0.0f;
reward_components_per_sample[out_off * 7 + 3] = 0.0f;
reward_components_per_sample[out_off * 7 + 4] = 0.0f;
reward_components_per_sample[out_off * 7 + 5] = 0.0f;
/* SP22 H6 Phase 3 β (2026-05-13): 7th reward component slot.
* Overwritten below at segment_complete with `scale_β × alignment
* × profit_pos` when policy direction aligns with aux conviction
* AND trade was profitable; otherwise stays at 0 (β no-op). */
reward_components_per_sample[out_off * 7 + 6] = 0.0f;
/* Trade-cumulative running scalar — tracks the in-progress reward
* across the if/else cascade. C.4/D.4b bonus blocks read this to
@@ -3504,7 +3539,7 @@ extern "C" __global__ void experience_env_step(
* trail-fire branch fires (was a structural placeholder pre-B1b).
* The reward_component_ema kernel will pick up real trail magnitude
* for ISV-driven component-mag-ratio canaries. */
reward_components_per_sample[out_off * 6 + 2] = r_trail;
reward_components_per_sample[out_off * 7 +2] = r_trail;
/* SP12 v3 Change 3: lump-sum opportunity cost on trade close.
*
@@ -3545,7 +3580,52 @@ extern "C" __global__ void experience_env_step(
r_opp_cost = compute_lump_sum_opp_cost(
shaping_scale, holding_cost_rate, pre_trade_position, segment_hold_time
);
reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
reward_components_per_sample[out_off * 7 +4] = r_opp_cost;
/* ── SP22 H6 Phase 3 β (2026-05-13): aux-aligned trade-close bonus ──
*
* Fires at segment_complete when policy direction aligns with aux
* conviction AND trade was profitable. Non-negative-only (no
* anti-alignment penalty; let r_trail / step_ret_core's negative
* carry the loss signal — let r_opp_cost's negative carry
* cost-dominance). Bounded by `pearl_one_unbounded_signal_per_reward`:
* `alignment ∈ [0, 1]`, `scale_β` controller-bound to [0.05, 2.0],
* `step_ret_core` is the natural per-trade scale (single unbounded
* multiplicand per the pearl).
*
* NULL-safe via two fallbacks: `aux_dir_prob_per_env == NULL` (test
* scaffolds) → `aux_at_close = 0` → `alignment = 0` → β no-op;
* `isv_signals_ptr == NULL` (same path) → `scale_β = 0` → β no-op.
*
* `pre_trade_position` is the position held BEFORE the exit
* (`unified_env_step_core` mutates `position` in-place; the saved
* pre-step value is what defines the direction of the trade that
* just closed). `step_ret_core` is the realized P&L of this
* trade-close (the segment return).
*
* Pearl notes:
* - `pearl_event_driven_reward_density_alignment` — β fires
* ONLY at segment_complete (event-driven density matches the
* trade-close objective; not per-step shaping).
* - `pearl_one_unbounded_signal_per_reward` — only step_ret_core
* is unbounded; alignment + scale_β are bounded.
* - `pearl_first_observation_bootstrap` — scale_β cold-starts at
* 0 (β no-op until SP11 controller's first emit). */
{
float position_sign = (pre_trade_position > 0.001f) ? 1.0f
: (pre_trade_position < -0.001f) ? -1.0f
: 0.0f;
float aux_at_close = (aux_dir_prob_per_env != NULL)
? aux_dir_prob_per_env[i]
: 0.0f;
float alignment = fmaxf(0.0f, aux_at_close * position_sign);
float profit_pos = fmaxf(0.0f, step_ret_core);
float scale_beta_isv = (isv_signals_ptr != NULL)
? isv_signals_ptr[aux_align_scale_idx]
: 0.0f;
float r_aux_align = scale_beta_isv * alignment * profit_pos;
reward_components_per_sample[out_off * 7 + 6] = r_aux_align;
}
/* ── Layer 3: CEA counterfactual loop REMOVED (4-branch: direction(4) replaces exposure(9)).
* Dense micro-reward (Gem 1) below provides directional gradient signal. ── */
@@ -3606,7 +3686,7 @@ extern "C" __global__ void experience_env_step(
r_bonus += timing_bonus;
reward += timing_bonus; /* keep cascade scalar in lockstep for Q-cap reads below */
/* C.2 attribution — accumulate into rc[5] bonus slot (see above). */
reward_components_per_sample[out_off * 6 + 5] += timing_bonus;
reward_components_per_sample[out_off * 7 +5] += timing_bonus;
}
/* ── D.4a (Plan 3 Task 6a): Persistence credit ──
@@ -3627,7 +3707,7 @@ extern "C" __global__ void experience_env_step(
* tanhf(reward / fmaxf(1e-4f, drawdown_depth));
r_bonus += persist_bonus;
reward += persist_bonus; /* keep cascade scalar in lockstep for Q-cap reads below */
reward_components_per_sample[out_off * 6 + 5] += persist_bonus;
reward_components_per_sample[out_off * 7 +5] += persist_bonus;
}
}
@@ -3679,7 +3759,7 @@ extern "C" __global__ void experience_env_step(
* bars_late_frac * reward_capped;
r_bonus -= penalty;
reward -= penalty; /* keep cascade scalar in lockstep for Q-cap reads below */
reward_components_per_sample[out_off * 6 + 5] -= penalty;
reward_components_per_sample[out_off * 7 +5] -= penalty;
}
/* Reset shift-bar for the NEXT trade (in addition to Step 3's lifecycle resets,
* this handles the exit → next-entry transition cleanly). */
@@ -3881,7 +3961,7 @@ extern "C" __global__ void experience_env_step(
}
reward = r_micro;
micro_reward_per_sample[out_off] = r_micro;
reward_components_per_sample[out_off * 6 + 3] = r_micro;
reward_components_per_sample[out_off * 7 +3] = r_micro;
} else if (!segment_complete) {
/* Plan 3 D.4c: Welford-style EMA of conviction_core during Flat bars.
* Updates running mean and squared-deviation EMAs that the consumer
@@ -3986,7 +4066,7 @@ extern "C" __global__ void experience_env_step(
}
}
reward = r_opp_cost;
reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
reward_components_per_sample[out_off * 7 +4] = r_opp_cost;
}
/* ── B.2 (Plan 3 Task 3): ISV-driven trade-attempt bonus ──
@@ -4033,7 +4113,7 @@ extern "C" __global__ void experience_env_step(
r_bonus += bonus_scaled;
reward += bonus_scaled; /* keep cascade scalar in lockstep */
/* C.2 attribution: record bonus component magnitude at rc[5]. */
reward_components_per_sample[out_off * 6 + 5] = bonus_scaled;
reward_components_per_sample[out_off * 7 +5] = bonus_scaled;
}
}
@@ -4083,7 +4163,7 @@ extern "C" __global__ void experience_env_step(
const float bonus = shaping_scale * vol_proxy_c * stability * conviction_core;
r_bonus += bonus;
reward += bonus; /* keep cascade scalar in lockstep */
reward_components_per_sample[out_off * 6 + 5] += bonus;
reward_components_per_sample[out_off * 7 +5] += bonus;
}
}
/* Defensive reset of D.4c EMA slots at entry — redundant with the
@@ -4355,7 +4435,7 @@ extern "C" __global__ void experience_env_step(
* to c51_loss_kernel.cu::block_bellman_project_f. */
total_reward_per_sample[out_off] = reward;
/* C.2: component [0] = final on-policy reward (PopArt input denominator). */
reward_components_per_sample[out_off * 6 + 0] = reward;
reward_components_per_sample[out_off * 7 +0] = reward;
/* SP11 Fix 39 (A1.2): per-bar saboteur engagement signal — see param
* docstring above. Computed after `total_reward_per_sample` is
@@ -4580,7 +4660,7 @@ extern "C" __global__ void experience_env_step(
* (popart-component) is fed by raw r_popart via
* popart_component_per_sample. Only rc[+0] is intentionally
* post-composition (PopArt input). */
reward_components_per_sample[out_off * 6 + 1] = cf_reward;
reward_components_per_sample[out_off * 7 +1] = cf_reward;
}
/* ---- Advance timestep counter ---- */

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

@@ -2042,8 +2042,13 @@ impl GpuExperienceCollector {
// [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).
// Consumer: reward_component_ema GPU kernel → ISV[63..69).
// SP22 H6 Phase 3 (2026-05-13): bumped 6 → 7 for the new aux_align
// reward component (β producer at trade-close in experience_env_step).
// Stride must match `experience_kernels.cu`'s `out_off * 7 +` writes;
// partial migration would produce OOB writes (kernel writes slot 6
// past row end of a 6-wide buffer). Per `feedback_no_partial_refactor`.
let reward_components_per_sample = stream
.alloc_zeros::<f32>(total_output * 6)
.alloc_zeros::<f32>(total_output * 7)
.map_err(|e| MLError::ModelError(format!("alloc reward_components_per_sample: {e}")))?;
// B.2 Plan 3 Task 3: per-sample Flat→Positioned transition flag.
// 1 iff (prev_sign == 0 && curr_sign != 0) for this (i,t), 0 otherwise.
@@ -6605,6 +6610,25 @@ impl GpuExperienceCollector {
// end of `collect_experiences_gpu`, mirroring the
// `aux_sign_labels` `[N × L × 2]` final shape).
.arg(&mut self.aux_conf_per_sample)
// SP22 H6 Phase 3 β (2026-05-13): two new args for the
// event-driven aux-aligned trade-close bonus.
// - aux_dir_prob_per_env: same per-env recentered p_up
// buffer that experience_state_gather reads (Phase 1+2
// H6 wiring); β reads at segment_complete to compute
// alignment. Reuses the existing prev_aux_dir_prob
// buffer; no new allocation needed.
// - aux_align_scale_idx: ISV slot index for the
// SP11-controller-emitted adaptive scale_β. In
// Phase 3a (this commit) the SP11 controller has not
// yet been extended, so the slot stays at sentinel 0
// and β is functionally a no-op (r_aux_align = 0).
// Phase 3b lands the controller extension to bring
// scale_β non-zero.
.arg(&self.prev_aux_dir_prob.raw_ptr())
.arg(&{
use crate::cuda_pipeline::sp22_isv_slots::SP22_AUX_ALIGN_SCALE_INDEX;
SP22_AUX_ALIGN_SCALE_INDEX as i32
})
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_env_step t={t}: {e}"

15
crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu
View File

@@ -14,7 +14,11 @@
* `popart_component_per_sample` GPU buffer (B1b fix-up).
*
* Reads per-sample per-component rewards from `reward_components_per_sample`
* [n_samples * 6], where the 6 components are laid out as:
* [n_samples * 7] (SP22 H6 Phase 3 stride change; this kernel still
* processes components 0..5 only, the 7th component aux_align is
* handled by the separate `reward_aux_align_ema_kernel` writing
* directly to ISV[REWARD_AUX_ALIGN_EMA_INDEX=536]), where the
* 6 components processed here are laid out as:
*
* [c=0] popart — final on-policy reward (pre-PopArt; denominator signal)
* [c=1] cf — counterfactual replay reward written to the CF slot
@@ -22,6 +26,9 @@
* [c=3] micro — dense OFI-alignment micro-reward (R5)
* [c=4] opp_cost — Flat opportunity cost (Plan 3 Task 2 B.1; ISV[21]-scaled)
* [c=5] bonus — 0.0 placeholder (future trade-attempt/timing bonus)
* [c=6] aux_align — SP22 H6 Phase 3 (NOT processed by this kernel;
* see `reward_aux_align_ema_kernel` for the EMA
* into ISV[REWARD_AUX_ALIGN_EMA_INDEX=536])
*
* Reduces mean |r_c| over the batch per component, then writes the step
* observation to `scratch_buf[scratch_first_index + c]` for c in 0..6.
@@ -54,7 +61,7 @@
*/
extern "C" __global__ void reward_component_ema(
const float* __restrict__ reward_components, /* [n_samples * 6] row-major */
const float* __restrict__ reward_components, /* [n_samples * 7] row-major (SP22 H6 Phase 3 stride; kernel reads only c=0..5) */
int n_samples,
float* __restrict__ scratch_buf, /* producer_step_scratch_buf */
int scratch_first_index, /* slot 63 — first of 6 contiguous mag slots */
@@ -75,7 +82,7 @@ extern "C" __global__ void reward_component_ema(
/* ── Pass 1: mean |r| over batch for this component. ────────────── */
float acc = 0.0f;
for (int idx = 0; idx < n_samples; ++idx) {
acc += fabsf(reward_components[idx * 6 + c]);
acc += fabsf(reward_components[idx * 7 + c]);
}
float mean_abs = acc / fmaxf(1.0f, (float)n_samples);
@@ -88,7 +95,7 @@ extern "C" __global__ void reward_component_ema(
if (c >= 1 && c <= 5) {
float v_acc = 0.0f;
for (int idx = 0; idx < n_samples; ++idx) {
float d = fabsf(reward_components[idx * 6 + c]) - mean_abs;
float d = fabsf(reward_components[idx * 7 + c]) - mean_abs;
v_acc += d * d;
}
float var_abs = v_acc / fmaxf(1.0f, (float)n_samples);

11
crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
View File

@@ -18,6 +18,17 @@
// axis 4 (opp_cost) = isv[cf_others_base_slot + 3] ← slot 67
// axis 5 (bonus) = isv[cf_others_base_slot + 4] ← slot 68
//
// SP22 H6 Phase 3 (2026-05-13): aux_align (component idx 6 in the
// 7-stride reward_components_per_sample buffer) is INTENTIONALLY
// excluded from this kernel's 6-axis ratio computation. The aux_align
// EMA lives at ISV[REWARD_AUX_ALIGN_EMA_INDEX=536] (non-contiguous
// with cf_others_base_slot at 64..68), and the SP11 controller's
// `w_aux_align` adaptive weight has its own anchor on that slot
// directly (see reward_subsystem_controller_kernel.cu Phase 3 extension).
// Keeping aux_align out of this kernel's ratio preserves the
// pre-Phase-3 6-axis mag/var z-score machinery bit-for-bit; only the
// SP11 controller's emit list gains the 7th weight.
//
// And per-component variance sources (B1b smoke-recovery):
//
// axis 0 var = isv[popart_var_slot] ← slot 361

6
crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu
View File

@@ -95,7 +95,11 @@
#define RCP_POPART_IDX 0
#define RCP_MICRO_IDX 3
#define RCP_OPP_COST_IDX 4
#define RCP_NUM_COMPS 6
/* SP22 H6 Phase 3 (2026-05-13): bumped 6 → 7 for the new aux_align
* reward component (c=6). The kernel's abs-sum aggregate (col 3
* `mean(|reward|)`) now naturally includes r_aux_align in the per-bin
* total — popart/micro/opp_cost per-bin means remain unchanged. */
#define RCP_NUM_COMPS 7
/* Output layout: 4 bins × 5 cols = 20 floats. */
#define SP18_DECOMP_NUM_BINS 4