feat(sp13): B1.0 — ISV[117] retirement + scale-free MSE bridge
Retires ISV[117]=AUX_LABEL_SCALE_EMA_INDEX together with its producer
kernel (aux_label_scale_ema_update), launch site, backward pass-through,
StateResetRegistry entry, HEALTH_DIAG snapshot field, and unit test.
Why: labels at the data layer are z-normalised, so the
mean(|label|) EMA tracked by ISV[117] sits at ~1.0 empirically.
Dividing by max(scale, 1e-6) before the residual `(pred - label)`
reduces to `(pred - label)` within rounding. The divisor was a
defensive scaffold from when the data layer carried mixed-scale
labels (1e-3 log returns vs 5000 raw prices); z-normalisation made
that scaffold redundant.
This is a numerical bridge, NOT the final fix. B1.1 lands on top:
- Aux head 1→2 dim (next-bar regression → 2-class direction logit)
- MSE → CE loss flip
- aux_dir_acc reads softmax over the 2 logits
- aux_pred_to_isv_tanh rewrite as logit-diff
- 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
Cascade (atomic per feedback_no_partial_refactor):
- aux_heads_kernel.cu: aux_next_bar_loss_reduce + aux_next_bar_backward
drop `isv` + `isv_label_scale_index` params; residual is (pred - label)
- aux_heads_loss_ema_kernel.cu: aux_label_scale_ema_update kernel deleted
- gpu_aux_heads.rs: kernel field/loader + launch_label_scale_ema +
isv_* args from next_bar_loss_reduce / backward_next_bar all dropped
- gpu_dqn_trainer.rs: Step 2b producer launch + ISV slot uses dropped;
AUX_LABEL_SCALE_EMA=117 line retained in fingerprint seed
(no fingerprint bump in B1.0; B1.1 will bump on head-dim flip)
- gpu_health_diag.rs + health_diag.rs: aux_label_scale snapshot field
dropped; aux block 4→3 floats, downstream offsets shift down by 1,
WORD_TOTAL 150→149, snapshot_size_is_stable test 150*4 → 149*4
- health_diag_kernel.cu: WORD_AUX_LABEL_SCALE removed, downstream
offsets shift, static_assert(WORD_TOTAL == 149)
- state_reset_registry.rs: isv_aux_label_scale_ema FoldReset dropped
- training_loop.rs: reset_named_state arm + HEALTH_DIAG read +
aux line label_scale field all dropped
- sp4_producer_unit_tests.rs: load_aux_label_scale_ema_kernel helper +
sp4_aux_label_scale_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d
test dropped
Hard rules upheld:
- feedback_no_partial_refactor: every consumer of ISV[117] migrates
atomically — kernel + Rust orchestrator + producer launch + backward
+ HEALTH_DIAG + reset registry + unit test all in this commit
- feedback_no_stubs: not a stub — divisor is removed at every site,
not aliased through a 1.0_const shim
- feedback_no_legacy_aliases: no legacy AUX_LABEL_SCALE_EMA_INDEX → 1.0
alias function
- feedback_no_hiding: doc comments forward to B1.1 explicitly; no
underscore suppression or #[allow(dead_code)]
Build: cargo check --workspace --tests clean.
Tests: snapshot_size_is_stable passes at 149*4=596 bytes.
cargo test -p ml --lib + cargo test -p ml-dqn --lib compile.
Net delta: 10 files, −288 LOC.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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) {
|
||||
|
||||
@@ -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();
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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<CudaStream>,
|
||||
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<CudaStream>,
|
||||
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::<f32>() 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)
|
||||
|
||||
@@ -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(),
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -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::<HealthDiagSnapshot>(),
|
||||
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,
|
||||
);
|
||||
}
|
||||
|
||||
@@ -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. */
|
||||
|
||||
@@ -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 {
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -1526,15 +1526,10 @@ fn load_aux_heads_loss_ema_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
|
||||
.expect("load aux_heads_loss_ema_update function")
|
||||
}
|
||||
|
||||
fn load_aux_label_scale_ema_kernel(stream: &Arc<CudaStream>) -> 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::<f32>() as u32;
|
||||
|
||||
// Stationary mixed-sign labels with absolute mean = 3.0.
|
||||
let labels: Vec<f32> = (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 ─────────────────────────────
|
||||
|
||||
|
||||
@@ -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 `<TBD>` by adding zero-init `CudaSlice<i32>` 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<i32>` 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.
|
||||
|
||||
Reference in New Issue
Block a user