diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 4eee7ac42..31cf224f3 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -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 diff --git a/crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu b/crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu new file mode 100644 index 000000000..dd6a0ea51 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu @@ -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 + +#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; +} diff --git a/crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu b/crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu new file mode 100644 index 000000000..d2da6177d --- /dev/null +++ b/crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu @@ -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 + +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; +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 499247d6a..77d0883ae 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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;\ diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 07ab428cb..bd3382311 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -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, diff --git a/crates/ml/src/cuda_pipeline/sp22_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp22_isv_slots.rs new file mode 100644 index 000000000..682cd3040 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp22_isv_slots.rs @@ -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; diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 870de5c46..7014abc8b 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -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` (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 } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 61419dd3b..f50002aeb 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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.