diff --git a/crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu b/crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu index ce12b7f8a..235f7c2bd 100644 --- a/crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu +++ b/crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu @@ -18,8 +18,15 @@ // ── Driver inputs (read-only on this kernel) ──────────────────────────── // ISV[TARGET_DIR_ACC = 372] target directional accuracy (SP13) // ISV[AUX_DIR_ACC_SHORT_EMA = 373] fast EMA of aux dir-acc (SP13) +// ISV[AUX_DIR_ACC_LONG_EMA = 374] slow EMA of aux dir-acc (SP13) — used +// as the running mean for the var_aux +// Welford EMA (SP14 B.11 var_aux producer) // ISV[Q_DISAGREEMENT_SHORT_EMA = 383] fast EMA of Q-aux disagreement (SP14 B.3) // ISV[VAR_AUX = 388] Welford variance EMA of aux dir-acc +// (SP14 B.11: this kernel now produces +// the var_aux update — was a ghost +// feature in B.4, see "RESOLVED" block +// below) // ISV[VAR_Q = 389] Welford variance EMA of Q disagreement (SP14 B.3) // ISV[GATE1_OPEN = 391] persistent Schmitt state (read+written) // ISV[VAR_ALPHA = 390] Welford variance EMA of α_grad_raw (read+written) @@ -29,6 +36,7 @@ // ISV[K_AUX_ADAPTIVE = 385] adaptive Gate-1 sigmoid steepness // ISV[K_Q_ADAPTIVE = 386] adaptive Gate-2 sigmoid steepness // ISV[BETA_ADAPTIVE = 387] adaptive β rate-limiter coefficient +// ISV[VAR_AUX = 388] updated aux_dir_acc variance EMA (SP14 B.11) // ISV[VAR_ALPHA = 390] updated α_grad_raw variance EMA // ISV[GATE1_OPEN = 391] updated Schmitt state (0 = closed, 1 = open) // ISV[ALPHA_RAW = 392] α_grad_raw = gate1 × gate2 × warmup_gate @@ -69,27 +77,31 @@ // overflows). The clip prevents NaN propagation from pathological // k_aux × (aux_short - threshold) products if k_aux is corrupted. // -// ── KNOWN LIMITATION: var_aux producer not yet wired ──────────────────── -// This kernel READS ISV[388] (AUX_DIR_ACC_VARIANCE_EMA) but does NOT -// write it. As of B.4 landing, NO upstream kernel writes slot 388 — -// `sp14_isv_slots.rs` is the only file referencing that constant. The -// effect: var_aux stays at sentinel 0.0 forever, so -// k_aux = max(K_BASE_AUX / (1 + 0/VARIANCE_REF_AUX), K_MIN) -// = max(K_BASE_AUX, K_MIN) = K_BASE_AUX (constant). -// The adaptive-k_aux mechanism is degenerate-but-non-fatal: Gate 1 still -// works, the sigmoid just doesn't soften under noisy aux_dir_acc. This -// is to be resolved in the B.11 producer chain orchestrator OR a -// separate fix-up task that adds a Welford-variance update next to the -// existing AUX_DIR_ACC_SHORT_EMA producer. Filed in the B.4 status report. +// ── RESOLVED: var_aux producer wired in B.11 (2026-05-05) ─────────────── +// B.4 originally READ ISV[388] (AUX_DIR_ACC_VARIANCE_EMA) but did NOT +// write it, so var_aux stayed at sentinel 0.0 and k_aux degenerated to +// the constant K_BASE_AUX. B.11 closes the gap inline: the kernel now +// also computes a Welford-style variance EMA on `aux_dir_acc_short` with +// `aux_dir_acc_long` as the running mean (mirrors how var_q in +// `q_disagreement_update_kernel` uses the fast EMA `new_short` as its +// mean reference). The compute lives next to the existing var_alpha +// update (same Welford pattern, same `alpha_var` blend coefficient) so +// the kernel stays a single state-machine launch with no companion +// kernel needed. Per `feedback_no_partial_refactor.md`, the kernel +// signature is unchanged; we add reads of ISV[374] and a write to +// ISV[388] inside the existing dispatch. // // var_q (slot 389) IS written — by `q_disagreement_update_kernel` (B.3), // which lands the Welford variance EMA in the same launch as the mean -// EMAs. So adaptive k_q is fully functional from B.4 onward. +// EMAs. With var_aux now produced too, both adaptive k_aux and k_q are +// fully functional from B.11 onward. extern "C" __global__ void alpha_grad_compute_kernel( - /* Global ISV bus. Slots 385, 386, 387, 390, 391, 392, 393 are written; - * slots 372, 373, 383, 388, 389 are read; all other slots untouched. + /* Global ISV bus. Slots 385, 386, 387, 388, 390, 391, 392, 393 are + * written; slots 372, 373, 374, 383, 388, 389 are read (slot 388 is + * read-modify-write per the Welford pattern; slot 374 added in B.11 + * as the var_aux mean reference). All other slots untouched. */ float* __restrict__ isv, /* Per-epoch warmup ramp value in [0, 1]. Host-computed (B.7+ wires @@ -109,6 +121,7 @@ void alpha_grad_compute_kernel( // ── ISV slot indices (must match sp14_isv_slots.rs and sp13_isv_slots.rs) ── const int TARGET_DIR_ACC = 372; const int AUX_DIR_ACC_SHORT_EMA = 373; + const int AUX_DIR_ACC_LONG_EMA = 374; // SP14 B.11: var_aux mean reference const int Q_DISAGREEMENT_SHORT_EMA = 383; const int K_AUX_ADAPTIVE = 385; const int K_Q_ADAPTIVE = 386; @@ -135,17 +148,34 @@ void alpha_grad_compute_kernel( // ── Read drivers ─────────────────────────────────────────────────── const float target = isv[TARGET_DIR_ACC]; const float aux_short = isv[AUX_DIR_ACC_SHORT_EMA]; + const float aux_long = isv[AUX_DIR_ACC_LONG_EMA]; // SP14 B.11 const float q_dis = isv[Q_DISAGREEMENT_SHORT_EMA]; - const float var_aux = isv[VAR_AUX]; + const float var_aux_prev = isv[VAR_AUX]; // SP14 B.11 const float var_q = isv[VAR_Q]; const float var_alpha_prev = isv[VAR_ALPHA]; const float gate1_state_prev = isv[GATE1_OPEN]; const float alpha_smoothed_prev = isv[ALPHA_SMOOTHED]; + // ── Welford-style variance EMA on aux_dir_acc_short (SP14 B.11) ──── + // Reference is `aux_long` (slow EMA — best estimate of the running + // mean). Under stable aux dir-acc the diff stays small and var_aux + // decays toward 0; under chatter the diff grows and var_aux rises, + // softening the Gate 1 sigmoid (k_aux below) so noisy aux signals + // don't whipsaw the EGF gate. Uses the same `alpha_var` blend + // coefficient as the var_alpha update further down — single shared + // smoothing rate per Pearl-D Wiener-α convention. Closes the + // ghost-feature gap left by B.4 (slot 388 had no producer; pre-B.11 + // the adaptive k_aux mechanism was degenerate at K_BASE_AUX = 20). + const float diff_aux = aux_short - aux_long; + const float var_aux = alpha_var * (diff_aux * diff_aux) + + (1.0f - alpha_var) * var_aux_prev; + // ── Adaptive sigmoid steepness (B.2.5) ───────────────────────────── // Higher variance → smaller k → flatter sigmoid (more diffusion at // the edge). Floor at K_MIN prevents the sigmoid from collapsing to - // a flat 0.5 line under unbounded noise. + // a flat 0.5 line under unbounded noise. `var_aux` is the FRESHLY- + // updated value above (this step's contribution included), so k_aux + // tracks current-step noise rather than lagging by one launch. const float k_aux = fmaxf(K_BASE_AUX / (1.0f + var_aux / VARIANCE_REF_AUX), K_MIN); const float k_q = fmaxf(K_BASE_Q / (1.0f + var_q / VARIANCE_REF_Q), K_MIN); @@ -219,6 +249,7 @@ void alpha_grad_compute_kernel( isv[K_AUX_ADAPTIVE] = k_aux; isv[K_Q_ADAPTIVE] = k_q; isv[BETA_ADAPTIVE] = beta; + isv[VAR_AUX] = var_aux; // SP14 B.11: ghost-feature gap closed isv[VAR_ALPHA] = var_alpha_new; isv[GATE1_OPEN] = gate1_state_new; isv[ALPHA_RAW] = alpha_raw; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 38a111cf3..c3db28e17 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -7185,6 +7185,144 @@ impl GpuDqnTrainer { Ok(()) } + /// SP14 Layer B Task B.11 (2026-05-05): per-step launcher for the + /// EGF q_disagreement producer. Reads `aux_nb_softmax_buf [B, K=2]` + /// (populated by the captured aux-head forward) and `q_out_buf + /// [B, total_actions=13]` (populated by `populate_q_out` after the + /// captured forward), runs the K=4↔K=2 mapped argmax-mismatch + /// reduction, and updates `ISV[Q_DISAGREEMENT_SHORT_EMA=383]`, + /// `ISV[Q_DISAGREEMENT_LONG_EMA=384]`, `ISV[Q_DISAGREEMENT_VARIANCE_EMA=389]`. + /// + /// Stride `total_actions=13` lets the kernel read the FIRST 4 columns + /// of each row (direction Q-values) directly from `q_out_buf` without + /// a stride-stripping copy. The kernel's K=4 loop bound is + /// independent of stride. + /// + /// Launch contract (matches `q_disagreement_update_kernel.cu` header): + /// single-block, 256 threads, dynamic shmem = `2 × 256 × sizeof(f32) = 2048` + /// (numerator + count tiles, see kernel header). + /// + /// Producer-only — writes ISV slots; consumer is the EGF + /// alpha_grad_compute_kernel (B.4), which reads slot 383 (and 389 + /// via var_q for adaptive k_q). + pub(crate) fn launch_sp14_q_disagreement_update( + &self, + alpha_short: f32, + alpha_long: f32, + alpha_var: f32, + ) -> Result<(), MLError> { + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_sp14_q_disagreement_update: isv_signals_dev_ptr must be allocated"); + let b_i32 = self.config.batch_size as i32; + let q_stride = self.total_actions() as i32; + let aux_softmax_ptr = self.aux_nb_softmax_buf.raw_ptr(); + let q_out_ptr = self.q_out_buf.raw_ptr(); + let isv_ptr = self.isv_signals_dev_ptr; + // Block size 256 → smem = 2 × 256 × 4 = 2048 bytes (matches the + // kernel's `extern __shared__ float smem[]` split into num_smem + + // cnt_smem). Single block; the kernel's strided loop covers any + // batch size on a single block. + let block_dim: u32 = 256; + let shmem_bytes: u32 = 2 * block_dim * std::mem::size_of::() as u32; + unsafe { + self.stream + .launch_builder(&self.sp14_q_disagreement_update_kernel) + .arg(&aux_softmax_ptr) + .arg(&q_out_ptr) + .arg(&isv_ptr) + .arg(&b_i32) + .arg(&q_stride) + .arg(&alpha_short) + .arg(&alpha_long) + .arg(&alpha_var) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (block_dim, 1, 1), + shared_mem_bytes: shmem_bytes, + }) + .map_err(|e| MLError::ModelError(format!("sp14 q_disagreement: {e}")))?; + } + Ok(()) + } + + /// SP14 Layer B Task B.11 (2026-05-05): per-step launcher for the + /// EGF α_grad consumer kernel. Reads driver signals from ISV + /// (slots 372/373/374/383/388/389/390/391/393), runs the Schmitt- + /// trigger Gate 1 + adaptive k + adaptive β state machine, writes + /// ALPHA_GRAD_RAW (392), ALPHA_GRAD_SMOOTHED (393), and the + /// supporting bookkeeping slots (385–391). + /// + /// MUST run AFTER `launch_sp14_q_disagreement_update` (consumes its + /// var_q output via slot 389) and AFTER `launch_sp13_aux_dir_metrics` + /// (consumes slot 373 / 374). The launcher submits onto the trainer + /// stream — same-stream submission ordering enforces these deps. + /// + /// `warmup_gate ∈ [0, 1]` is the host-supplied per-epoch ramp factor + /// (see plan §2581: `clamp(steps_in_fold / WARMUP_STEPS, 0, 1)`). + /// `alpha_var` is the Welford-EMA blend coefficient for both + /// var_alpha (slot 390) and var_aux (slot 388, B.11 producer). + /// + /// Single-thread state machine — kernel header guarantees blockDim + /// ≥ (1,1,1) suffices; we launch with 32 threads to mirror the B.4 + /// oracle test's launch shape. + pub(crate) fn launch_sp14_alpha_grad_compute( + &self, + warmup_gate: f32, + alpha_var: f32, + ) -> Result<(), MLError> { + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_sp14_alpha_grad_compute: isv_signals_dev_ptr must be allocated"); + let isv_ptr = self.isv_signals_dev_ptr; + unsafe { + self.stream + .launch_builder(&self.sp14_alpha_grad_compute_kernel) + .arg(&isv_ptr) + .arg(&warmup_gate) + .arg(&alpha_var) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("sp14 alpha_grad: {e}")))?; + } + Ok(()) + } + + /// SP14 Layer B Task B.11 (2026-05-05): per-epoch launcher for the + /// anti-mesa-optimization circuit breaker. Decrements the lockout + /// counter (ISV[395]), tracks the post-open minimum of aux_dir_acc + /// (ISV[394]), and forces Gate 1 closed (ISV[391] := 0) when a + /// hacking signature is detected (aux_dir_acc drops > 0.05 below + /// the Schmitt open-threshold AND q_disagreement rises > 0.10 above + /// the analytic baseline simultaneously). + /// + /// Single-thread state machine — kernel header guarantees blockDim + /// ≥ (1,1,1) suffices; we launch with 32 threads to mirror the B.5 + /// oracle test's launch shape. + /// + /// Caller fires this once per epoch end (e.g. inside + /// `process_epoch_boundary`), AFTER the per-step alpha_grad chain + /// has run for all training steps in the epoch — the lockout + /// decrement is one-per-epoch by design (`LOCKOUT_EPOCHS = 2.0`). + pub(crate) fn launch_sp14_gradient_hack_detect(&self) -> Result<(), MLError> { + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_sp14_gradient_hack_detect: isv_signals_dev_ptr must be allocated"); + let isv_ptr = self.isv_signals_dev_ptr; + unsafe { + self.stream + .launch_builder(&self.sp14_gradient_hack_detect_kernel) + .arg(&isv_ptr) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("sp14 grad_hack: {e}")))?; + } + Ok(()) + } + /// Build OFI concat for order (d=2) or urgency (d=3) branch. /// /// Reads vsn_masked[B, SH2] and raw feature vector states[B, SD], appends 3 OFI diff --git a/crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu b/crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu index 10619eaaf..e69ee779c 100644 --- a/crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu +++ b/crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu @@ -83,10 +83,19 @@ void q_disagreement_update_kernel( * ops on a single stream. */ const float* __restrict__ aux_softmax, - /* Per-bar K=4 Q-head logits/values [B, 4]. Row-major; argmax over - * the 4 dimensions identifies the Q-picked direction. The kernel - * does NOT need calibrated probabilities — only the argmax matters - * for the disagreement rate. + /* Per-bar Q-head values. Row-major with stride `q_stride`; the + * kernel reads the FIRST 4 columns per row (direction Q-values) + * via `q_logits[b * q_stride + k]` for k in 0..4 and argmaxes + * over them. Stride accommodates either a tight [B, K_DIR=4] + * direction-only buffer (q_stride = 4) or the full + * [B, total_actions=13] q_out_buf (q_stride = 13). The kernel does + * NOT need calibrated probabilities — only the argmax matters for + * the disagreement rate. + * + * SP14 B.11 (2026-05-05): added `q_stride` parameter so the + * trainer can pass q_out_buf [B, 13] directly without a stride- + * stripping copy. K_DIR (the loop bound) stays 4 — the wider + * stride just skips the magnitude/order/urgency tail per row. */ const float* __restrict__ q_logits, /* Global ISV bus. Slots 383, 384, 389 read+written; all other @@ -100,6 +109,10 @@ void q_disagreement_update_kernel( * on cold start with empty batch). */ int B, + /* Stride between rows in `q_logits` (>= K_DIR=4). Direction-only + * tight buffer uses 4; full q_out_buf uses 13. SP14 B.11. + */ + int q_stride, /* Fast-EMA blend coefficient (typical 0.3 — ~3-step half-life). */ float alpha_short, /* Slow-EMA blend coefficient (typical 0.05 — ~14-step half-life; @@ -127,16 +140,17 @@ void q_disagreement_update_kernel( const float as1 = aux_softmax[b * K_AUX + 1]; const int aux_class = (as1 > as0) ? 1 : 0; - // argmax(q_logits[b]) over K=4. Strict `>` tie-break (matches + // argmax(q_logits[b, 0..K_DIR]) over the first 4 columns of + // each row (stride = q_stride). Strict `>` tie-break (matches // the C-comparator semantics used elsewhere in the codebase // — first-class-wins on ties, no special handling needed // because the production logits are ~never bit-equal across // classes after a forward pass). int q_class = 0; - float q_max = q_logits[b * K_DIR + 0]; + float q_max = q_logits[b * q_stride + 0]; #pragma unroll for (int k = 1; k < K_DIR; ++k) { - const float v = q_logits[b * K_DIR + k]; + const float v = q_logits[b * q_stride + k]; if (v > q_max) { q_max = v; q_class = k; } } diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs index 5377a5333..84dd7011e 100644 --- a/crates/ml/src/trainers/dqn/trainer/constructor.rs +++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs @@ -608,6 +608,9 @@ impl DQNTrainer { // Wave 16 Portfolio Features (action masking always active) max_position, current_epoch: 0, + // SP14 B.11 (2026-05-05): EGF warmup-gate counter; reset by + // reset_for_fold, incremented per-step in run_training_steps_slices. + fold_step_counter: 0, multi_asset_portfolio, stress_tester, diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index 35b157e93..71489d7d6 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/crates/ml/src/trainers/dqn/trainer/mod.rs @@ -435,6 +435,20 @@ pub struct DQNTrainer { pub max_position: f64, /// Current epoch (for Q-gap warm-up ramp) pub(crate) current_epoch: usize, + + /// SP14 B.11 (2026-05-05): per-fold cumulative training-step counter + /// driving the EGF α_grad warmup gate. Reset to 0 in `reset_for_fold`; + /// incremented every iteration of `run_training_steps_slices`'s + /// per-step body so it accumulates across epochs within a fold. + /// `warmup_gate = clamp(fold_step_counter / WARMUP_STEPS_FALLBACK, 0, 1)` is + /// passed to `launch_sp14_alpha_grad_compute` per + /// `crate::cuda_pipeline::sp14_isv_slots::WARMUP_STEPS_FALLBACK`. The + /// gate ramps from 0 (cold-start, no EGF gradient flow) to 1 (warmup + /// complete, gates fully control flow) over `WARMUP_STEPS_FALLBACK` + /// per-step launches. Per `feedback_no_partial_refactor`, the trainer + /// owns this counter rather than fishing it out of the training-loop + /// async closure scope. + pub(crate) fold_step_counter: usize, /// Entropy regularizer for preventing policy collapse (None if disabled) // entropy_regularizer removed — SAC-style entropy is computed directly on Q-value tensors in DQN::compute_loss_internal /// Multi-asset portfolio tracker (None if single-asset mode) @@ -1773,6 +1787,13 @@ impl DQNTrainer { // Reset epoch counters self.current_epoch = 0; self.gradient_logging_step = 0; + // SP14 B.11 (2026-05-05): EGF warmup-gate counter resets at fold + // boundary so the new fold cold-starts with α_grad force-closed + // and ramps over WARMUP_STEPS_FALLBACK steps. Without this, fold N + // would inherit fold N-1's saturated warmup gate and the EGF wire + // would be hot from step 0 — defeating the cold-start safety + // contract documented in `sp14_isv_slots::WARMUP_STEPS_FALLBACK`. + self.fold_step_counter = 0; // Reset loss/Q-value history self.loss_history.clear(); diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index cef4421b1..c9ee3e54f 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -4079,6 +4079,79 @@ impl DQNTrainer { } } + // SP14 Layer B Task B.11 (2026-05-05): EGF producer chain. + // + // Three launches; the first two run per training step here, the + // third (gradient_hack_detect) fires once per epoch in + // `process_epoch_boundary`. Order is load-bearing: + // + // 1. q_disagreement_update_kernel — reads the captured-graph + // `aux_nb_softmax_buf [B, 2]` (aux head forward this step) + // + `q_out_buf [B, 13]` (populated by the captured + // `populate_q_out` in `submit_aux_ops`) and updates + // ISV[Q_DISAGREEMENT_SHORT/LONG_EMA = 383/384] + + // ISV[Q_DISAGREEMENT_VARIANCE_EMA = 389]. Same launch + // cadence as `launch_sp13_aux_dir_metrics` above — + // producer-only. + // + // 2. alpha_grad_compute_kernel — consumes #1's ISV[383] / + // ISV[389] outputs AND `launch_sp13_aux_dir_metrics`'s + // ISV[373] / ISV[374] outputs to compute the EGF gate + // ISV[ALPHA_GRAD_SMOOTHED = 393]. This is the kernel that + // activates the EGF wire — pre-B.11 the slot stayed at + // sentinel 0.0 so the backward wire-column scaling + // (B.10) effectively zeroed gradient flow. Post-B.11 the + // slot tracks the live gate value ∈ [0, 1] each step. + // + // The backward consumer (`launch_sp14_scale_wire_col` inside + // `launch_cublas_backward_to`) reads ISV[393] inside the + // CAPTURED graph. Because these two producer launches run + // OUTSIDE the captured graph (after `run_full_step` returns, + // identical position to `launch_sp13_aux_dir_metrics`), the + // backward consumer at training step N reads the value + // written by step N-1's producer chain (one-step lag). This + // is the same one-step-lag pattern as the forward wire + // (`launch_sp14_dir_concat_qaux` consumes the prior step's + // `aux_nb_softmax_buf`). Bootstrap effect: at step 0, ISV[393] + // is still 0.0 → wire column zeroed; from step 1 onward the + // gate value accrues. + // + // Fold-cold-start contract: `fold_step_counter` is reset to 0 + // in `reset_for_fold`, ensuring each fold's first step sees + // `warmup_gate = 0` (closed gate) and ramps over + // `WARMUP_STEPS_FALLBACK = 1000` steps to 1. + // + // Constants per plan §2581-2585: + // * α_short = 0.3 (fast EMA, ~3-step half-life) + // * α_long = 0.05 (slow EMA, ~14-step half-life) + // * α_var = 0.05 (Welford EMA blend; shared by var_q + // in this kernel and var_alpha + var_aux + // in alpha_grad_compute) + // + // var_aux producer gap (from B.4 status report) is now closed: + // alpha_grad_compute_kernel writes ISV[VAR_AUX = 388] inline + // (Welford on aux_short vs aux_long), so the adaptive k_aux + // branch is no longer degenerate at K_BASE_AUX = 20. + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::WARMUP_STEPS_FALLBACK; + let alpha_short_qd: f32 = 0.3; + let alpha_long_qd: f32 = 0.05; + let alpha_var: f32 = 0.05; + if let Err(e) = fused.trainer().launch_sp14_q_disagreement_update( + alpha_short_qd, alpha_long_qd, alpha_var, + ) { + tracing::warn!("SP14 B.11 q_disagreement_update launch failed: {e}"); + } + let warmup_gate = ((self.fold_step_counter as f32) + / (WARMUP_STEPS_FALLBACK as f32)).min(1.0); + if let Err(e) = fused.trainer().launch_sp14_alpha_grad_compute( + warmup_gate, alpha_var, + ) { + tracing::warn!("SP14 B.11 alpha_grad_compute launch failed: {e}"); + } + } + self.fold_step_counter += 1; + // Phase 3 T3.5: MoE expert-utilisation EMA + gate entropy EMA. // Reads `moe_gate_softmax_buf [B, K]` (valid after the captured // forward graph ran) and EMA-updates ISV[118..127). @@ -5058,6 +5131,32 @@ impl DQNTrainer { // Q-stats reset moved to reset_epoch_state() — runs at the START of the // next epoch, so log_epoch_metrics_and_financials can read the final values. + // SP14 Layer B Task B.11 (2026-05-05): EGF anti-mesa-optimization + // circuit breaker. Single-launch state machine that decrements the + // lockout counter (ISV[GRADIENT_HACK_LOCKOUT_REMAINING = 395]) by + // 1.0 per epoch, tracks the post-Schmitt-open minimum of + // aux_dir_acc (ISV[AUX_DIR_ACC_POST_OPEN_MIN = 394]), and forces + // ISV[GATE1_OPEN_STATE = 391] to 0 if the joint hacking signature + // fires (aux drops > LOCKOUT_TRIGGER_DROP=0.05 below open-thresh + // AND q_disagreement rises > LOCKOUT_TRIGGER_DIS_RISE=0.10 above + // baseline). See `gradient_hack_detect_kernel.cu` header for full + // detection rationale. + // + // MUST run AFTER all per-step alpha_grad_compute launches in the + // epoch (the circuit breaker reads the epoch-final ISV[391] gate + // state). Stream order is enforced by sequential same-stream + // submission — alpha_grad runs in `run_training_steps_slices` on + // the trainer stream; this launch goes to the same stream. + // + // Producer-only at the trainer level — the kernel mutates ISV + // state consumed by NEXT epoch's alpha_grad_compute via the + // persistent slot 391 + 395 + 394 reads. No host readback. + if let Some(ref fused) = self.fused_ctx { + if let Err(e) = fused.trainer().launch_sp14_gradient_hack_detect() { + tracing::warn!("SP14 B.11 gradient_hack_detect launch failed: {e}"); + } + } + // Return raw accumulated totals for epoch metric computation // The caller computes per-step averages from epoch_loss/n, etc. Ok(EpochBoundaryMetrics { diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index c6f3e8665..1f96a8934 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -82,6 +82,7 @@ mod gpu { use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; use ml::cuda_pipeline::sp13_isv_slots::{ + AUX_DIR_ACC_LONG_EMA_INDEX, AUX_DIR_ACC_SHORT_EMA_INDEX, TARGET_DIR_ACC_INDEX, }; @@ -194,6 +195,10 @@ mod gpu { isv_buf.write_from_slice(&isv); let bsi: i32 = B as i32; + // SP14 B.11 added a `q_stride` parameter (int) so the kernel can + // consume either tight [B, K_DIR=4] or wider [B, total_actions=13] + // q_logits buffers without a copy. Tight test buffer → stride = 4. + let q_stride: i32 = K_DIR as i32; let alpha_short: f32 = 0.3; let alpha_long: f32 = 0.05; let alpha_var: f32 = 0.05; @@ -205,6 +210,7 @@ mod gpu { .arg(&q_buf.dev_ptr) .arg(&isv_buf.dev_ptr) .arg(&bsi) + .arg(&q_stride) .arg(&alpha_short) .arg(&alpha_long) .arg(&alpha_var) @@ -282,6 +288,9 @@ mod gpu { isv_buf.write_from_slice(&isv); let bsi: i32 = B as i32; + // SP14 B.11: `q_stride` argument (see q_disagreement_k4_k2_mapping + // for rationale). Tight test buffer → stride = K_DIR = 4. + let q_stride: i32 = K_DIR as i32; let alpha_short: f32 = 0.3; let alpha_long: f32 = 0.05; let alpha_var: f32 = 0.05; @@ -293,6 +302,7 @@ mod gpu { .arg(&q_buf.dev_ptr) .arg(&isv_buf.dev_ptr) .arg(&bsi) + .arg(&q_stride) .arg(&alpha_short) .arg(&alpha_long) .arg(&alpha_var) @@ -362,12 +372,18 @@ mod gpu { // SP13 driver slots (no shift). isv[TARGET_DIR_ACC_INDEX] = 0.55; isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.55; // exactly at target → below threshold_open=0.58 + isv[AUX_DIR_ACC_LONG_EMA_INDEX] = 0.55; // SP14 B.11: var_aux mean + // reference; matches short + // → diff = 0 → var_aux stays + // at sentinel 0 on step 1. // SP14 driver slots (post-+2 shift). isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = 0.55; // > 0.5 baseline → gate2 ≈ 0.66 isv[K_AUX_ADAPTIVE_INDEX] = 20.0; isv[K_Q_ADAPTIVE_INDEX] = 15.0; isv[BETA_RATE_LIMITER_ADAPTIVE_INDEX] = 0.5; isv[AUX_DIR_ACC_VARIANCE_EMA_INDEX] = 0.0; // → k_aux = K_BASE_AUX = 20.0 + // (var_aux now produced by + // this kernel, see B.11) isv[Q_DISAGREEMENT_VARIANCE_EMA_INDEX] = 0.0; // → k_q = K_BASE_Q = 15.0 isv[ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX] = 0.0; // → β = β_base = 0.5 isv[GATE1_OPEN_STATE_INDEX] = 0.0; // closed (sentinel) @@ -515,6 +531,18 @@ mod gpu { let mut isv = vec![0.0_f32; ISV_DIM]; isv[TARGET_DIR_ACC_INDEX] = 0.55; isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.55; + isv[AUX_DIR_ACC_LONG_EMA_INDEX] = 0.55; // SP14 B.11: var_aux mean + // reference. Test holds + // long_ema fixed at 0.55 + // while oscillating short + // ±0.05 — diff = ±0.05 + // produces var_aux ≈ 0.0025 + // by end of 20 steps, + // softening k_aux to ≈11.4 + // (still well above K_MIN=1). + // Adaptive β still grows + // because α_raw chatter + // persists. isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = 0.6; isv[K_AUX_ADAPTIVE_INDEX] = 20.0; isv[K_Q_ADAPTIVE_INDEX] = 15.0; diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c60eae9b3..9d2bc4e32 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -6642,3 +6642,31 @@ B.8/B.9 grew `w_b0fc` to `[adv_h, SH2 + 1]` end-to-end across forward dispatch + - **Reverse dependencies**: `apply_iqn_trunk_gradient` and aux-paths use the `bw_d_h_s2` that this path writes — the new accumulator pipeline (memset → backward_full → wire-col scale → strided accumulate from d_dir_qaux_concat → mag/ord/urg accumulators → value-FC inside backward_full) leaves `bw_d_h_s2` with the same algebraic value as pre-B.10 *except* for the gated wire-col contribution from the direction-Q's first FC. Pre-B.11 (α=0) the gated contribution is zero → bit-identical to pre-B.10. - **CudaSlice wrapper path**: passes `0u64` for both `dir_qaux_concat_ptr` and `d_dir_qaux_concat_ptr`, falling back to the legacy K=SH2 path. This is consistent with the forward CudaSlice wrapper (`dir_qaux_concat_ptr: 0`); the wrapper-based callers (causal intervention, DDQN argmax) are diagnostic-only paths whose direction-Q outputs are downstream-bounded per the B.9 residual-path analysis. +## SP14 Layer B Task B.11 — Producer chain orchestrator (2026-05-05) + +**Goal:** Wire the 3 EGF producer kernels (B.3 q_disagreement, B.4 alpha_grad, B.5 gradient_hack_detect) into per-step / per-epoch hooks. After this commit, `α_grad_smoothed` is computed every step from real driver signals (vs. holding sentinel 0.0 force-closed pre-B.11). The EGF wire becomes ACTIVE. + +### Wire status + +- **q_disagreement_update**: per-step launch after action select. Reads `aux_nb_softmax_buf [B, K=2]` + `q_dir_logits [B, K=4]`; writes ISV[383] (short EMA), ISV[384] (long EMA), ISV[389] (variance EMA). α_short=0.3, α_long=0.05, α_var=0.05 (plan-spec'd). +- **alpha_grad_compute**: per-step launch after q_disagreement, before backward. Reads driver signals (target_dir_acc, aux_dir_acc_short, q_disagreement_short, var_aux, var_q, var_alpha_prev, gate1_state); writes ISV[385..395] (k_aux, k_q, β, var_alpha, gate1_state, α_raw, α_smoothed). Warmup gate computed from steps_in_fold / WARMUP_STEPS_FALLBACK. +- **gradient_hack_detect**: per-epoch launch at end of epoch (process_epoch_boundary). Decrements lockout, evaluates trigger conditions, force-closes gate1 if triggered. +- **var_aux producer gap closed (option C from B.4)**: alpha_grad_compute_kernel extended to also write ISV[VAR_AUX_INDEX=388] using Welford EMA against `aux_dir_acc_short - aux_dir_acc_long`. Adaptive `k_aux` is now functional (was degenerate at K_BASE_AUX=20.0 pre-B.11). + +### Files changed +- `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu` — var_aux Welford write added +- `crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu` — minor adjustments +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 3 launcher methods (`launch_sp14_q_disagreement_update`, `launch_sp14_alpha_grad_compute`, `launch_sp14_gradient_hack_detect`) +- `crates/ml/src/trainers/dqn/trainer/{constructor,mod,training_loop}.rs` — per-step + per-epoch hooks +- `crates/ml/tests/sp14_oracle_tests.rs` — updated test expectations for var_aux + +### Verification +- `cargo check -p ml` — clean, 18 warnings (pre-existing baseline) +- `cargo test -p ml --lib aux_w` — 4/4 P0b tests pass (no regression) + +### Wire status (cumulative post-B.11) +- **Forward**: B.9 active (concat → SGEMM with K=SH2+1) +- **Backward**: B.10 active (wire-col scale by ISV[393]; dW unchanged) +- **Producers**: B.11 active (3 launches every step + 1 per epoch) +- **EGF gate**: now responsive to live signals (was force-closed at sentinel 0.0) +