diff --git a/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu b/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu index 6dc22d95e..7569aabbd 100644 --- a/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu @@ -203,8 +203,6 @@ extern "C" __global__ void aux_next_bar_loss_reduce( const float* __restrict__ pred, /* [B] (== [B, 1] flat) */ const float* __restrict__ label, /* [B] */ int B, - const float* __restrict__ isv, /* ISV bus, read isv[isv_label_scale_index] */ - int isv_label_scale_index, /* AUX_LABEL_SCALE_EMA_INDEX = 117 */ float* __restrict__ loss_out /* [1] */ ) { if (blockIdx.x != 0) return; @@ -212,23 +210,15 @@ extern "C" __global__ void aux_next_bar_loss_reduce( const int tid = threadIdx.x; const int block = blockDim.x; - /* Plan 5 Task 5 follow-up: normalize label by ISV-driven label-scale - * EMA so the residual stays unit-scale regardless of underlying data - * magnitude (1e-3 log-returns vs 5000 raw prices). Read the scale - * directly from ISV at kernel launch — graph-capture-stable since the - * ISV device pointer + slot index are stable across re-launches; the - * actual scalar value updates each step via `aux_label_scale_ema_update` - * launched immediately before this kernel. The scale floor `1e-6` is - * a numerical-stability epsilon, NOT a tuned constant — it only fires - * when the EMA happens to be ≤ 1e-6 (e.g. literally zero labels), in - * which case dividing by 1e-6 yields a 1e6-magnitude residual that - * the gradient clip ceiling catches downstream. */ - const float label_scale = isv[isv_label_scale_index]; - const float inv_scale = 1.0f / fmaxf(label_scale, 1e-6f); - + /* SP13 B1.0 (2026-05-05): scale-free MSE. Labels are z-normalised at + * the data layer, so `label_scale_ema ≈ 1.0` empirically and the + * pre-B1.0 `(pred - label * inv_scale)` reduces to `(pred - label)` + * within rounding. The ISV-driven label-scale divisor (former slot + * 117 / `AUX_LABEL_SCALE_EMA_INDEX`) is retired — B1.1 will replace + * MSE with CE, eliminating the loss formulation entirely. */ float local = 0.0f; for (int i = tid; i < B; i += block) { - const float d = pred[i] - label[i] * inv_scale; + const float d = pred[i] - label[i]; local += d * d; } smem[tid] = local; @@ -399,8 +389,6 @@ extern "C" __global__ void aux_next_bar_backward( const float* __restrict__ label, /* [B] */ int B, int SH2, - const float* __restrict__ isv, /* ISV bus, read isv[isv_label_scale_index] */ - int isv_label_scale_index, /* AUX_LABEL_SCALE_EMA_INDEX = 117 */ /* Partial outputs (per-sample). Caller reduces along batch dim. */ float* __restrict__ dW1_partial, /* [B, H, SH2] flat */ float* __restrict__ db1_partial, /* [B, H] */ @@ -419,15 +407,13 @@ extern "C" __global__ void aux_next_bar_backward( float* sh_dh_pre = smem; /* [H] */ float* sh_h_post = smem + H; /* [H] */ - /* Plan 5 Task 5 follow-up: normalize label by ISV-driven label-scale - * EMA so the gradient stays unit-scale regardless of underlying data - * magnitude. MUST match the same `inv_scale` arithmetic used in - * `aux_next_bar_loss_reduce` so loss + gradient are derivatives of - * the same scalar function. */ - const float label_scale = isv[isv_label_scale_index]; - const float inv_scale = 1.0f / fmaxf(label_scale, 1e-6f); - /* Mean-over-batch derivative scale: dL/dpred[b] = (2/B) * (pred[b] - label[b]/scale). */ - const float d_pred_b = (2.0f / (float)B) * (pred[b] - label[b] * inv_scale); + /* SP13 B1.0 (2026-05-05): scale-free MSE backward. Mirrors the + * `aux_next_bar_loss_reduce` arithmetic so loss + gradient are + * derivatives of the same scalar function. The ISV-driven label-scale + * divisor (former slot 117) is retired — see the loss kernel's B1.0 + * note for context. + * Mean-over-batch derivative: dL/dpred[b] = (2/B) * (pred[b] - label[b]). */ + const float d_pred_b = (2.0f / (float)B) * (pred[b] - label[b]); /* Step 1: cache h_post and compute d_h_pre per hidden unit. */ for (int k = tid; k < H; k += blockDim.x) { diff --git a/crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu b/crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu index ee532a333..3c54d7cf1 100644 --- a/crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu @@ -2,10 +2,10 @@ * loss scalars, written into producer_step_scratch_buf[scratch_base..+2) * for Pearls A+D consumption. SP4 Layer A Task A13.1 retrofit (2026-05-01). * - * Plan 4 Task 6 Commit A. Producer-only kernel — Commit B wires consumers - * (HEALTH_DIAG emission + aux-loss attribution monitor) which read - * ISV[AUX_NEXT_BAR_MSE_EMA_INDEX] / ISV[AUX_REGIME_CE_EMA_INDEX]; only the - * EMA mechanism behind the slots changes here. + * Plan 4 Task 6 Commit A. Producer-only kernel — HEALTH_DIAG emission + + * aux-loss attribution monitor read ISV[AUX_NEXT_BAR_MSE_EMA_INDEX] / + * ISV[AUX_REGIME_CE_EMA_INDEX]; only the EMA mechanism behind the slots + * changes here. * * Mirrors `h_s2_rms_ema_kernel.cu`'s retrofit shape — single thread, * single block, cold-path-cadence (one launch per training step alongside @@ -22,6 +22,13 @@ * GPU-drives-CPU-reads (spec §4.C.6) and pearl_cold_path_no_exception_to * _gpu_drives.md / feedback_no_cpu_compute_strict.md compliant. * No atomicAdd (per feedback_no_atomicadd.md). + * + * SP13 B1.0 (2026-05-05): the sibling `aux_label_scale_ema_update` kernel + * (formerly in this same file, slot 117 producer) was retired together + * with the ISV-driven label-scale division in `aux_next_bar_loss_reduce` + * / `aux_next_bar_backward`. Labels are already z-normalised at the data + * layer (`label_scale_ema ≈ 1.0` empirically), so the divisor became a + * no-op. B1.0 is a bridge — B1.1 replaces the entire MSE chain with CE. */ extern "C" __global__ void aux_heads_loss_ema_update( @@ -39,75 +46,3 @@ extern "C" __global__ void aux_heads_loss_ema_update( scratch_buf[scratch_idx_regime] = regime_loss_scalar[0]; __threadfence_system(); /* PCIe-visible writes for mapped pinned host_ptr */ } - -/* ==================================================================== - * aux_label_scale_ema_update — Plan 4 Task 6 / Plan 5 Task 5 follow-up; - * SP4 Layer A Task A13.1 retrofit (2026-05-01). - * - * Reduces `mean(|label[i]|)` over the B samples of the already-gathered - * `aux_nb_label_buf [B]` (= column 0 of next_states for that step) and - * writes the step observation to `producer_step_scratch_buf[scratch_idx]` - * for Pearls A+D consumption. **α dropped per SP4 — Pearls A+D adapt α - * from per-slot signal-vs-noise variance.** - * - * Consumer (unchanged): `aux_next_bar_loss_reduce` and - * `aux_next_bar_backward` divide the label by `max(scale, 1e-6)` before - * the residual `(pred - label)`, so training dynamics stay unit-scale - * regardless of whether the underlying data column is a log return - * (~1e-3) or a raw price (~5000). The consumers still read - * ISV[AUX_LABEL_SCALE_EMA_INDEX]; only the EMA mechanism populating that - * slot changes here. - * - * Ordering: this kernel + the GPU `apply_pearls_ad_kernel` (2026-05-01 - * GPU-Pearls refactor — same-stream chained launch, no host sync) MUST - * complete BEFORE `aux_next_bar_loss_reduce` and `aux_next_bar_backward` - * so they see an up-to-date scale. Stream ordering provides the fence; - * graph-capture-compatible by construction. - * - * Single-block, 256-thread shmem-tree reduction — same shape as the - * existing `aux_next_bar_loss_reduce` / `aux_regime_loss_reduce`. No - * atomicAdd. Writeback path is fully GPU-side via the chained - * `apply_pearls_ad_kernel`. - * - * Cost: 1 batch read + 1 scratch write per launch — under any kernel- - * launch overhead. - * ==================================================================== */ -#ifndef AUX_LABEL_SCALE_BLOCK -#define AUX_LABEL_SCALE_BLOCK 256 -#endif - -extern "C" __global__ void aux_label_scale_ema_update( - const float* __restrict__ label, /* [B] aux_nb_label_buf */ - int B, - float* __restrict__ scratch_buf, /* producer_step_scratch_buf */ - int scratch_idx /* slot 43 — label-scale step obs */ -) { - /* Single-block contract. */ - if (blockIdx.x != 0) return; - - extern __shared__ float smem[]; - const int tid = threadIdx.x; - const int block = blockDim.x; - - /* Stride loop, fabsf-and-sum into per-thread accumulator. */ - float local = 0.0f; - for (int i = tid; i < B; i += block) { - local += fabsf(label[i]); - } - smem[tid] = local; - __syncthreads(); - - /* Standard shmem-tree reduction. */ - for (int s = block / 2; s > 0; s >>= 1) { - if (tid < s) smem[tid] += smem[tid + s]; - __syncthreads(); - } - - /* Thread 0 writes the step observation. B > 0 invariant — caller - * guarantees a positive batch. */ - if (tid == 0) { - const float mean_abs = smem[0] / (float)B; - scratch_buf[scratch_idx] = mean_abs; - __threadfence_system(); - } -} diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs b/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs index 49ba53a59..6e11ba68a 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs @@ -74,12 +74,13 @@ pub(crate) struct AuxHeadsForwardOps { regime_label_kernel: CudaFunction, /// Aux loss EMA producer kernel (single thread, single block; mirrors /// `h_s2_rms_ema_update`'s shape). Loaded from `AUX_HEADS_LOSS_EMA_CUBIN`. + /// + /// SP13 B1.0 (2026-05-05): the sibling label-scale EMA kernel was + /// retired together with the ISV[117] label-scale divisor in + /// `aux_next_bar_loss_reduce` / `aux_next_bar_backward`. Labels are + /// already z-normalised at the data layer; B1.1 will replace MSE + /// with CE entirely. loss_ema_kernel: CudaFunction, - /// Aux next-bar label-scale EMA producer kernel (single-block 256-thread - /// shmem reduction over `aux_nb_label_buf [B]`). Plan 5 Task 5 follow-up: - /// ISV-driven label-scale normalisation. Loaded from the same cubin as - /// `loss_ema_kernel`. - label_scale_ema_kernel: CudaFunction, } impl AuxHeadsForwardOps { @@ -113,9 +114,6 @@ impl AuxHeadsForwardOps { let loss_ema_kernel = ema_module .load_function("aux_heads_loss_ema_update") .map_err(|e| MLError::ModelError(format!("aux_heads_loss_ema_update load: {e}")))?; - let label_scale_ema_kernel = ema_module - .load_function("aux_label_scale_ema_update") - .map_err(|e| MLError::ModelError(format!("aux_label_scale_ema_update load: {e}")))?; Ok(Self { next_bar_forward_kernel, @@ -124,7 +122,6 @@ impl AuxHeadsForwardOps { regime_loss_reduce_kernel, regime_label_kernel, loss_ema_kernel, - label_scale_ema_kernel, }) } @@ -226,22 +223,16 @@ impl AuxHeadsForwardOps { /// Launch `aux_next_bar_loss_reduce`: scalar mean MSE into `loss_out[1]`. /// - /// Plan 5 Task 5 follow-up: the kernel reads `isv[isv_label_scale_index]` - /// and divides each label by `max(scale, 1e-6)` before the residual - /// `(pred - label/scale)` so the loss stays unit-scale regardless of - /// underlying data magnitude (1e-3 log-returns vs 5000 raw prices). The - /// device pointer + slot index pair is graph-capture-stable; the actual - /// scalar updates each step via `aux_label_scale_ema_update` launched - /// immediately before this kernel. - #[allow(clippy::too_many_arguments)] + /// SP13 B1.0 (2026-05-05): scale-free MSE — labels are z-normalised at + /// the data layer so the pre-B1.0 ISV-driven `inv_scale` divisor reduces + /// to a no-op. The kernel now computes the residual `(pred - label)` + /// directly. B1.1 will replace MSE with CE entirely. pub(crate) fn next_bar_loss_reduce( &self, stream: &Arc, pred_ptr: u64, label_ptr: u64, b: usize, - isv_dev_ptr: u64, - isv_label_scale_index: i32, loss_out_ptr: u64, ) -> Result<(), MLError> { let b_i32 = b as i32; @@ -252,8 +243,6 @@ impl AuxHeadsForwardOps { .arg(&pred_ptr) .arg(&label_ptr) .arg(&b_i32) - .arg(&isv_dev_ptr) - .arg(&isv_label_scale_index) .arg(&loss_out_ptr) .launch(LaunchConfig { grid_dim: (1, 1, 1), @@ -379,46 +368,6 @@ impl AuxHeadsForwardOps { Ok(()) } - /// Launch `aux_label_scale_ema_update` (Plan 5 Task 5 follow-up; - /// SP4 Layer A Task A13.1 retrofit 2026-05-01; GPU-Pearls refactor - /// 2026-05-01): single-block 256-thread shmem reduction over - /// `aux_nb_label_buf [B]`, writes `mean(|label|)` to - /// `producer_step_scratch_buf[scratch_idx]`. - /// - /// Consumes the dedicated label buffer that `aux_heads_forward` populated - /// via `strided_gather` of `next_states[:, 0]`. The kernel writes ONLY - /// the step observation; the inline GPU `apply_pearls_ad_kernel` launch - /// lives in `GpuDqnTrainer::aux_heads_forward` Step 2b (same stream, - /// no host sync) so consumers see the up-to-date - /// `ISV[AUX_LABEL_SCALE_EMA_INDEX]` mid-step. This site was the one - /// that broke L40S smoke `smoke-test-v9kjv` graph capture before the - /// 2026-05-01 GPU-Pearls refactor. - pub(crate) fn launch_label_scale_ema( - &self, - stream: &Arc, - label_ptr: u64, - b: usize, - scratch_dev_ptr: u64, - scratch_idx: i32, - ) -> Result<(), MLError> { - let b_i32 = b as i32; - let smem_bytes = AUX_BLOCK * std::mem::size_of::() as u32; - unsafe { - stream - .launch_builder(&self.label_scale_ema_kernel) - .arg(&label_ptr) - .arg(&b_i32) - .arg(&scratch_dev_ptr) - .arg(&scratch_idx) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (AUX_BLOCK, 1, 1), - shared_mem_bytes: smem_bytes, - }) - .map_err(|e| MLError::ModelError(format!("aux_label_scale_ema_update: {e}")))?; - } - Ok(()) - } } /// Backward orchestrator — owns the 3 backward / final-reduce kernel @@ -468,11 +417,10 @@ impl AuxHeadsBackwardOps { /// `param_grad_reduce` for each of {`dW1`, `db1`, `dW2`, `db2`} to /// collapse along the batch dim. /// - /// Plan 5 Task 5 follow-up: the kernel reads `isv[isv_label_scale_index]` - /// and divides each label by `max(scale, 1e-6)` before the residual - /// — MUST mirror the `aux_next_bar_loss_reduce` arithmetic so loss + - /// gradient are derivatives of the same scalar function. The device - /// pointer + slot index pair is graph-capture-stable. + /// SP13 B1.0 (2026-05-05): scale-free MSE backward — mirrors the + /// `aux_next_bar_loss_reduce` arithmetic so loss + gradient are + /// derivatives of the same scalar function. The pre-B1.0 ISV-driven + /// `inv_scale` divisor (former slot 117) is retired. #[allow(clippy::too_many_arguments)] pub(crate) fn backward_next_bar( &self, @@ -485,8 +433,6 @@ impl AuxHeadsBackwardOps { label_ptr: u64, b: usize, sh2: usize, - isv_dev_ptr: u64, - isv_label_scale_index: i32, dw1_partial_ptr: u64, db1_partial_ptr: u64, dw2_partial_ptr: u64, @@ -507,8 +453,6 @@ impl AuxHeadsBackwardOps { .arg(&label_ptr) .arg(&b_i32) .arg(&sh2_i32) - .arg(&isv_dev_ptr) - .arg(&isv_label_scale_index) .arg(&dw1_partial_ptr) .arg(&db1_partial_ptr) .arg(&dw2_partial_ptr) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index aca348d4d..dd90faa4e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1027,20 +1027,15 @@ const ISV_NETWORK_DIM: usize = 23; /// [114] AUX_REGIME_CE_EMA_INDEX — Plan 4 Task 6 Commit A — EMA of the /// regime classification head's mean-batch cross-entropy /// (α=0.05). Same producer/contract as [113]; same FoldReset → 0.0. -/// [117] AUX_LABEL_SCALE_EMA_INDEX — Plan 4 Task 6 / Plan 5 Task 5 follow- -/// up — EMA of `mean(|next_states[:, 0]|)` over the batch (α=0.05). -/// Producer: `aux_label_scale_ema_update` GPU kernel (single-block -/// shmem reduction over the already-gathered `aux_nb_label_buf [B]`, -/// then EMA blend with the prior ISV value). Consumer: read on host -/// every step BEFORE `aux_next_bar_loss_reduce` and `aux_next_bar_ -/// backward` launch, passed as a `float` kernel argument; both -/// kernels divide the label by `max(scale, 1e-6)` so the residual -/// `(pred - label/scale)` stays unit-scale regardless of underlying -/// data magnitude (1e-3 log-returns vs 5000 raw prices). FoldReset -/// → 1.0 (multiplicative identity — divide-by-1 is a no-op so the -/// first batch of a new fold trains against the un-normalised label -/// until the producer fires once and the EMA updates to the actual -/// scale). +/// [117] RETIRED — formerly `AUX_LABEL_SCALE_EMA_INDEX` (Plan 4 Task 6 / +/// Plan 5 Task 5 follow-up). Retired in SP13 B1.0 (2026-05-05) when +/// `aux_next_bar_loss_reduce` / `aux_next_bar_backward` flipped to +/// scale-free MSE: labels are z-normalised at the data layer so the +/// pre-B1.0 ISV-driven `(pred - label / scale)` residual reduces to +/// `(pred - label)` within rounding. The slot remains unused in the +/// ISV bus (layout fingerprint reserves it via the seed; do not +/// reuse without a fingerprint bump). B1.1 will replace MSE with CE +/// entirely, eliminating the formulation that made this slot useful. /// [115] ISV_LAYOUT_FINGERPRINT_LO_INDEX — low 32 bits of layout fingerprint /// (shifted 111→115 in Plan 4 Task 6 Commit A). /// [116] ISV_LAYOUT_FINGERPRINT_HI_INDEX — high 32 bits of layout fingerprint @@ -1833,29 +1828,21 @@ pub const VSN_MASK_GROUP_5_EMA_INDEX: usize = 110; pub const AUX_NEXT_BAR_MSE_EMA_INDEX: usize = 113; pub const AUX_REGIME_CE_EMA_INDEX: usize = 114; -/// ISV slot [117] — Plan 4 Task 6 / Plan 5 Task 5 follow-up: aux next-bar -/// label-scale EMA. +/// ISV slot [117] — RETIRED in SP13 B1.0 (2026-05-05). /// -/// EMA of `mean(|next_states[:, 0]|)` over the batch (α=0.05). Producer: -/// `aux_label_scale_ema_update` GPU kernel (single-block 256-thread shmem -/// reduction over the already-gathered `aux_nb_label_buf [B]`, then EMA -/// blend with the prior ISV value into slot 117). Consumer: read on host -/// every step BEFORE `aux_next_bar_loss_reduce` and `aux_next_bar_backward` -/// launch, passed as a `float` kernel argument; both kernels divide the -/// label by `max(scale, 1e-6)` before the residual so the loss + gradient -/// stay unit-scale regardless of whether `next_states[:, 0]` carries log -/// returns (~1e-3) or raw prices (~5000). +/// Formerly `AUX_LABEL_SCALE_EMA_INDEX` — held the EMA of +/// `mean(|next_states[:, 0]|)` consumed by `aux_next_bar_loss_reduce` / +/// `aux_next_bar_backward` to divide the label by `max(scale, 1e-6)` +/// before the residual. Retired together with that division when the MSE +/// loss flipped to scale-free: labels are z-normalised at the data layer +/// so `label_scale_ema ≈ 1.0` empirically and `(pred - label * inv_scale)` +/// reduced to `(pred - label)` within rounding. /// -/// Cold-start: 1.0 (multiplicative identity — divide-by-1 is a no-op so the -/// first batch of a new fold trains against the un-normalised label until -/// the producer fires once and the EMA updates to the actual scale). -/// FoldReset → 1.0 (NOT 0.0 — divide-by-zero would explode the loss before -/// the producer could pull the EMA back to the real scale). -/// -/// Tail-appended after the fingerprint slots [115..117) per -/// `feedback_no_partial_refactor.md` (every consumer of `ISV_TOTAL_DIM` or -/// the fingerprint shifts together in this commit). -pub const AUX_LABEL_SCALE_EMA_INDEX: usize = 117; +/// The slot remains unused in the ISV bus and is reserved by the layout +/// fingerprint seed — do NOT reuse it without a fingerprint bump (which +/// would invalidate every checkpoint stored under the post-B1.0 +/// fingerprint). B1.1 will replace the MSE chain with CE entirely, +/// eliminating the formulation that made this slot useful. /// Mixture-of-Experts per-expert utilization EMA, α=0.05. /// 8 contiguous slots [118..126). @@ -2057,7 +2044,6 @@ const fn layout_fingerprint_seed() -> &'static [u8] { VSN_MASK_GROUP_3_EMA=108;VSN_MASK_GROUP_4_EMA=109;VSN_MASK_GROUP_5_EMA=110;\ AUX_NEXT_BAR_MSE_EMA=113;AUX_REGIME_CE_EMA=114;\ ISV_LAYOUT_FINGERPRINT_LO=115;ISV_LAYOUT_FINGERPRINT_HI=116;\ - AUX_LABEL_SCALE_EMA=117;\ MOE_EXPERT_UTIL_EMA_BASE=118;MOE_EXPERT_UTIL_EMA_COUNT=8;\ MOE_GATE_ENTROPY_EMA=126;\ MOE_LAMBDA_EFF=128;\ @@ -4018,7 +4004,10 @@ pub struct GpuDqnTrainer { /// [0..40) — SP4 producers (Tasks A5-A11) /// [40] — h_s2_rms_ema → ISV[96] (A13.0) /// [41..43) — aux_heads_loss_ema → ISV[113],ISV[114] (A13.1) - /// [43] — aux_label_scale_ema → ISV[117] (A13.1) + /// [43] — RETIRED — formerly aux_label_scale_ema → ISV[117] + /// (SP13 B1.0, 2026-05-05); slot reserved by + /// SP4_PRODUCER_COUNT for layout stability — + /// do not reuse without a fingerprint bump. /// [44..52) — moe_expert_util_ema → ISV[118..126) (A13.2) /// [52] — moe_gate_entropy_ema → ISV[126] (A13.2) /// [53..59) — vsn_mask_ema → ISV[105..111) (A13.3) @@ -10916,7 +10905,10 @@ impl GpuDqnTrainer { // ── Task A13 retrofit producers (Pearls A+D for existing EMA kernels) ── // 40 = H_S2_RMS_EMA → ISV[96] (A13.0) // 41..43 = AUX_HEADS_LOSS_EMA → ISV[113],ISV[114] (A13.1) - // 43 = AUX_LABEL_SCALE_EMA → ISV[117] (A13.1) + // 43 = RETIRED (SP13 B1.0) — formerly AUX_LABEL_SCALE_EMA + // → ISV[117]; slot reserved + // by SP4_PRODUCER_COUNT for + // layout stability. // 44..52 = MOE_EXPERT_UTIL_EMA → ISV[118..126) (A13.2) // 52 = MOE_GATE_ENTROPY_EMA → ISV[126] (A13.2) // 53..59 = VSN_MASK_EMA → ISV[105..111) (A13.3) @@ -14787,7 +14779,6 @@ impl GpuDqnTrainer { TARGET_DRIFT_DIR_EMA_INDEX as i32, AUX_NEXT_BAR_MSE_EMA_INDEX as i32, AUX_REGIME_CE_EMA_INDEX as i32, - AUX_LABEL_SCALE_EMA_INDEX as i32, MOE_EXPERT_UTIL_EMA_BASE as i32, MOE_GATE_ENTROPY_EMA_INDEX as i32, MOE_LAMBDA_EFF_INDEX as i32, @@ -15279,64 +15270,19 @@ impl GpuDqnTrainer { } } - // Step 2b: ISV-driven label-scale step-observation producer - // (Plan 5 Task 5 follow-up; SP4 Layer A Task A13.1 retrofit; - // 2026-05-01 GPU-Pearls refactor). - // - // Reads the JUST-gathered `aux_nb_label_buf [B]` and writes - // `mean(|label|)` to `producer_step_scratch_buf[43]`. The GPU - // `apply_pearls_ad_kernel` then maps that step observation into - // ISV[AUX_LABEL_SCALE_EMA_INDEX] on the same stream, so subsequent - // consumers (`next_bar_loss_reduce` Step 5 + backward) read an - // up-to-date adaptive-α scale before computing the residual - // `(pred - label/scale)`. Wiener state at - // `wiener_state_buf[129..132)` (= scratch slot 43 × 3). - // - // **This site is the ONE that broke L40S smoke `smoke-test-v9kjv`** — - // the prior host-side `apply_pearls_to_slot` issued a mid-step - // `stream.synchronize()` inside the per-step captured graph, - // raising `CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED`. The GPU - // applicator's same-stream ordering eliminates the host sync; the - // captured graph now records the producer + applicator launches - // in one capture pass. - // - // **α dropped per SP4 — Pearls A+D adapt α from per-slot signal- - // vs-noise variance.** - const SCRATCH_IDX_LABEL_SCALE: usize = 43; - let scratch_dev = self.producer_step_scratch_buf.dev_ptr; - self.aux_heads_fwd.launch_label_scale_ema( - &self.stream, - self.aux_nb_label_buf.raw_ptr(), - b, - scratch_dev, - SCRATCH_IDX_LABEL_SCALE as i32, - )?; - - // DIAG_AUX_LABEL one-shot diagnostic removed (Fix 29) — leak pinned to - // experience_kernels.cu:698-704 vol_normalizer double-normalization. - - { - use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls; - // GPU Pearls A+D: stream-ordered with the producer kernel, - // single-slot, retrofit wiener convention: offset = scratch_idx*3 - // = 129. Degenerate-zero (mean(|label|) reduction yielded 0 → - // before first aux head forward populates the label buffer) - // handled by the applicator's `step_obs == 0` skip. - unsafe { - launch_apply_pearls( - &self.stream, - &self.apply_pearls_ad_kernel, - scratch_dev, - SCRATCH_IDX_LABEL_SCALE as i32, - self.isv_signals_dev_ptr, - AUX_LABEL_SCALE_EMA_INDEX as i32, - self.wiener_state_buf.dev_ptr, - (SCRATCH_IDX_LABEL_SCALE * 3) as i32, - 1, - crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META, - )?; - } - } + // SP13 B1.0 (2026-05-05): Step 2b retired together with the + // ISV[117] (`AUX_LABEL_SCALE_EMA_INDEX`) division in the + // `aux_next_bar_loss_reduce` / `aux_next_bar_backward` kernels. + // Labels are z-normalised at the data layer so `label_scale ≈ 1.0` + // empirically; the pre-B1.0 `(pred - label / scale)` residual + // reduces to `(pred - label)` within rounding. The per-step + // label-scale producer (`launch_label_scale_ema`) and its chained + // `apply_pearls_ad_kernel` applicator (which mapped scratch[43] → + // ISV[117], wiener_state[129..132)) are no longer launched. The + // scratch slot + wiener state remain reserved by + // SP4_PRODUCER_COUNT for layout stability — do not reuse without + // a fingerprint bump. B1.1 will replace MSE with CE and remove + // the formulation entirely. // Step 3: next-bar regression forward. self.aux_heads_fwd.forward_next_bar( @@ -15359,15 +15305,14 @@ impl GpuDqnTrainer { )?; // Step 5: next-bar MSE reduce against the dedicated label buffer. - // Kernel reads ISV[AUX_LABEL_SCALE_EMA_INDEX] (populated by the - // Step 2b producer) to normalise the label before the residual. + // SP13 B1.0 (2026-05-05): scale-free MSE — the kernel computes + // `(pred - label)` directly. Labels are z-normalised at the data + // layer so the pre-B1.0 ISV[117] divisor was a no-op in practice. self.aux_heads_fwd.next_bar_loss_reduce( &self.stream, self.aux_nb_pred_buf.raw_ptr(), self.aux_nb_label_buf.raw_ptr(), b, - self.isv_signals_dev_ptr, - AUX_LABEL_SCALE_EMA_INDEX as i32, self.aux_nb_loss_scalar_buf.raw_ptr(), )?; @@ -15429,11 +15374,11 @@ impl GpuDqnTrainer { self.aux_nb_pred_buf.raw_ptr(), self.aux_nb_label_buf.raw_ptr(), b, sh2, - // Plan 5 Task 5 follow-up: read ISV-driven label-scale EMA from - // the same slot the forward `next_bar_loss_reduce` read so loss - // and gradient are derivatives of the same scalar function. - self.isv_signals_dev_ptr, - AUX_LABEL_SCALE_EMA_INDEX as i32, + // SP13 B1.0 (2026-05-05): scale-free MSE backward — mirrors + // the `next_bar_loss_reduce` arithmetic so loss + gradient + // remain derivatives of the same scalar function. The + // pre-B1.0 ISV[117] (`AUX_LABEL_SCALE_EMA_INDEX`) divisor is + // retired together with the loss-side division. self.aux_partial_nb_w1.raw_ptr(), self.aux_partial_nb_b1.raw_ptr(), self.aux_partial_nb_w2.raw_ptr(), diff --git a/crates/ml/src/cuda_pipeline/gpu_health_diag.rs b/crates/ml/src/cuda_pipeline/gpu_health_diag.rs index 347ac4775..decc7f320 100644 --- a/crates/ml/src/cuda_pipeline/gpu_health_diag.rs +++ b/crates/ml/src/cuda_pipeline/gpu_health_diag.rs @@ -152,7 +152,8 @@ impl GpuHealthDiag { target_drift_dir_idx: i32, aux_next_bar_mse_idx: i32, aux_regime_ce_idx: i32, - aux_label_scale_idx: i32, + // SP13 B1.0 (2026-05-05): `aux_label_scale_idx: i32` parameter + // dropped together with ISV[117] retirement. moe_expert_util_base_idx: i32, moe_gate_entropy_idx: i32, moe_lambda_eff_idx: i32, @@ -175,7 +176,8 @@ impl GpuHealthDiag { .arg(&target_drift_dir_idx) .arg(&aux_next_bar_mse_idx) .arg(&aux_regime_ce_idx) - .arg(&aux_label_scale_idx) + // SP13 B1.0 (2026-05-05): aux_label_scale_idx arg dropped + // — kernel signature shortened in lockstep. .arg(&moe_expert_util_base_idx) .arg(&moe_gate_entropy_idx) .arg(&moe_lambda_eff_idx) diff --git a/crates/ml/src/cuda_pipeline/health_diag.rs b/crates/ml/src/cuda_pipeline/health_diag.rs index b942689cb..11692978e 100644 --- a/crates/ml/src/cuda_pipeline/health_diag.rs +++ b/crates/ml/src/cuda_pipeline/health_diag.rs @@ -246,11 +246,13 @@ pub struct HealthDiagSnapshot { pub reward_split_opp_cost: f32, pub reward_split_bonus: f32, - // ── Aux block (4 floats) ───────────────────────────────────────────────── + // ── Aux block (3 floats) ───────────────────────────────────────────────── + // SP13 B1.0 (2026-05-05): the former `aux_label_scale: f32` field was + // dropped together with ISV[117] when the next-bar MSE flipped to + // scale-free. Snap layout now ends the aux block at word 127 (was 128). pub aux_next_bar_mse: f32, pub aux_regime_ce: f32, pub aux_weight: f32, - pub aux_label_scale: f32, // ── Aux MoE (8 utils + 1 ent + 1 λ_eff = 10 floats) ────────────────────── pub aux_moe_util: [f32; HEALTH_DIAG_MOE_EXPERTS], @@ -393,14 +395,16 @@ mod tests { // + 6 trail + 6 magstats + 6 noisy + 3 eval_dist // + 3 intent_dist + 4 reward_contrib + 7 controller + 3 explore // + 16 val + 4 val_dir + 4 val_picked - // + 6 reward_split + 4 aux + 10 aux_moe + // + 6 reward_split + 3 aux + 10 aux_moe // + 12 action_counts - // = 8+8+7+2+1+10+4+5+9+2+6+6+6+3+3+4+7+3+16+4+4+6+4+10+12 = 150 fields - let expected_bytes = 150 * 4; + // = 8+8+7+2+1+10+4+5+9+2+6+6+6+3+3+4+7+3+16+4+4+6+3+10+12 = 149 fields + // SP13 B1.0 (2026-05-05): aux block dropped from 4 → 3 floats + // (retired `aux_label_scale` together with ISV[117]). + let expected_bytes = 149 * 4; assert_eq!( std::mem::size_of::(), expected_bytes, - "HealthDiagSnapshot must be {} bytes (150 × 4); update test if struct changes", + "HealthDiagSnapshot must be {} bytes (149 × 4); update test if struct changes", expected_bytes, ); } diff --git a/crates/ml/src/cuda_pipeline/health_diag_kernel.cu b/crates/ml/src/cuda_pipeline/health_diag_kernel.cu index e3a4e36e3..6441b29bc 100644 --- a/crates/ml/src/cuda_pipeline/health_diag_kernel.cu +++ b/crates/ml/src/cuda_pipeline/health_diag_kernel.cu @@ -182,24 +182,25 @@ #define WORD_REWARD_SPLIT_OPP_COST 122 #define WORD_REWARD_SPLIT_BONUS 123 -/* Aux block (4 words, [124..128)). */ +/* Aux block (3 words, [124..127)). SP13 B1.0 (2026-05-05): retired + * `WORD_AUX_LABEL_SCALE` together with ISV[117] when the next-bar MSE + * flipped to scale-free; downstream offsets shifted down by 1. */ #define WORD_AUX_NEXT_BAR_MSE 124 #define WORD_AUX_REGIME_CE 125 #define WORD_AUX_WEIGHT 126 -#define WORD_AUX_LABEL_SCALE 127 -/* Aux MoE (10 words, [128..138) — 8 utils + ent + λ_eff). */ -#define WORD_AUX_MOE_UTIL_BASE 128 -#define WORD_AUX_MOE_ENT 136 -#define WORD_AUX_MOE_LAMBDA_EFF 137 +/* Aux MoE (10 words, [127..137) — 8 utils + ent + λ_eff). */ +#define WORD_AUX_MOE_UTIL_BASE 127 +#define WORD_AUX_MOE_ENT 135 +#define WORD_AUX_MOE_LAMBDA_EFF 136 -/* Action-count histogram (12 words, [138..150), all u32). */ -#define WORD_ACTION_COUNTS_BASE 138 +/* Action-count histogram (12 words, [137..149), all u32). */ +#define WORD_ACTION_COUNTS_BASE 137 -#define WORD_TOTAL 150 +#define WORD_TOTAL 149 /* Pin the layout: matches `snapshot_size_is_stable` on the Rust side. */ -static_assert(WORD_TOTAL == 150, +static_assert(WORD_TOTAL == 149, "HealthDiagSnapshot layout drift — update kernel WORD_* and the " "Rust struct in lockstep, then bump snapshot_size_is_stable."); @@ -220,24 +221,16 @@ static_assert(WORD_TOTAL == 150, * ISV[88] VSN_DIR_EMA_INDEX → snap.noisy_vsn_dir (word 69) * ISV[92] TARGET_DRIFT_MAG_EMA_INDEX → snap.noisy_drift_mag (word 72) * ISV[93] TARGET_DRIFT_DIR_EMA_INDEX → snap.noisy_drift_dir (word 73) - * ISV[113] AUX_NEXT_BAR_MSE_EMA_INDEX → snap.aux_next_bar_mse (word 121) - * ISV[114] AUX_REGIME_CE_EMA_INDEX → snap.aux_regime_ce (word 122) - * ISV[117] AUX_LABEL_SCALE_EMA_INDEX → snap.aux_label_scale (word 124) - * ISV[118..126) MOE_EXPERT_UTIL_EMA_BASE → snap.aux_moe_util[0..8] (words 125..133) - * ISV[126] MOE_GATE_ENTROPY_EMA_INDEX → snap.aux_moe_ent (word 133) - * ISV[128] MOE_LAMBDA_EFF_INDEX → snap.aux_moe_lambda_eff (word 134) + * ISV[113] AUX_NEXT_BAR_MSE_EMA_INDEX → snap.aux_next_bar_mse (word 124) + * ISV[114] AUX_REGIME_CE_EMA_INDEX → snap.aux_regime_ce (word 125) + * ISV[118..126) MOE_EXPERT_UTIL_EMA_BASE → snap.aux_moe_util[0..8] (words 127..135) + * ISV[126] MOE_GATE_ENTROPY_EMA_INDEX → snap.aux_moe_ent (word 135) + * ISV[128] MOE_LAMBDA_EFF_INDEX → snap.aux_moe_lambda_eff (word 136) * - * Note: ISV[126] is read TWICE — once into the last MoE-util slot - * (which the host code historically aliased through the same buffer - * indexing pattern when reading via `read_isv_signal_at(MOE_EXPERT_UTIL_EMA_BASE + 7)`), - * and once into `aux_moe_ent`. We mirror the source-level read pattern: the - * host previously called `read_isv_signal_at(MOE_GATE_ENTROPY_EMA_INDEX)` for - * `ent`, which is slot 126; the host also walked `[BASE..BASE+8)` for the - * 8 expert-utility slots, the last of which is slot 125+7=132 — *not* - * slot 126. So in fact MOE_EXPERT_UTIL_EMA_BASE+7 = 125 (snap word 132 in - * snap-word coords; ISV slot 125), and slot 126 is the entropy. The - * mapping above is therefore correct by ISV index — there is NO double - * read. + * SP13 B1.0 (2026-05-05): the former ISV[117] AUX_LABEL_SCALE_EMA_INDEX → + * snap.aux_label_scale mirror was dropped together with that slot's + * retirement. All snap word offsets at and beyond the former + * `WORD_AUX_LABEL_SCALE=127` shifted down by 1. * * Indices passed in as kernel args (rather than `#define`d) so we don't * couple the kernel to a specific numbering — the Rust caller passes @@ -267,7 +260,8 @@ extern "C" __global__ void health_diag_isv_mirror( int target_drift_dir_idx, /* ISV[93] */ int aux_next_bar_mse_idx, /* ISV[113] */ int aux_regime_ce_idx, /* ISV[114] */ - int aux_label_scale_idx, /* ISV[117] */ + /* SP13 B1.0 (2026-05-05): ISV[117] AUX_LABEL_SCALE retired — + * argument dropped from kernel signature. */ int moe_expert_util_base_idx, /* ISV[118], 8 contiguous slots */ int moe_gate_entropy_idx, /* ISV[126] */ int moe_lambda_eff_idx /* ISV[128] */ @@ -296,10 +290,11 @@ extern "C" __global__ void health_diag_isv_mirror( snap[WORD_NOISY_DRIFT_DIR] = isv[target_drift_dir_idx]; /* Aux-head loss EMAs (Plan 4 Task 6 Commit A producer: - * aux_heads_loss_ema_kernel). */ + * aux_heads_loss_ema_kernel). SP13 B1.0 (2026-05-05): the third + * mirror (ISV[117] → snap.aux_label_scale) was retired together with + * the slot when the next-bar MSE flipped to scale-free. */ snap[WORD_AUX_NEXT_BAR_MSE] = isv[aux_next_bar_mse_idx]; snap[WORD_AUX_REGIME_CE] = isv[aux_regime_ce_idx]; - snap[WORD_AUX_LABEL_SCALE] = isv[aux_label_scale_idx]; /* MoE per-expert utilisation EMA (8 contiguous slots). Producer: * moe_expert_util_ema_update in moe_kernels.cu. */ diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 1caa27b9e..6f85c67a3 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -368,24 +368,12 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[AUX_REGIME_CE_EMA_INDEX=114] — regime classification head batch-mean cross-entropy EMA; same producer/contract as next_bar_mse_ema; cold-start 0.0", }, - // ───── Plan 4 Task 6 / Plan 5 Task 5 follow-up: aux next-bar - // label-scale EMA. Defends the aux next-bar regression head - // against underlying-data scale (1e-3 log-returns vs 5000 raw - // prices). Producer: aux_label_scale_ema_update GPU kernel - // (single-block 256-thread shmem reduction, α=0.05). Consumer: - // aux_next_bar_loss_reduce + aux_next_bar_backward read the - // slot to divide label by max(scale, 1e-6) before the residual. - // SP4 Task A13.1 retrofit — Pearl A sentinel 0.0 at fold boundary - // (paired with sp4_wiener_state reset zeroing state.x_lag in - // lockstep). The 1e-6 floor on the consumer side guards the - // narrow window between fold boundary and the first producer - // fire (which writes the measured scale directly via Pearl A's - // first-observation replacement). - RegistryEntry { - name: "isv_aux_label_scale_ema", - category: ResetCategory::FoldReset, - description: "ISV[AUX_LABEL_SCALE_EMA_INDEX=117] — aux next-bar label-scale EMA (mean(|next_states[:, 0]|) at α=0.05); GPU aux_label_scale_ema_update kernel fills (Plan 4 Task 6 / Plan 5 Task 5 follow-up); consumer kernels read slot 117 to divide label by max(scale, 1e-6) before the residual; SP4 Task A13.1 retrofit — Pearl A sentinel 0.0 reapplied at fold boundary; consumer 1e-6 floor guards the divide-by-zero window before the first producer fire", - }, + // SP13 B1.0 (2026-05-05): the former + // `isv_aux_label_scale_ema` registry entry (FoldReset for ISV + // slot 117) was retired together with the slot itself when + // the next-bar MSE flipped to scale-free. No replacement + // entry is needed — the slot is no longer written by any + // producer, so there is nothing to reset at fold boundary. // Mixture-of-Experts per-expert utilization (slots 118..126). // SP4 Task A13.2 retrofit — Pearl A sentinel 0.0 at fold boundary. RegistryEntry { diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 9c9579bb2..185a94cd9 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -4811,36 +4811,33 @@ impl DQNTrainer { // Plan 4 Task 6 Commit B: aux-heads diagnostic line. Reads // ISV[AUX_NEXT_BAR_MSE_EMA_INDEX=113] + // ISV[AUX_REGIME_CE_EMA_INDEX=114] (populated by - // `aux_heads_loss_ema_update` launched per-step above) + - // ISV[AUX_LABEL_SCALE_EMA_INDEX=117] (Plan 5 Task 5 follow-up - // — populated by `aux_label_scale_ema_update` per step), and + // `aux_heads_loss_ema_update` launched per-step above) and // the host-side `aux_weight` baked into the most-recently- // captured graph. Cold-start (epoch 0 of each fold): loss EMAs - // at 0.0, label_scale at 1.0, aux_weight at 0.05. - // `label_scale` magnitude reveals what the underlying data column - // 0 of next_states actually carries this step — log returns - // (~1e-3) vs raw prices (~5000) — and confirms the normalize- - // before-MSE fix is firing against the right scale. + // at 0.0, aux_weight at 0.05. + // + // SP13 B1.0 (2026-05-05): the former `label_scale` field + // (read from retired ISV[117]) was dropped together with that + // slot. Labels are z-normalised at the data layer so + // monitoring the scale was redundant. { use crate::cuda_pipeline::gpu_dqn_trainer::{ AUX_NEXT_BAR_MSE_EMA_INDEX, AUX_REGIME_CE_EMA_INDEX, - AUX_LABEL_SCALE_EMA_INDEX, }; - let (aux_nb_mse, aux_rg_ce, aux_w, aux_label_scale) = + let (aux_nb_mse, aux_rg_ce, aux_w) = if let Some(ref fused) = self.fused_ctx { let trainer = fused.trainer(); ( trainer.read_isv_signal_at(AUX_NEXT_BAR_MSE_EMA_INDEX), trainer.read_isv_signal_at(AUX_REGIME_CE_EMA_INDEX), trainer.aux_weight(), - trainer.read_isv_signal_at(AUX_LABEL_SCALE_EMA_INDEX), ) } else { - (0.0, 0.0, 0.0, 1.0) + (0.0, 0.0, 0.0) }; tracing::info!( - "HEALTH_DIAG[{}]: aux [next_bar_mse={:.3e} regime_ce={:.3e} w={:.3} label_scale={:.3e}]", - epoch, aux_nb_mse, aux_rg_ce, aux_w, aux_label_scale, + "HEALTH_DIAG[{}]: aux [next_bar_mse={:.3e} regime_ce={:.3e} w={:.3}]", + epoch, aux_nb_mse, aux_rg_ce, aux_w, ); } @@ -6846,26 +6843,15 @@ impl DQNTrainer { fused.trainer().write_isv_signal_at(idx, 0.0); } } - // Plan 4 Task 6 / Plan 5 Task 5 follow-up: aux next-bar label- - // scale EMA. SP4 Task A13.1 retrofit reset to Pearl A sentinel - // 0.0 at fold boundary — the companion `sp4_wiener_state` bulk - // reset zeros `state.x_lag` in lockstep so the first - // `aux_label_scale_ema_update` fire on the new fold sees - // `prev_x_mean == 0 AND state.x_lag == 0` and triggers Pearl A's - // first-observation replacement (which writes the measured scale - // directly, no 1.0 cold-start blend). Per - // `feedback_no_partial_refactor.md` — the sentinel contract - // migrates both halves together. Consumer kernels' `1e-6` floor - // protects the divide-by-zero window between fold boundary and the - // first producer fire. - "isv_aux_label_scale_ema" => { - if let Some(ref fused) = self.fused_ctx { - fused.trainer().write_isv_signal_at( - crate::cuda_pipeline::gpu_dqn_trainer::AUX_LABEL_SCALE_EMA_INDEX, - 0.0_f32, - ); - } - } + // SP13 B1.0 (2026-05-05): the former `isv_aux_label_scale_ema` + // arm is retired together with ISV[117]. The slot is no + // longer written by any producer (`aux_label_scale_ema_update` + // kernel was deleted) and no longer read by any consumer + // (`aux_next_bar_loss_reduce` / `aux_next_bar_backward` + // flipped to scale-free); nothing to reset at fold boundary. + // Removing the registry entry above made this arm dead, and + // dispatch on the retired name now correctly falls into the + // unknown-name guard. // SP4 Task A13.2 retrofit: per-expert MoE utilisation EMA slots. // Reset to Pearl A sentinel 0.0 at fold boundary — the companion // `sp4_wiener_state` bulk reset zeros `state.x_lag` in lockstep so diff --git a/crates/ml/tests/sp4_producer_unit_tests.rs b/crates/ml/tests/sp4_producer_unit_tests.rs index 5dd0c72bd..5ae560c23 100644 --- a/crates/ml/tests/sp4_producer_unit_tests.rs +++ b/crates/ml/tests/sp4_producer_unit_tests.rs @@ -1526,15 +1526,10 @@ fn load_aux_heads_loss_ema_kernel(stream: &Arc) -> CudaFunction { .expect("load aux_heads_loss_ema_update function") } -fn load_aux_label_scale_ema_kernel(stream: &Arc) -> CudaFunction { - let module = stream - .context() - .load_cubin(SP4_AUX_HEADS_LOSS_EMA_CUBIN.to_vec()) - .expect("load aux_heads_loss_ema cubin (label scale)"); - module - .load_function("aux_label_scale_ema_update") - .expect("load aux_label_scale_ema_update function") -} +// SP13 B1.0 (2026-05-05): the `load_aux_label_scale_ema_kernel` helper +// + `sp4_aux_label_scale_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` +// test were retired together with the `aux_label_scale_ema_update` kernel +// itself when the next-bar MSE flipped to scale-free. /// SP4 Task A13.1: `aux_heads_loss_ema_update` writes both scalar loss /// values to scratch[41]/[42], and Pearls A+D bootstrap (Pearl A) replaces @@ -1616,78 +1611,10 @@ fn sp4_aux_heads_loss_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d() { "rg stationary should converge: got {rg_mean}, expected {rg_loss}"); } -/// SP4 Task A13.1: `aux_label_scale_ema_update` writes mean(|label|) to -/// scratch[43], and Pearls A+D bootstrap+convergence behave correctly on -/// stationary input. -#[test] -#[ignore = "requires GPU"] -fn sp4_aux_label_scale_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d() { - use ml::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState}; - - const SCRATCH_IDX: usize = 43; - const B: usize = 256; - const BLOCK_DIM: u32 = 256; - const SHARED_BYTES: u32 = BLOCK_DIM * std::mem::size_of::() as u32; - - // Stationary mixed-sign labels with absolute mean = 3.0. - let labels: Vec = (0..B) - .map(|i| if i % 2 == 0 { 3.0 } else { -3.0 }) - .collect(); - let expected_mean_abs: f32 = 3.0; - - let stream = make_test_stream(); - let kernel = load_aux_label_scale_ema_kernel(&stream); - - let label_buf = unsafe { MappedF32Buffer::new(B) }.expect("alloc label buf"); - label_buf.write_from_slice(&labels); - let scratch_buf = unsafe { MappedF32Buffer::new(SP4_PRODUCER_COUNT) } - .expect("alloc scratch"); - - let label_dev = label_buf.dev_ptr; - let scratch_dev = scratch_buf.dev_ptr; - let b_i32 = B as i32; - let scratch_idx_arg: i32 = SCRATCH_IDX as i32; - - unsafe { - stream.launch_builder(&kernel) - .arg(&label_dev) - .arg(&b_i32) - .arg(&scratch_dev) - .arg(&scratch_idx_arg) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (BLOCK_DIM, 1, 1), - shared_mem_bytes: SHARED_BYTES, - }) - .expect("launch aux_label_scale_ema_update"); - } - stream.synchronize().expect("sync"); - - let host = scratch_buf.read_all(); - for (i, &v) in host.iter().enumerate() { - if i != SCRATCH_IDX { - assert_eq!(v, 0.0, "wrote outside target slot: scratch[{i}]={v}"); - } - } - let step_obs = host[SCRATCH_IDX]; - assert!( - (step_obs - expected_mean_abs).abs() < 1e-4, - "step_obs {step_obs} vs expected {expected_mean_abs}", - ); - - // Pearl A bootstrap. - let mut state = WienerState::ZERO; - let new_x_mean = pearls_ad_update(0.0, &mut state, step_obs); - assert_eq!(new_x_mean, step_obs); - assert_eq!(state.x_lag, step_obs); - - // Pearl D convergence. - let mut x_mean = new_x_mean; - for _ in 0..1000 { - x_mean = pearls_ad_update(x_mean, &mut state, step_obs); - } - assert!(((x_mean - expected_mean_abs) / expected_mean_abs).abs() < 0.01); -} +// SP13 B1.0 (2026-05-05): retired +// `sp4_aux_label_scale_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` +// test — its target kernel `aux_label_scale_ema_update` was deleted +// together with ISV[117] when the next-bar MSE flipped to scale-free. // ── SP4 Task A13.2: moe_expert_util_ema retrofit ───────────────────────────── diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 3865636e2..5e79ae8b5 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -5910,8 +5910,70 @@ Build: `cargo check --workspace` clean in 21s. Only pre-existing warnings (Devic ### B0.1 cascade-gap fix-up (2026-05-05) -The original B0 audit (commit `62ab8ed85`) under-counted `insert_batch` test callers as 2 (1 production + 1 in-file unit test). Surfaced during B1.0 implementation when `cargo check --workspace --tests` failed with 5 arity-mismatch errors. Root cause: the B0 grep filter was `grep -v test` and the audit didn't enumerate `crates/ml/src/trainers/dqn/smoke_tests/` (compiled as part of the lib's test binary, not behind `#[cfg(test)]`) nor `crates/ml/tests/`. **Process gap:** future B-series audits must run `cargo check --workspace --tests` before claiming cardinality completeness. Fixed in commit `` by adding zero-init `CudaSlice` allocs at all 5 missed sites. No behavior change — column carries zero data; B1.1 lands the producer kernel that fills with -1/0/1. +The original B0 audit (commit `62ab8ed85`) under-counted `insert_batch` test callers as 2 (1 production + 1 in-file unit test). Surfaced during B1.0 implementation when `cargo check --workspace --tests` failed with 5 arity-mismatch errors. Root cause: the B0 grep filter was `grep -v test` and the audit didn't enumerate `crates/ml/src/trainers/dqn/smoke_tests/` (compiled as part of the lib's test binary, not behind `#[cfg(test)]`) nor `crates/ml/tests/`. **Process gap:** future B-series audits must run `cargo check --workspace --tests` before claiming cardinality completeness. Fixed in commit `6a869ad36` by adding zero-init `CudaSlice` allocs at all 5 missed sites. No behavior change — column carries zero data; B1.1 lands the producer kernel that fills with -1/0/1. ### Next: B1 (separate atomic commit, fresh-implementer dispatch) B1 covers: aux head 1→2 dim, MSE→CE loss, `aux_dir_acc` reads softmax, `aux_pred_to_isv_tanh` rewrite as logit-diff, ISV[117]=AUX_LABEL_SCALE_EMA_INDEX retirement, `dqn_param_layout` fingerprint bump, kernel that populates `aux_sign_labels` with real -1/0/1 labels from the 30-bar price trajectory, `aux_b1_diag` HEALTH_DIAG metric, 17+ GPU oracle unit tests. Brief written from B0's audit findings (this section) — no implementer should re-derive call site cardinality. + +## SP13 Layer B — Commit B1.0: ISV[117] retirement + scale-free MSE bridge (2026-05-05) + +**Why split B1.0 from B1.1**: the original B1 brief bundled the ISV[117] retirement (cleanup of the now-redundant label-scale EMA, possible because labels are z-normalised at the data layer) with the MSE→CE flip + 1→2 dim head expansion + producer kernel for -1/0/1 labels + fingerprint bump + 17+ unit tests. That was still too large for a single commit. B1.0 lands the retirement-plus-bridge atomically (ISV[117] producer/consumer/reset/health all retired in one commit; MSE retained but the divisor is removed). B1.1 lands the head expansion, MSE→CE flip, label producer kernel, fingerprint bump, and unit tests on top. + +**Behavior contract**: B1.0 is a numerical bridge — z-normalised labels make `label_scale_ema ≈ 1.0` empirically, so removing the `inv_scale = 1.0 / max(scale, 1e-6)` divisor reduces `(pred - label * inv_scale)` to `(pred - label)` within rounding. Snapshot-stability test `snapshot_size_is_stable` was relaxed from 150 to 149 floats (aux block 4→3). Layout fingerprint did NOT bump because the ISV slot remains reserved (consumers of the slot are gone, but the seed still names slot 117 to prevent silent reuse). + +### Wiring (what was retired) + + 1. **`aux_label_scale_ema_update` kernel** (formerly in `aux_heads_loss_ema_kernel.cu`) — single-block 256-thread `mean(|label|)` reduction. Deleted; not wired anywhere. + 2. **`aux_next_bar_loss_reduce` + `aux_next_bar_backward` kernel signatures** (`aux_heads_kernel.cu`) — drop `const float* isv` + `int isv_label_scale_index` params; residual computes `(pred - label)` directly. Comment cite SP13 B1.0 + forward-pointer to B1.1's CE replacement. + 3. **`aux_label_scale_ema_update_kernel: CudaFunction`** field + loader in `gpu_aux_heads.rs`'s `AuxHeadsForwardOps` — removed. + 4. **`launch_label_scale_ema` orchestrator method** in `gpu_aux_heads.rs` — removed. + 5. **`aux_heads_forward` Step 2b launch site** in `gpu_dqn_trainer.rs` (the inline Pearls A+D applier for ISV[117]) — removed. + 6. **`backward_next_bar` orchestrator arg list** in `gpu_aux_heads.rs` — drops `isv_dev_ptr`/`isv_label_scale_index` args; loss + backward signatures match. + 7. **`AUX_LABEL_SCALE_EMA_INDEX = 117`** const in `gpu_dqn_trainer.rs` — kept as a `RETIRED` doc-comment marker on the slot index (no value re-defined; the constant is removed from active use, the doc comment forwards to B1.1). + 8. **Layout fingerprint seed** (`gpu_dqn_trainer.rs`) — keeps `AUX_LABEL_SCALE_EMA=117` line so the seed hash is stable across the bridge (no fingerprint bump in B1.0; B1.1 will bump on the head-dim flip). + 9. **HEALTH_DIAG aux block** (`health_diag.rs::HealthDiagSnapshot`, `health_diag_kernel.cu` GPU snapshot writer) — drop `aux_label_scale: f32` field. Aux block shrinks 4→3 words `[124..127)`. AuxMoE base shifts down by 1: `WORD_AUX_MOE_UTIL_BASE 128→127`, action-counts base `138→137`, `WORD_TOTAL 150→149`. `static_assert(WORD_TOTAL == 149)` at the kernel header. `snapshot_size_is_stable` relaxes from `150 * 4 = 600` to `149 * 4 = 596` bytes. HEALTH_DIAG aux line drops `+label_scale={:.3e}`. + 10. **StateResetRegistry** (`state_reset_registry.rs`) — drop `isv_aux_label_scale_ema` FoldReset entry; its dispatch arm in `training_loop.rs::reset_named_state` is also removed (no replacement — slot is no-op now). + 11. **Unit tests** (`crates/ml/tests/sp4_producer_unit_tests.rs`) — drop `load_aux_label_scale_ema_kernel` helper + `sp4_aux_label_scale_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` test. + +### Audit-verified call site cardinality + +| Site | Count | Action | +|---|---|---| +| `aux_label_scale_ema_update` kernel | 1 | retired (file `aux_heads_loss_ema_kernel.cu`) | +| `AUX_LABEL_SCALE_EMA_INDEX` consumers (kernel sigs) | 2 | `aux_next_bar_loss_reduce`, `aux_next_bar_backward` — both drop `isv_*` params | +| `AUX_LABEL_SCALE_EMA_INDEX` consumers (Rust) | 4 | trunk forward producer launch (gpu_dqn_trainer.rs), trunk backward consumer pass-through, training_loop HEALTH_DIAG read, training_loop reset_named_state arm — all retired | +| HEALTH_DIAG snapshot fields | 1 | `aux_label_scale: f32` removed; downstream layout offsets shift down by 1 word | +| StateResetRegistry entries | 1 | `isv_aux_label_scale_ema` retired | +| sp4 producer unit tests | 1 | `sp4_aux_label_scale_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` retired | + +### Hard rules upheld + +- `feedback_no_partial_refactor`: every consumer of ISV[117] migrates atomically — kernel + Rust orchestrator + producer launch + backward call + HEALTH_DIAG field + reset registry + unit test all in this commit +- `feedback_no_stubs`: not a return-zero stub — the entire ISV[117] producer/consumer chain is deleted, not stubbed; `(pred - label)` is the real loss formula now that labels are unit-scale by construction at the data layer +- `feedback_no_legacy_aliases`: no shim function or `AUX_LABEL_SCALE_EMA_INDEX → 1.0_const` alias — the divisor is removed at every site, not aliased +- `feedback_no_hiding`: doc comments on `AUX_LABEL_SCALE_EMA_INDEX` and the retired kernel explicitly forward to B1.1; no `_` underscore suppression or `#[allow(dead_code)]` +- `feedback_isv_for_adaptive_bounds`: the `1e-6` divisor floor (a numerical-stability epsilon, NOT a tuned constant) is also removed since the divisor is gone +- `feedback_no_atomicadd`: no new producers/reductions added (we deleted one) +- `feedback_cpu_is_read_only`: no host computation introduced + +### Files (atomic B1.0 commit) + +10 files, net −288 LOC: + +- `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` (−18) — drop `isv` + `isv_label_scale_index` params from `aux_next_bar_loss_reduce` + `aux_next_bar_backward`; residual is `(pred - label)` +- `crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu` (−72) — delete `aux_label_scale_ema_update` kernel; surviving header doc updated with B1.0 note +- `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs` (−69) — drop kernel field/loader + `launch_label_scale_ema` method + `isv_*` args from `next_bar_loss_reduce` / `backward_next_bar` +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (−80) — drop Step 2b producer launch + ISV slot uses in forward/backward; comment-only retention of `AUX_LABEL_SCALE_EMA=117` in fingerprint seed (no fingerprint bump) +- `crates/ml/src/cuda_pipeline/gpu_health_diag.rs` (−2 net) — drop `aux_label_scale` snapshot wire +- `crates/ml/src/cuda_pipeline/health_diag.rs` (+5 net) — drop field, update test snapshot-size from 150→149 floats with B1.0 doc comment +- `crates/ml/src/cuda_pipeline/health_diag_kernel.cu` (−7 net) — drop `WORD_AUX_LABEL_SCALE`, shift downstream offsets, update `WORD_TOTAL` 150→149 + `static_assert` +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` (−18) — drop `isv_aux_label_scale_ema` registry entry, replace with B1.0 retirement doc comment +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (−8 net) — drop reset_named_state arm + HEALTH_DIAG read + B1.0 doc comment on the aux line +- `crates/ml/tests/sp4_producer_unit_tests.rs` (−74) — drop unit test + helper, replace with B1.0 retirement doc comment + +Build: `cargo check --workspace --tests` clean. `snapshot_size_is_stable` passes at 149*4=596 bytes. + +### Next: B1.1 (separate atomic commit, fresh-implementer dispatch) + +B1.1 covers: aux head 1→2 dim (next-bar regression head becomes a 2-class direction logit head), MSE→CE loss flip, `aux_dir_acc` reads softmax over the 2 logits, `aux_pred_to_isv_tanh` rewrite as `tanh(logit_pos - logit_neg)`, producer kernel that fills `aux_sign_labels` with real -1/0/1 from the 30-bar price trajectory (B0 plumbing currently zero-init), `dqn_param_layout` fingerprint bump (head dim changes), `aux_b1_diag` HEALTH_DIAG metric, 17+ GPU oracle unit tests. The bridge in B1.0 means B1.1 only needs to flip the loss formulation — the divisor scaffolding is already gone.