feat(sp14): B.11 — orchestrator wire-up for 3 EGF producer kernels + var_aux gap closure
Per-step launches (in graph capture order): 1. Forward (existing) 2. Action select (existing) → q_dir_logits available 3. launch_sp14_q_disagreement_update → ISV[383, 384, 389] 4. launch_sp14_alpha_grad_compute → ISV[385..395] (consumes q_disagreement) 5. Backward (existing) — wire-col scale at B.10 reads ISV[393] Per-epoch launch (end of epoch): 6. launch_sp14_gradient_hack_detect → circuit breaker α_short=0.3, α_long=0.05, α_var=0.05 per spec; warmup_gate derived from steps_in_fold / WARMUP_STEPS_FALLBACK. Var_aux producer gap closed (option C from B.4): alpha_grad_compute_kernel now also writes ISV[VAR_AUX_INDEX=388] via 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 constant pre-B.11). Closes the "adaptive_k_aux currently degenerate" concern flagged in B.4 commit. After this commit, the EGF pearl is FULLY ACTIVE end-to-end: - Forward: aux signal feeds direction Q-head input (B.8/B.9) - Backward: wire-col gradient gated by α_grad_smoothed (B.10) - Producers: α_grad computed every step from real driver signals (B.11) - Pre-B.11 force-closed gate (sentinel 0.0) → post-B.11 responsive gate Build clean: 18 warnings pre-existing baseline, 0 new. Tests: 4/4 P0b aux_w tests pass (no regression). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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;
|
||||
|
||||
@@ -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::<f32>() 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
|
||||
|
||||
@@ -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; }
|
||||
}
|
||||
|
||||
|
||||
@@ -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,
|
||||
|
||||
|
||||
@@ -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();
|
||||
|
||||
@@ -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 {
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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)
|
||||
|
||||
|
||||
Reference in New Issue
Block a user