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

feat(sp22): H6 Phase 3 — Phase A foundation (WIP checkpoint, additive)

Browse Source 
Phase A of the SP22 H6 Phase 3 implementation runbook
(docs/plans/2026-05-13-sp22-h6-phase3-alpha-beta-runbook.md). Purely
additive: new ISV slots + new kernel files + build.rs registration
+ state_reset_registry entries. Nothing in the running training or
eval pipeline consumes this infrastructure yet — the existing
6-component contract is untouched. Phase A is committable as a clean
WIP foundation; Phases B-F (7-component contract migration, α
plumbing, A2 eval-side wiring, verification, smoke) resume in a
future session.

Files
─────
- crates/ml/src/cuda_pipeline/sp22_isv_slots.rs (NEW)
    REWARD_AUX_ALIGN_EMA_INDEX = 536 — EMA of 7th reward component
    SP22_AUX_ALIGN_SCALE_INDEX  = 537 — SP11-controller scale_β

- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
    ISV_TOTAL_DIM bumped 536 → 538 (SP22 H6 Phase 3 +2 slots)
    Layout fingerprint extended with SLOT_536 + SLOT_537 entries

- crates/ml/src/cuda_pipeline/mod.rs
    pub mod sp22_isv_slots

- crates/ml/src/trainers/dqn/state_reset_registry.rs
    sp22_reward_aux_align_ema (FoldReset, sentinel 0)
    sp22_aux_align_scale       (FoldReset, sentinel 0)

- crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu (NEW)
    α forward: Q_dir[b, a] += W_aux[a] * state_121[b]
    Reads state_121 from encoder INPUT (batch_states), bypassing
    the cold encoder weight for slot 121. No atomicAdd, no host
    branches, capture-safe.

- crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu (NEW)
    Two kernels in one .cu:
    - aux_to_q_dir_bias_backward_dw: per-action block tree-reduce
      computing dW[a] = Σ_b state_121[b] * dq_dir[b, a]
    - aux_to_q_dir_bias_backward_dstate: per-sample sum
      dstate_121[b] += Σ_a W[a] * dq_dir[b, a] (option-(i) gradient
      routing: both α and encoder paths get signal)

- crates/ml/build.rs
    Both new cubins registered in kernels_with_common (compiled
    successfully by nvcc; cubins exist at OUT_DIR/aux_to_q_dir_bias_*.cubin)

- docs/dqn-wire-up-audit.md
    Phase A entry documenting the checkpoint + remaining-work
    breakdown for next session.

Verification
────────────
- cargo check -p ml --features cuda: 0 errors, 21 pre-existing
  warnings (Phase 2 baseline parity)
- nvcc compiles both new cubins (3.6 KB fwd, 10.6 KB bwd)
- No runtime impact: no consumer wires the new infrastructure yet

Resumes in
──────────
Future session executes Phases B-F per the runbook:
- B: 11 tasks — 7-component contract migration cascade
- C: 1 task — α collector rollout-side wiring
- D: 7 tasks — A2 eval-side aux trunk + α + state-gather
- E: 7 tasks — verification gates
- F: 5 tasks — audit doc append + atomic commit (incl. this Phase A
  + Phase B-F changes) + push + smoke + verdict

Estimated remaining: ~28-43 hr engineering + ~37 min smoke wall-clock.

Phase B partial work (300-line diff for experience_kernels.cu + reward_component_ema_kernel.cu — 7-stride migration + β producer + preamble update) saved at /tmp/sp22-h6-phase3-b-partial.patch (local-only).

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)
- pearl_no_partial_refactor (Phase A is additive; safe to commit alone)
- pearl_no_atomicadd (backward block tree-reduce, no atomics)
- pearl_no_host_branches_in_captured_graph (kernel capture-safety)

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

13
crates/ml/build.rs
View File

@@ -805,6 +805,19 @@ fn main() {
// stream right after `aux_next_bar_forward` in
// `gpu_experience_collector.rs`.
"aux_softmax_to_per_env_kernel.cu",
// SP22 H6 Phase 3 α (2026-05-13): post-encoder Q_dir bias.
// Forward kernel adds `W_aux_to_Q_dir[a] * state_121[b]` to
// Q_dir[b, a]. Reads state_121 from encoder INPUT (batch_states),
// bypassing the cold encoder weight for slot 121. Backward
// kernel has TWO entry points: `aux_to_q_dir_bias_backward_dw`
// (per-action block tree-reduce; one block per action) and
// `aux_to_q_dir_bias_backward_dstate` (per-sample sum across
// 4 actions, accumulating into encoder-input gradient slot 121).
// No atomicAdd per `feedback_no_atomicadd`; no host branches
// per `pearl_no_host_branches_in_captured_graph` (capture-safe
// for both training and rollout forward graphs).
"aux_to_q_dir_bias_kernel.cu",
"aux_to_q_dir_bias_backward_kernel.cu",
// SP14 Layer C Phase C.4b (2026-05-08): adaptive aux prediction
// horizon producer. Single-thread kernel writing
// `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` from

132
crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu Normal file
View File

@@ -0,0 +1,132 @@
// crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu
//
// SP22 H6 Phase 3 alpha backward (2026-05-13): gradients for the
// `aux_to_q_dir_bias` forward kernel.
//
// Forward (from aux_to_q_dir_bias_kernel.cu):
// q_dir[b, a] += w_aux[a] * state_121[b]
//
// Backward (this file, two kernels in one .cu):
// dW_aux[a] = Sum_b state_121[b] * dq_dir[b, a] (per-action reduce)
// dstate_121[b] += Sum_a w_aux[a] * dq_dir[b, a] (per-sample sum)
//
// Discipline
// ----------
// - No atomicAdd per `feedback_no_atomicadd.md`. The dW reduce uses
// per-action block tree-reduce (one block per action, 256 threads
// each). The dstate accumulation is per-sample (no reduction; each
// thread writes its own b slot, possibly accumulating with the
// encoder backward chain's existing write via `+=`).
// - No host branches per `pearl_no_host_branches_in_captured_graph`.
// - Pure GPU compute per `feedback_no_cpu_compute_strict.md`.
// - Same-stream contract with the upstream backward producer that
// wrote dq_dir.
//
// Encoder-state[121] gradient routing
// ------------------------------------
// dstate_121 ACCUMULATES (+=) into dbatch_states[b, AUX_DIR_PROB_INDEX]
// per spec option (i): both alpha's W and the encoder's column for
// slot 121 get gradient signal. Long-term robust; alpha dominates
// initially because encoder weights for slot 121 are near-zero.
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
/* -------------------------------------------------------------------- */
/* Kernel: aux_to_q_dir_bias_backward_dw */
/* */
/* Computes dW_aux[a] = Sum_b state_121[b] * dq_dir[b, a]. */
/* */
/* Launch: grid=(b0_size, 1, 1), block=(256, 1, 1). One block per */
/* action; each block reduces the per-action column of dq_dir × state. */
/* -------------------------------------------------------------------- */
extern "C" __global__ void aux_to_q_dir_bias_backward_dw(
/* Upstream dQ_dir [B, b0_size]. */
const float* __restrict__ dq_dir,
/* Encoder INPUT buffer (same buffer the forward kernel read). */
const float* __restrict__ batch_states,
/* Output: dW_aux [b0_size] f32. Each block writes one slot. */
float* __restrict__ dw_aux,
/* AUX_DIR_PROB_INDEX (= 121 in canonical Phase-2 layout). */
int aux_dir_prob_index,
/* State dim (= 128 in canonical layout). */
int state_dim,
/* Batch size. */
int batch_size,
/* Number of direction actions. */
int b0_size
) {
int a = blockIdx.x;
if (a >= b0_size) return;
int tid = threadIdx.x;
__shared__ float sh[BLOCK_SIZE];
/* Per-thread strided accumulation across [0, B). Each thread folds
* its slice of (state_121[b] * dq_dir[b, a]) into a local sum. */
float local_sum = 0.0f;
for (int b = tid; b < batch_size; b += BLOCK_SIZE) {
float state_121 = batch_states[(long long)b * state_dim + aux_dir_prob_index];
float grad = dq_dir[(long long)b * b0_size + a];
local_sum += state_121 * grad;
}
sh[tid] = local_sum;
__syncthreads();
/* log2(BLOCK_SIZE)=8 step tree reduce. */
for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
if (tid < s) sh[tid] += sh[tid + s];
__syncthreads();
}
/* Thread 0 writes the per-action gradient. */
if (tid == 0) {
dw_aux[a] = sh[0];
}
}
/* -------------------------------------------------------------------- */
/* Kernel: aux_to_q_dir_bias_backward_dstate */
/* */
/* Accumulates dstate_121[b] += Sum_a w_aux[a] * dq_dir[b, a] into */
/* dbatch_states[b * state_dim + aux_dir_prob_index]. */
/* */
/* Launch: grid=((B + 255) / 256, 1, 1), block=(256, 1, 1). One thread */
/* per sample. The b0_size inner loop is unrolled by the compiler for */
/* b0_size=4. */
/* -------------------------------------------------------------------- */
extern "C" __global__ void aux_to_q_dir_bias_backward_dstate(
/* Upstream dQ_dir [B, b0_size]. */
const float* __restrict__ dq_dir,
/* alpha weight vector [b0_size]. Same buffer as the forward kernel. */
const float* __restrict__ w_aux,
/* INOUT: encoder-input gradient [B, state_dim]. This kernel
* touches only slot aux_dir_prob_index, ACCUMULATING (+=) into
* whatever the encoder backward chain has already written there.
* Option (i) gradient routing per spec: both alpha and encoder
* paths get signal. */
float* __restrict__ dbatch_states,
/* AUX_DIR_PROB_INDEX (= 121). */
int aux_dir_prob_index,
/* State dim (= 128). */
int state_dim,
/* Batch size. */
int batch_size,
/* Number of direction actions (= b0_size). */
int b0_size
) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= batch_size) return;
/* Sum_a w_aux[a] * dq_dir[b, a]. b0_size=4 in practice. */
float sum = 0.0f;
for (int a = 0; a < b0_size; a++) {
sum += w_aux[a] * dq_dir[(long long)b * b0_size + a];
}
/* Accumulate into the encoder-input gradient slot 121. The encoder
* backward chain may have already written zero or already-computed
* values here; we add (option (i) gradient routing per spec). */
dbatch_states[(long long)b * state_dim + aux_dir_prob_index] += sum;
}

84
crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu Normal file
View File

@@ -0,0 +1,84 @@
// crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu
//
// SP22 H6 Phase 3 alpha forward (2026-05-13): post-encoder Q_dir bias
// from the aux signal. Adds `W_aux_to_Q_dir[a] * state_121[b]` to
// `Q_dir[b, a]` for each sample b in [0, B) and action a in
// [0, b0_size=4).
//
// Architectural role
// -------------------
// Routes the aux signal around the cold-encoder bottleneck: state[121]
// is read directly from the encoder INPUT buffer (`batch_states`), not
// from the encoder output, so this bias bypasses the encoder's slow
// weight initialization for slot 121. The encoder STILL gets gradient
// via option-(i) routing (see spec "Encoder-state[121] gradient
// routing"); alpha is a parallel skip connection.
//
// Discipline
// ----------
// - No atomicAdd per `feedback_no_atomicadd.md` (no reduction; just
// per-(b, a) in-place add).
// - No host branches per `pearl_no_host_branches_in_captured_graph` --
// safe for capture inside training + rollout forward graphs.
// - Pure GPU compute per `feedback_no_cpu_compute_strict.md`.
// - Stream-ordered with the producer that wrote Q_dir (Q-head forward
// completion); same-stream barrier suffices.
//
// Launch
// ------
// Grid: ((B * b0_size + 255) / 256, 1, 1) -- 1D thread mapping with
// `tid = b * b0_size + a`. Block: (256, 1, 1). Each thread handles one
// (b, a) pair.
//
// Cost per launch: one f32 load (state_121[b]), one f32 load (W[a]),
// one fused multiply-add, one f32 store (Q_dir[b, a]). At B=8192,
// b0_size=4 -> 32768 threads -> ~128 warps -> negligible (~us).
#include <cuda_runtime.h>
extern "C" __global__ void aux_to_q_dir_bias(
/* Q_dir buffer to modify in-place. Layout [B, b0_size]. */
float* __restrict__ q_dir,
/* alpha weight vector [b0_size]. Trainer-owned; collector + eval
* read via shared device pointer (same source-of-truth pattern as
* the network's GEMM weights). */
const float* __restrict__ w_aux,
/* Encoder INPUT buffer [B, state_dim]. The kernel reads
* `batch_states[b * state_dim + aux_dir_prob_index]` -- specifically
* the encoder-input slot 121 (Phase 2 recentered to [-1, +1]),
* NOT the encoder output. This is what makes alpha a bypass. */
const float* __restrict__ batch_states,
/* AUX_DIR_PROB_INDEX (= SL_PADDING_START + 0 = 121 in the canonical
* Phase-2 layout). Passed runtime so the kernel ABI tolerates
* future state-layout shifts. */
int aux_dir_prob_index,
/* State dim (canonical 128 under Phase 1+2). Used as the row
* stride for batch_states[b * state_dim + ...]. */
int state_dim,
/* Number of samples in the batch. */
int batch_size,
/* Number of direction actions (= b0_size = 4 under the 4-3-3-3
* factored space). */
int b0_size
) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int total = batch_size * b0_size;
if (tid >= total) return;
int b = tid / b0_size;
int a = tid - b * b0_size;
/* Encoder-input slot 121 (Phase 2 recentered: `2*p_up - 1` in
* [-1, +1]; 0 = neutral; +1 = up with full conviction; -1 = down
* with full conviction). */
float state_121 = batch_states[(long long)b * state_dim + aux_dir_prob_index];
/* Per-action bias contribution. W_aux is small (4 elements);
* Adam-trained from cold-start zero. */
float bias = w_aux[a] * state_121;
/* In-place add to Q_dir. The Q-head forward upstream (cublas SGEMM
* + activation) wrote q_dir[b, a]; we add the aux bias here,
* before action selection / TD computation downstream. */
q_dir[(long long)b * b0_size + a] += bias;
}

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

@@ -2634,7 +2634,7 @@ const ISV_NETWORK_DIM: usize = 23;
/// (and the C-side mirrors in `state_layout.cuh`). Bump this constant in the
/// SAME commit that adds the new range, and update `layout_fingerprint_seed`
/// to register the new slot names.
pub(crate) const ISV_TOTAL_DIM: usize = 536; // SP21 Phase 7.5 adds 8 slots [528..536) for CURRICULUM_WEIGHT_{0..8} (per_insert_pa per-segment priority boost); Phase 7 added 1 slot [527..528); Phase 5+6 added 2 slots [525..527); Phase 4 added 4 slots [521..525); Phase 3 added 1 slot [520..521)
pub(crate) const ISV_TOTAL_DIM: usize = 538; // SP22 H6 Phase 3 (2026-05-13) adds 2 slots [536..538) for the aux→policy bypass: REWARD_AUX_ALIGN_EMA_INDEX (7th reward component EMA, anchor for SP11 controller's w_aux_align) and SP22_AUX_ALIGN_SCALE_INDEX (controller-emitted adaptive scale_β consumed by β producers at training and eval). SP21 Phase 7.5 added 8 slots [528..536) for CURRICULUM_WEIGHT_{0..8} (per_insert_pa per-segment priority boost); Phase 7 added 1 slot [527..528); Phase 5+6 added 2 slots [525..527); Phase 4 added 4 slots [521..525); Phase 3 added 1 slot [520..521)
/// Legacy alias preserved for call sites that haven't been audited for the
/// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight
/// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer).
@@ -3776,7 +3776,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
SLOT_533_CURRICULUM_WEIGHT_5=533;\
SLOT_534_CURRICULUM_WEIGHT_6=534;\
SLOT_535_CURRICULUM_WEIGHT_7=535;\
ISV_TOTAL_DIM=536;\
SLOT_536_REWARD_AUX_ALIGN_EMA=536;\
SLOT_537_SP22_AUX_ALIGN_SCALE=537;\
ISV_TOTAL_DIM=538;\
SP19_PRODUCER_HARDCODED_HORIZON_BLEND=sp19_path_b;\
SP14_C_AUX_TRUNK_CONTROL_PLANE=sp14_c_phase_1;\
ALPHA_MACHINERY_DELETED=sp14_c_phase_1;\

1
crates/ml/src/cuda_pipeline/mod.rs
View File

@@ -66,6 +66,7 @@ pub mod sp13_isv_slots;
pub mod sp14_isv_slots;
pub mod sp15_isv_slots;
pub mod sp21_isv_slots;
pub mod sp22_isv_slots;
pub mod lob_bar;
pub use sp4_isv_slots::{
TARGET_Q_BOUND_INDEX, ATOM_POS_BOUND_BASE, WEIGHT_BOUND_BASE,

70
crates/ml/src/cuda_pipeline/sp22_isv_slots.rs Normal file
View File

@@ -0,0 +1,70 @@
//! SP22 ISV slot reservations.
//!
//! SP22 H6 Phase 3 (2026-05-13) introduces the aux→policy bypass routing —
//! α (post-encoder Q_dir bias) and β (event-driven aux-aligned trade-close
//! bonus). β requires two new ISV slots: one for the EMA of the new 7th
//! reward component, and one for the SP11-controller-emitted adaptive
//! scale_β.
//!
//! ## Slot allocation
//!
//! | Index | Name | Producer | Consumer |
//! |-------|------|----------|----------|
//! | 536 | `REWARD_AUX_ALIGN_EMA_INDEX` | `reward_component_ema_kernel` (extended 0..7 iteration) | SP11 `reward_subsystem_controller_kernel` (anchor for `w_aux_align` adaptive weight emit); HEALTH_DIAG `reward_split` printer |
//! | 537 | `SP22_AUX_ALIGN_SCALE_INDEX` | SP11 `reward_subsystem_controller_kernel` (extended 7-weight output: anchor on `REWARD_AUX_ALIGN_EMA_INDEX`, target ~10% of total reward magnitude, bounds [0.05, 2.0]) | `experience_env_step` (training-side β producer at segment_complete); `backtest_env_kernel` (eval-side β producer); HEALTH_DIAG `sp11_reward` printer (`w_aux` column) |
//!
//! ## Bus extension
//!
//! These two slots bump `ISV_TOTAL_DIM` from 536 → 538 in
//! `gpu_dqn_trainer.rs`. They also extend the layout fingerprint emitted
//! by the slot-metadata format string (also in `gpu_dqn_trainer.rs`).
//! All three changes (this file, ISV_TOTAL_DIM, fingerprint) move
//! atomically per `feedback_no_partial_refactor`.
//!
//! ## Cold-start sentinel discipline
//!
//! Both slots cold-start at sentinel 0 per `pearl_first_observation_bootstrap`.
//! For `SP22_AUX_ALIGN_SCALE_INDEX`, the cold-start zero makes β a no-op
//! until the SP11 controller's first emit (sometime in the first epoch
//! after the EMA producer has accumulated a non-zero signal). For
//! `REWARD_AUX_ALIGN_EMA_INDEX`, the producer's first observation replaces
//! the sentinel directly (Pearl A bootstrap).
//!
//! Registered in `state_reset_registry.rs` as FoldReset category so both
//! slots reset to 0 at fold boundaries.
/// SP22 H6 Phase 3 (2026-05-13) — EMA of the 7th reward component
/// `r_aux_align`. Produced by the extended (0..7 iteration)
/// `reward_component_ema_kernel`. Consumed by:
///
/// - `reward_subsystem_controller_kernel` (anchor signal for the
/// `w_aux_align` adaptive weight emit at index `SP22_AUX_ALIGN_SCALE_INDEX`).
/// - HEALTH_DIAG `reward_split` printer (the `aux_align=…` column).
/// - `health_diag_kernel` snap layout (`snap.reward_split_aux_align`).
///
/// Cold-start sentinel 0 per `pearl_first_observation_bootstrap`; first
/// non-zero observation replaces directly. FoldReset.
pub const REWARD_AUX_ALIGN_EMA_INDEX: usize = 536;
/// SP22 H6 Phase 3 (2026-05-13) — SP11-controller-emitted adaptive
/// `scale_β` for the β producer at trade-close.
///
/// Producer: extended `reward_subsystem_controller_kernel` (7-weight
/// output). The new 7th weight is anchored on `REWARD_AUX_ALIGN_EMA_INDEX`,
/// targets `0.10 × Σ_c |reward_components_ema[c]|` so β contributes
/// ~10% of total reward signal at steady state, and is bounded to
/// `[0.05, 2.0]` (floor prevents β vanishing; ceiling prevents
/// β dominating per `pearl_one_unbounded_signal_per_reward`'s
/// component-balancing intent).
///
/// Consumers:
///
/// - `experience_env_step::segment_complete` (training-side β producer).
/// - `backtest_env_kernel::segment_complete` (eval-side β producer).
/// - HEALTH_DIAG `sp11_reward` printer (the `w_aux=…` column).
///
/// Cold-start sentinel 0 per `pearl_first_observation_bootstrap` → β
/// is a no-op until the controller's first non-zero emit (typically
/// within the first epoch after the EMA producer has accumulated a
/// non-zero signal). FoldReset.
pub const SP22_AUX_ALIGN_SCALE_INDEX: usize = 537;

14
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -2152,6 +2152,20 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "GpuExperienceCollector.hold_baseline_buffer [alloc_episodes × HOLD_BASELINE_BUFFER_SIZE = 30] f32 device-resident scratch — SP20 Phase 3 Task 3.2 (2026-05-10) per-env circular Hold opportunity-cost buffer per spec §4.2 (Component 2). Layout: row-major, `[env, current_t % 30]` slot per per-bar write. Producer: `experience_env_step` per-bar branches (positioned-non-event ~line 3650 + flat-non-event ~line 3737) — both write `hold_baseline_buffer[env, current_t % 30] = -aux_conf × cost_scale` UNCONDITIONALLY (Path 2 counterfactual emission, regardless of action; spec §4.2 line 225 `hold_baseline_buffer[i] = R_per_bar_i // uncentered, for Component 1`). Consumer: `experience_env_step` `is_close && segment_hold_time > 0.0f` branch — `hold_baseline = sp20_sum_hold_baseline_over_trade(buf, env, 30, current_t, segment_hold_time)` walking backwards `min(30, segment_hold_time)` slots from `(current_t - 1) % 30`. Replaces the Phase 2 Task 2.2 placeholder `hold_baseline = 0.0f` atomically per `feedback_no_partial_refactor`. FoldReset sentinel 0.0 across all `alloc_episodes × 30` slots — leftover per-bar costs from the previous fold's open trade would corrupt the new fold's first trade-close summation (the new fold's portfolio_states is also FoldReset so any trade open at fold boundary is impossible, but if a buffer slot retained a stale value it could leak into the first trade-close after fold reset via the modulo-arithmetic indexing). Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice<f32>` (device-resident, no host buffer to fill). Per `pearl_first_observation_bootstrap`: cold-start with `ISV[HOLD_COST_SCALE_INDEX=513]` at sentinel 0 means `per_bar_opp_cost = 0` (Path 2 writes zeros to the buffer); the trade-close consumer sums zeros and lands `hold_baseline = 0.0f` (matches Phase 2 Task 2.2 placeholder behaviour bit-for-bit on the very first trade close after fold reset). The controller's first observation lifts cost_scale, then per-bar writes become non-zero, then the next trade-close summation engages — gracefully bootstrapped, no warmup gate needed.",
},
// SP22 H6 Phase 3 (2026-05-13) — aux→policy bypass routing.
// Two new ISV slots backing β: the 7th reward component EMA and
// the SP11-controller-emitted adaptive scale_β. Both FoldReset
// sentinel 0 per `pearl_first_observation_bootstrap`.
RegistryEntry {
name: "sp22_reward_aux_align_ema",
category: ResetCategory::FoldReset,
description: "ISV[REWARD_AUX_ALIGN_EMA_INDEX=536] — SP22 H6 Phase 3 (2026-05-13) EMA of the 7th reward component `r_aux_align` (β producer at trade-close). Producer: extended `reward_component_ema_kernel.cu` (0..7 iteration after the 7-component contract migration). Consumers: `reward_subsystem_controller_kernel.cu` (anchor signal for the new `w_aux_align` adaptive weight at ISV[SP22_AUX_ALIGN_SCALE_INDEX=537]); HEALTH_DIAG `reward_split` printer (the `aux_align=…` column); `health_diag_kernel.cu` snap layout (`snap.reward_split_aux_align`). FoldReset sentinel 0 — Pearl A bootstrap: first non-zero observation from the EMA producer replaces the sentinel directly. The β producer fires at `segment_complete` only when (a) policy direction aligns with aux conviction AND (b) trade was profitable, so cold-start with no aligned-profitable trades yet means the EMA stays at 0 until the first such trade closes; this is the correct cold-start behavior — no warmup gate needed.",
},
RegistryEntry {
name: "sp22_aux_align_scale",
category: ResetCategory::FoldReset,
description: "ISV[SP22_AUX_ALIGN_SCALE_INDEX=537] — SP22 H6 Phase 3 (2026-05-13) SP11-controller-emitted adaptive scale_β for the β producer at trade-close. Producer: extended `reward_subsystem_controller_kernel.cu` (7-weight output, anchor on `REWARD_AUX_ALIGN_EMA_INDEX=536`, target ~10% of total reward magnitude, bounds [0.05, 2.0]). Consumers: `experience_env_step::segment_complete` (training-side β: `r_aux_align = scale_β × alignment × profit_pos`) and `backtest_env_kernel::segment_complete` (eval-side β with the same formula); HEALTH_DIAG `sp11_reward` printer (the `w_aux=…` column). FoldReset sentinel 0 — when scale_β = 0, β is a no-op (no reward signal added), which is the correct behavior pre-first-controller-emit. The controller's first emit replaces the sentinel directly per `pearl_first_observation_bootstrap`; subsequent EMA-driven adaptation tracks the anchor signal within the bounds.",
},
];
Self { entries }
}

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

@@ -16667,3 +16667,79 @@ combining into one atomic plan.
If combined α+β fails: amplitude scaling or per-dim adaptive LR are
fallbacks; the action-distribution bit-identity finding suggests the
problem is structural, not numerical.
### Phase 3 Phase A foundation (2026-05-13) — WIP checkpoint, additive
Phase A of the H6 Phase 3 implementation runbook
(`docs/plans/2026-05-13-sp22-h6-phase3-alpha-beta-runbook.md`) lands
as a purely additive foundation. **No consumer wires the new
infrastructure yet** — the existing 6-component reward decomposition
contract is unchanged. Phase A is committable as a clean checkpoint;
Phases B–F (7-component contract migration, α plumbing, A2 eval-side
wiring, verification gates, atomic-commit + smoke + verdict) resume
in a future session.
Phase A files (7):
- `crates/ml/src/cuda_pipeline/sp22_isv_slots.rs` (new) —
`REWARD_AUX_ALIGN_EMA_INDEX = 536`, `SP22_AUX_ALIGN_SCALE_INDEX = 537`.
- `crates/ml/src/cuda_pipeline/mod.rs` — `pub mod sp22_isv_slots`.
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — `ISV_TOTAL_DIM`
bumped 536 → 538; layout fingerprint extended with the two new
`SLOT_536_REWARD_AUX_ALIGN_EMA` + `SLOT_537_SP22_AUX_ALIGN_SCALE`
entries.
- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — two new
`RegistryEntry` rows (FoldReset, sentinel 0).
- `crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu` (new) —
α forward: `Q_dir[b, a] += W_aux[a] * state_121[b]`. Reads
`state_121` from the encoder INPUT (`batch_states`), bypassing the
cold encoder weight for slot 121.
- `crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu`
(new) — two kernels in one `.cu`: `aux_to_q_dir_bias_backward_dw`
(per-action block tree-reduce; no atomicAdd) and
`aux_to_q_dir_bias_backward_dstate` (per-sample sum accumulating
into encoder-input gradient slot 121, option-(i) gradient routing
per spec).
- `crates/ml/build.rs` — both new cubins registered in
`kernels_with_common`.
Verification (Phase A, partial — final gates run at Phase F):
| Gate | Result |
|---|---|
| `cargo check -p ml --features cuda` | 0 errors, 21 pre-existing warnings |
| nvcc cubins | both compile cleanly (3.6 KB fwd, 10.6 KB bwd) |
| No new consumer wiring | confirmed; running code identical to pre-Phase-A |
Remaining work for Phase 3 verdict (~28–43 hr engineering):
- **Phase B** (11 tasks): 7-component contract migration cascade
across `experience_kernels.cu`, `reward_component_ema_kernel.cu`,
`reward_decomp_diag_kernel.cu`,
`reward_component_mag_ratio_compute_kernel.cu`,
`reward_subsystem_controller_kernel.cu`, `health_diag_kernel.cu`,
`gpu_experience_collector.rs`, `gpu_dqn_trainer.rs`,
`reward_component_monitor.rs`, `training_loop.rs`.
- **Phase C** (1 task): α collector rollout-side captured-graph
wiring.
- **Phase D** (7 tasks): A2 eval-side — aux trunk weight loading,
per-window `prev_aux_dir_prob` buffer, per-step aux trunk forward
+ softmax-to-per-env launch, α post-Q_dir bias launch, non-NULL
`prev_aux_dir_prob.raw_ptr()` to all 3 backtest state-gather
launchers, W/aux-trunk ptr sharing trainer ↔ evaluator.
- **Phase E** (7 tasks): verification gates (cargo check,
gpu_backtest_validation, compute-sanitizer, captured-graph
integrity, α weight movement, 7-component contract sanity, A2
eval activity).
- **Phase F** (5 tasks): audit doc append (Phase B–F outcomes),
single atomic commit covering all Phase B–F changes (per
`feedback_no_partial_refactor`), push, smoke dispatch
(`./scripts/argo-train.sh dqn --baseline --branch sp20-aux-h-fixed
--epochs 3 --gpu-pool ci-training-l40s`), verdict monitoring.
Phase B partial work (file diffs for `experience_kernels.cu` +
`reward_component_ema_kernel.cu` — 7-component stride migration +
β producer + EMA kernel preamble update) saved at
`/tmp/sp22-h6-phase3-b-partial.patch` (300 lines, local-only). Next
session can apply if useful as a starting point, or discard and
re-derive directly from the runbook.