gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

feat(sp13): B1.1a — K=1→2 + softmax CE kernel rewrites + struct flips

Browse Source 
Flips the aux next-bar head from K=1 MSE regression to K=2 softmax
cross-entropy classification. Kernel ABIs, struct fields, partial-buf
shapes, and per-step ISV producers all migrate atomically; the producer
that fills `aux_sign_labels` with real -1/0/1 from the price trajectory
lands separately in B1.1b.

Why split B1.1a from B1.1b: the original B1 brief was decomposed (B1.0
+ B1.1) after six full-B1 dispatches confirmed agent-session capacity
is the bottleneck, not cascade understanding. B1.1 itself is now further
split into B1.1a (kernel/struct/test cascade — this commit) and B1.1b
(producer kernel + replay direct path + experience collector hoist +
remaining tests + Smoke A). B1.1a is contract-consistent atomic
kernel-side; B1.1b lands the producer + smoke.

Why labels stay zero-init: B1.1a is producer-less by design. The
`aux_nb_label_buf: CudaSlice<i32>` is `alloc_zeros<i32>` so every
sample receives label 0 ("down"). The model converges on "predict
class 0 (down) everywhere" until B1.1b lands the producer kernel that
fills real -1/0/1 from the 30-bar price trajectory. This degraded
training behavior is intentional and known — the cascade is internally
consistent (every consumer of K-flip / softmax tile / CE / i32 label
migrates atomically per `feedback_no_partial_refactor`); the labels are
placeholder. Local unit tests (CE correctness, dir_acc correctness,
isv_tanh correctness, fingerprint bump) validate B1.1a in isolation;
no L40S smoke runs between B1.1a and B1.1b.

Four contracts (atomic in this commit):
  1. K_NB flip 1 → 2: AUX_NEXT_BAR_K constant, compute_param_sizes
     ([121]/[122] grow), fingerprint seed rename
     (PARAM_AUX_NB_W2/B2 → PARAM_AUX_NB_W2_K2/B2_K2 — bumps the hash),
     forward + backward kernels, partial-buf allocs (nb_w2 [B,H]→[B,K,H],
     nb_b2 [B]→[B,K]), saxpy spec table, max_aux_tensor_len,
     aux_nb_pred_buf renamed to aux_nb_logits_buf per
     feedback_no_legacy_aliases.
  2. Softmax tile: aux_next_bar_forward writes [B, K] softmax via
     in-kernel stable softmax (max-shift form, K=2 single-thread fanout
     mirrors regime kernel); 4 consumers read the tile (loss, backward,
     dir-acc, isv-tanh). NEW field aux_nb_softmax_buf [B, K].
  3. MSE → CE: aux_next_bar_loss_reduce reads softmax + i32 labels,
     masks -1, divides by B_valid (mean-over-valid-rows), writes loss +
     B_valid scalar. aux_next_bar_backward reads B_valid so loss + grad
     share the same divisor — derivatives of the same scalar function.
     NEW field aux_nb_valid_count_buf [1]. All-skip batch produces
     loss = 0 (no NaN; fmaxf(valid, 1) divisor) and zero gradients.
     Numerical floor 1e-30 prevents -log(0) = +inf in extreme-logit path.
  4. i32 label dtype: aux_nb_label_buf flipped f32 → i32; -1 mask
     sentinel handled across loss + backward + dir-acc + isv-tanh.
     The strided_gather of next_states[:, 0] retired entirely.

Cascade (atomic per feedback_no_partial_refactor):
- aux_heads_kernel.cu: aux_next_bar_forward gains K + softmax tile
  output (via in-kernel stable softmax); aux_next_bar_loss_reduce
  ABI flipped (softmax + i32 labels, mean-over-valid CE,
  valid_count_out); aux_next_bar_backward ABI flipped (softmax +
  i32 labels + valid_count, K-fanout d_logits, masked rows zero
  across the K-vector)
- aux_dir_acc_reduce_kernel.cu: read softmax + i32 labels, argmax
  over K, output grew 3 → 6 floats (added n_down/n_up/n_skip);
  shmem 4 → 6 int arrays
- aux_pred_to_isv_tanh_kernel.cu: read softmax tile, compute
  mean(softmax[:, 1] - softmax[:, 0]); tanh transcend retired
  (structural [-1, +1] bound via softmax components per
  pearl_bounded_modifier_outputs_require_structural_activation)
- gpu_aux_heads.rs: AUX_NEXT_BAR_K 1 → 2; forward_next_bar gains K +
  logits_out + softmax_out args; next_bar_loss_reduce gains K +
  valid_count_out; backward_next_bar gains softmax_in + labels_i32_in
  + valid_count_in + K
- gpu_dqn_trainer.rs: compute_param_sizes ([121]/[122]); fingerprint
  seed rename (W2/B2 → W2_K2/B2_K2); struct fields (logits, softmax,
  i32 label, valid_count); aux_dir_acc_buf 3 → 6 floats; partial-buf
  allocs grow; max_aux_tensor_len extended; saxpy spec table updated;
  orchestrator launchers (launch_aux_dir_acc_reduce,
  launch_aux_pred_to_isv_tanh, launch_sp13_aux_dir_metrics) gain K
  arg; strided_gather block deleted entirely
- training_loop.rs: aux_b1_diag HEALTH_DIAG line reads
  aux_dir_acc_buf [3..6] for n_down / n_up / n_skip + mask_frac;
  doc comment update for the per-step aux dir-metrics block
- tests/sp13_phase0_oracle_tests.rs: 6 dir_acc + 3 isv_tanh tests
  rewritten in-place to new ABI (no shadow tests, per
  feedback_no_legacy_aliases)
- tests/sp13_layer_b_oracle_tests.rs (NEW): 11 B1.1a tests — 5 CE
  loss/backward (single-row, batch-mixed, all-skip-NaN, backward
  single-row, backward batch-mixed); 2 dir_acc (handcrafted argmax,
  all-skip NaN-safe); 2 isv_tanh (bounded fuzz, mean handcrafted);
  2 layout regression (fingerprint bump, HEALTH_DIAG snap stable)
- docs/dqn-wire-up-audit.md: B1.1a section

Hard rules upheld:
- feedback_no_partial_refactor: every consumer of K-flip / softmax tile
  / CE / i32 labels migrates atomically — kernels + orchestrators +
  struct + diag + existing oracle tests
- feedback_no_atomicadd: block tree-reduce only; CE loss reduce uses
  2 parallel partial-reduction strips; CE backward uses existing
  per-sample partial → final aux_param_grad_reduce pattern
- feedback_cpu_is_read_only: aux_nb_label_buf is GPU-resident
  CudaSlice<i32>; HEALTH_DIAG aux_b1_diag reads via mapped-pinned
  aux_dir_acc_buf (no DtoH)
- feedback_no_stubs: every new buffer + kernel arg wired through to
  a real consumer; CE forward / loss / backward chain executes
  end-to-end against placeholder labels (degraded behavior, not stub)
- feedback_no_legacy_aliases: aux_nb_pred_buf renamed in-place
  (no shim); PARAM_AUX_NB_W2/B2 renamed to _W2_K2/_B2_K2 in seed
  (no _DEPRECATED alias); 9 existing oracle tests rewritten in-place
- feedback_no_cpu_test_fallbacks: 9 GPU tests gated #[ignore]; 2
  layout-regression tests are CPU-only (pub const + size_of)
- feedback_no_htod_htoh_only_mapped_pinned: every CPU↔GPU buffer
  in tests + production is MappedF32Buffer / MappedI32Buffer
- feedback_isv_for_adaptive_bounds: no hardcoded thresholds added
  (1e-30 numerical floor on log is stability epsilon, not tunable)
- feedback_trust_code_not_docs: 8/8 Phase 0 anchors verified at
  HEAD 75e94858c before editing
- pearl_bounded_modifier_outputs_require_structural_activation:
  softmax IS the structural activation; runtime tanh squash retired

Build: cargo check --workspace --tests clean.
Tests:
  - snapshot_size_is_stable passes at 149*4=596 bytes (B1.1a
    doesn't touch HEALTH_DIAG snap-words)
  - fingerprint_bumped_from_pre_b1_1a passes (proves W2/B2 rename
    flipped seed hash)
  - health_diag_snap_size_stable_at_149_floats passes
  - 9 GPU oracle tests in sp13_layer_b_oracle_tests pass on local
    RTX 3050 Ti (B=1/4)
  - 14 tests in sp13_phase0_oracle_tests pass (6 dir_acc + 3
    isv_tanh rewrites + 5 unchanged hold/EMA tests)

Pre-existing test failures (14) on `cargo test -p ml --lib` are
environmental (OFI test data missing, etc.) and reproduce identically
at HEAD 75e94858c (B1.0) — verified via git stash round-trip.

Net delta: 9 files (8 source + 1 audit doc), +1039 / −539 LOC.

Next: B1.1b — producer kernel + replay direct-path 8th gather +
experience collector hoist + remaining 7 tests (6 producer + 1
end-to-end round-trip) + Smoke A on L40S.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-05 14:09:48 +02:00
parent 75e94858c5
commit 7d10ea8b3e
9 changed files with 1865 additions and 544 deletions
Download Patch File Download Diff File Expand all files Collapse all files

179
crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu
View File

@@ -1,66 +1,76 @@
// crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu
//
// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy reducer.
// SP13 B1.1a (2026-05-05): ABI flipped — reads `softmax [B, K]` + i32
// labels (replacing the regression-mode `(aux_pred f32, label_sign f32)`
// pair). Output grew from `[3]` to `[6]` to expose the per-class counts
// the new diagnostic line (`HEALTH_DIAG aux_b1_diag`) needs.
//
// Single-block tree-reduce kernel that scores the auxiliary next-bar
// regression head's per-bar prediction sign against the next-bar return
// label sign. Produces three batch-level fractions:
// out_3[0] = dir_acc — fraction of valid bars where pred_sign == label_sign
// out_3[1] = pos_pred_frac — fraction of valid bars where pred > 0
// out_3[2] = pos_label_frac — fraction of valid bars where label > 0
// 2-class softmax classifier head's per-bar argmax against the i32
// direction label. Produces six batch-level scalars:
// out_6[0] = dir_acc — fraction of valid bars where argmax(softmax) == label
// out_6[1] = pos_pred_frac — fraction of valid bars where argmax(softmax) == 1
// out_6[2] = pos_label_frac — fraction of valid bars where label == 1
// out_6[3] = n_down — count of valid bars where label == 0 (raw count)
// out_6[4] = n_up — count of valid bars where label == 1 (raw count)
// out_6[5] = n_skip — count of bars where label == -1 (raw count)
//
// "Valid" bars are those whose `next_bar_label != 0.0f` — a zero label is
// treated as "no signal to score against" (the producer of the label
// emits 0 when next-bar return rounds to flat) and excluded from the
// denominator entirely. When the entire batch is invalid (denom == 0),
// all three outputs are the random-baseline sentinel 0.5 (matches
// "Valid" bars are those whose `labels[i] != -1` — `-1` is the producer's
// mask sentinel ("no signal to score against"). Valid bars are split into
// down (label==0) and up (label==1) classes. When the entire batch is
// invalid (all labels == -1, i.e. denom == 0), `dir_acc / pos_pred_frac
// / pos_label_frac` are the random-baseline sentinel 0.5 (matches
// `DIR_ACC_EMA_SENTINEL` in `sp13_isv_slots.rs` per
// `pearl_first_observation_bootstrap` — the EMA's first observation
// replaces the sentinel directly).
// replaces the sentinel directly). The raw counts (n_down/n_up/n_skip)
// are returned as their actual integer values regardless — they're
// HEALTH_DIAG observability, not part of the EMA bootstrap.
//
// Sign convention for the prediction: `(p > 0.0f) ? 1 : 0`. Zero (and
// negative-zero) predict class 0 (negative). Tested by the
// `dir_acc_zero_pred_is_negative` GPU oracle.
// Argmax convention for the prediction: `(softmax[b, 1] > softmax[b, 0])
// ? 1 : 0`. Ties (both equal) predict class 0 (matches the regression-
// mode convention `(p > 0.0f) ? 1 : 0` where 0/-0 mapped to negative).
//
// Reads regression-mode aux scalar in P0a (the existing aux_heads
// next-bar regression head writes one f32 per bar). Layer B refactors the
// aux head to a 2-class softmax classifier and updates this kernel to
// read `(softmax[1] > softmax[0])` instead — the producer/consumer
// contract on the rest of the pipeline (3-element output, sentinel
// semantics, valid-bar masking) is unchanged.
//
// Single block, BLOCK_SIZE=256. Four shared-memory int arrays
// (correct, pos_pred, pos_label, valid) all participate in one fused
// log2(BLOCK_SIZE) tree reduce — no atomicAdd per
// Single block, BLOCK_SIZE=256. Six shared-memory int arrays (correct,
// pos_pred, pos_label, n_down, n_up, n_skip) all participate in one
// fused log2(BLOCK_SIZE) tree reduce — no atomicAdd per
// `feedback_no_atomicadd.md`. Pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Shared memory footprint:
// 4 × 256 × sizeof(int) = 4096 bytes.
// 6 × 256 × sizeof(int) = 6144 bytes (was 4 × 256 × sizeof(int) = 4096
// in the regression-mode version).
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
extern "C" __global__ void aux_dir_acc_reduce_kernel(
/* Per-bar aux-head prediction. P0a regression mode: one f32 per bar
* (the next-bar return regression scalar). Layer B switches the head
* to 2-class softmax and the launcher swaps the source pointer to
* `softmax[1] - softmax[0]` — same buffer shape, same kernel. */
const float* __restrict__ aux_pred,
/* Per-bar next-bar return label, sign-encoded. +1 / -1 for valid
* directional bars; 0 for "no signal" bars which are excluded from
* the denominator (see header comment). */
const float* __restrict__ next_bar_label,
/* Per-bar aux-head softmax tile [B, K] — produced inside the
* captured forward graph by `aux_next_bar_forward` (SP13 B1.1a),
* valid after the launcher's same-stream sync against the producer.
* Layout row-major; the kernel reads `softmax[b * K + 0]` and
* `softmax[b * K + 1]` per-bar to compute argmax. */
const float* __restrict__ aux_softmax,
/* Per-bar i32 direction label in {-1, 0, 1}. -1 = mask/skip;
* 0 = down (negative direction); 1 = up (positive direction). */
const int* __restrict__ labels,
/* Number of bars in the batch. The kernel strides each thread over
* `[0, batch_size)` so any batch_size up to ~2^31 is valid; in
* practice batches are O(1k-10k). */
int batch_size,
/* 3-element output buffer. Mapped-pinned in production (see the
/* Softmax K (= AUX_NEXT_BAR_K = 2 in B1.1a). Passed runtime so the
* kernel signature stays stable across head-dim flips; the only
* indices read are `b*K + 0` and `b*K + 1` so K must be ≥ 2. */
int K,
/* 6-element output buffer. Mapped-pinned in production (see the
* launcher in `gpu_dqn_trainer.rs`). The kernel writes
* out_3[0] = dir_acc
* out_3[1] = pos_pred_frac
* out_3[2] = pos_label_frac
* out_6[0] = dir_acc
* out_6[1] = pos_pred_frac
* out_6[2] = pos_label_frac
* out_6[3] = n_down
* out_6[4] = n_up
* out_6[5] = n_skip
* in thread 0 after the tree-reduce completes. */
float* __restrict__ out_3)
float* __restrict__ out_6)
{
/* Single-block reducer — guard against accidental multi-block launch. */
if (blockIdx.x != 0) return;
@@ -68,77 +78,104 @@ extern "C" __global__ void aux_dir_acc_reduce_kernel(
const int tid = threadIdx.x;
const int bdim = blockDim.x;
/* Four int arrays packed back-to-back in dynamic shared memory.
/* Six int arrays packed back-to-back in dynamic shared memory.
* `extern __shared__ int shared[]` is sized at launch via
* `shared_mem_bytes = 4 × bdim × sizeof(int)` — the launcher MUST
* `shared_mem_bytes = 6 × bdim × sizeof(int)` — the launcher MUST
* pass that exact value (the test scaffold and the production
* launcher both do). */
extern __shared__ int shared[];
int* sh_correct = &shared[0 * bdim];
int* sh_pos_pred = &shared[1 * bdim];
int* sh_pos_label = &shared[2 * bdim];
int* sh_valid = &shared[3 * bdim];
int* sh_n_down = &shared[3 * bdim];
int* sh_n_up = &shared[4 * bdim];
int* sh_n_skip = &shared[5 * bdim];
/* Per-thread strided accumulation. Each thread folds its slice of
* `[0, batch_size)` into four private int counters, then writes them
* into shared memory for the tree reduce. Bars whose label is zero
* (no directional signal to score against) are skipped entirely —
* `valid` does NOT increment, and `correct/pos_pred/pos_label` are
* not touched for those bars. */
* `[0, batch_size)` into six private int counters, then writes them
* into shared memory for the tree reduce. Bars whose label is -1
* (no directional signal to score against) increment `n_skip` only —
* `correct/pos_pred/pos_label/n_down/n_up` are not touched. */
int local_correct = 0;
int local_pos_pred = 0;
int local_pos_label = 0;
int local_valid = 0;
int local_n_down = 0;
int local_n_up = 0;
int local_n_skip = 0;
for (int i = tid; i < batch_size; i += bdim) {
const float p = aux_pred[i];
const float l = next_bar_label[i];
if (l == 0.0f) continue;
local_valid += 1;
const int pred_pos = (p > 0.0f) ? 1 : 0;
const int label_pos = (l > 0.0f) ? 1 : 0;
const int tgt = labels[i];
if (tgt == -1) {
local_n_skip += 1;
continue;
}
/* Defensive: producer (B1.1b) emits only -1/0/1, but guard
* against producer bugs by treating out-of-range labels as
* skip rather than reading garbage past the [B, K] softmax row. */
if (tgt < 0 || tgt >= K) {
local_n_skip += 1;
continue;
}
const float p0 = aux_softmax[(size_t)i * K + 0];
const float p1 = aux_softmax[(size_t)i * K + 1];
const int pred_pos = (p1 > p0) ? 1 : 0;
const int label_pos = (tgt == 1) ? 1 : 0;
if (pred_pos == label_pos) local_correct += 1;
local_pos_pred += pred_pos;
local_pos_label += label_pos;
if (tgt == 0) local_n_down += 1;
else local_n_up += 1;
}
sh_correct[tid] = local_correct;
sh_pos_pred[tid] = local_pos_pred;
sh_pos_label[tid] = local_pos_label;
sh_valid[tid] = local_valid;
sh_n_down[tid] = local_n_down;
sh_n_up[tid] = local_n_up;
sh_n_skip[tid] = local_n_skip;
__syncthreads();
/* Standard log2(BLOCK_SIZE) tree reduction over all four arrays in
/* Standard log2(BLOCK_SIZE) tree reduction over all six arrays in
* lockstep. One `__syncthreads()` per pass keeps every thread's view
* of all four arrays consistent — splitting them into separate
* loops would cost 4× the syncs without any compute saving. */
* of all six arrays consistent — splitting them into separate
* loops would cost 6× the syncs without any compute saving. */
for (int s = bdim / 2; s > 0; s >>= 1) {
if (tid < s) {
sh_correct[tid] += sh_correct[tid + s];
sh_pos_pred[tid] += sh_pos_pred[tid + s];
sh_pos_label[tid] += sh_pos_label[tid + s];
sh_valid[tid] += sh_valid[tid + s];
sh_n_down[tid] += sh_n_down[tid + s];
sh_n_up[tid] += sh_n_up[tid + s];
sh_n_skip[tid] += sh_n_skip[tid + s];
}
__syncthreads();
}
/* Thread 0 finalises. When the batch has no valid bars (empty batch
* or all-zero labels) the denom is 0; all three outputs fall back to
* 0.5 — the random-baseline sentinel matching DIR_ACC_EMA_SENTINEL
* in `sp13_isv_slots.rs`. The downstream EMA's first-observation
* replacement (Pearl A) consumes the sentinel naturally. */
* or all-mask labels) the denom is 0; dir_acc / pos_pred_frac /
* pos_label_frac fall back to 0.5 — the random-baseline sentinel
* matching DIR_ACC_EMA_SENTINEL in `sp13_isv_slots.rs`. The raw
* counts (n_down / n_up / n_skip) emit their actual values
* regardless (HEALTH_DIAG observability, not bootstrap-sensitive). */
if (tid == 0) {
const float denom = (float)sh_valid[0];
const float sentinel = 0.5f;
const int n_down = sh_n_down[0];
const int n_up = sh_n_up[0];
const int n_skip = sh_n_skip[0];
const int n_valid = n_down + n_up;
const float denom = (float)n_valid;
const float sentinel = 0.5f;
if (denom > 0.0f) {
out_3[0] = (float)sh_correct[0] / denom;
out_3[1] = (float)sh_pos_pred[0] / denom;
out_3[2] = (float)sh_pos_label[0] / denom;
out_6[0] = (float)sh_correct[0] / denom;
out_6[1] = (float)sh_pos_pred[0] / denom;
out_6[2] = (float)sh_pos_label[0] / denom;
} else {
out_3[0] = sentinel;
out_3[1] = sentinel;
out_3[2] = sentinel;
out_6[0] = sentinel;
out_6[1] = sentinel;
out_6[2] = sentinel;
}
out_6[3] = (float)n_down;
out_6[4] = (float)n_up;
out_6[5] = (float)n_skip;
__threadfence_system();
}
}

302
crates/ml/src/cuda_pipeline/aux_heads_kernel.cu
View File

@@ -1,12 +1,18 @@
/*
* aux_heads_kernel.cu — Plan 4 Task 6 Commit A.
* aux_heads_kernel.cu — Plan 4 Task 6 Commit A; SP13 B1.1a (2026-05-05)
* flipped the next-bar head from K=1 MSE regression to K=2 softmax CE
* classification. The remaining brief comments describing K_out = 1 / MSE
* are preserved BELOW where the rewrite history is relevant; the active
* design is documented at the top of each kernel.
*
* Multi-task auxiliary heads (E.6): two small MLPs branching off the
* trunk's `h_s2 [B, SH2]` post-GRN activation:
*
* (1) Next-bar return regression head:
* pred = Linear_2(ELU(Linear_1(h_s2))) where K_out = 1
* Loss = MSE against `next_close_pct` derived from `next_states`.
* (1) Next-bar direction classification head (SP13 B1.1a):
* logits = Linear_2(ELU(Linear_1(h_s2))) where K_out = 2
* softmax = stable_softmax(logits)
* Loss = cross-entropy against per-bar i32 direction label
* (0 = down, 1 = up, -1 = mask/skip). Producer wires in B1.1b.
*
* (2) Five-class regime classification head:
* logits = Linear_2(ELU(Linear_1(h_s2))) where K_out = 5
@@ -16,9 +22,16 @@
* Both heads share the `Linear(SH2 → 32) → ELU → Linear(32 → K)` shape
* with `H = AUX_HIDDEN_DIM = 32`.
*
* This commit lands the kernels + Rust orchestrator + params/ISV/registry
* scaffolding only — there are NO production callers in this commit.
* Commit B wires the forward/backward + training-loop loss accumulation.
* SP13 B1.1a (2026-05-05): the next-bar head's K_out flipped 1 → 2 and
* the loss formulation flipped MSE → softmax CE. The forward kernel now
* emits BOTH the [B, K] logits tile (saved for backward via mean-shift
* recompute is unnecessary — we save the softmax tile directly) AND a
* [B, K] softmax tile. Three downstream consumers read the softmax tile:
* the loss reduce, the backward kernel, and the dir-acc / isv-tanh
* producers. Labels are i32 in [0, K) with `-1` reserved for "mask/skip".
* Mean-over-batch CE divides by B_valid (count of non-masked rows), not
* B, so an all-skip batch produces loss = 0 and zero gradients without
* NaN. B1.1b lands the producer kernel that fills the labels.
*
* Pearls applied:
* - feedback_no_atomicadd.md — per-block shmem-tree reductions, no
@@ -74,36 +87,54 @@ __device__ __forceinline__ float aux_elu_bwd_from_post(float y) {
}
/* =====================================================================
* aux_next_bar_forward — Linear → ELU → Linear with K=1 output.
* aux_next_bar_forward — Linear → ELU → Linear with K-way output (SP13
* B1.1a: K_out = AUX_NEXT_BAR_K = 2, softmax classification).
*
* Per sample b (one block per sample, AUX_BLOCK threads):
* 1. h[k] = ELU( b1[k] + sum_j w1[k, j] * h_s2[b, j] ) for k ∈ [0, H)
* 2. pred[b] = b2 + sum_k w2[0, k] * h[k]
* 2. logits[b, kc] = b2[kc] + sum_k w2[kc, k] * h[k] for kc ∈ [0, K)
* 3. softmax[b, kc] = stable_softmax(logits[b, :])[kc] (max-shift)
*
* Hidden activations `h [B, H]` are SAVED for backward (consumed as the
* ELU post-activation buffer; backward derives `f'(x_pre)` from `h_post`
* via the standard ELU backward identity).
* via the standard ELU backward identity). The softmax tile `[B, K]` is
* also saved — backward re-reads it (rather than recomputing from logits)
* because the softmax tile has THREE downstream consumers in B1.1a:
*
* Block: AUX_BLOCK threads. Grid: (B, 1, 1). Shared memory: H floats
* (h cache for the second linear's reduction).
* - aux_next_bar_loss_reduce (CE numerator: -log(softmax[label]))
* - aux_next_bar_backward (d_logits = softmax - one_hot(label))
* - aux_dir_acc_reduce_kernel (argmax-vs-label compare)
* - aux_pred_to_isv_tanh_kernel (mean(softmax[1] - softmax[0]) → ISV[375])
*
* Mirroring `aux_regime_forward`'s K-fanout structure — thread 0 computes
* all K logits serially (K=2 / K=5; the matmul is small enough that a
* tree reduce wastes more cycles in shfl bookkeeping than the serial
* pass), then computes the stable softmax in-place on the same K-vector.
*
* Block: AUX_BLOCK threads. Grid: (B, 1, 1). Shared memory: H + K floats
* (h cache for Linear_2 + temporary K-vector for stable softmax).
* ===================================================================== */
extern "C" __global__ void aux_next_bar_forward(
const float* __restrict__ h_s2, /* [B, SH2] row-major */
const float* __restrict__ w1, /* [H, SH2] row-major */
const float* __restrict__ b1, /* [H] */
const float* __restrict__ w2, /* [1, H] (== [H] flat) */
const float* __restrict__ b2, /* [1] */
const float* __restrict__ w2, /* [K, H] row-major */
const float* __restrict__ b2, /* [K] */
int B,
int SH2,
int K,
float* __restrict__ hidden_out, /* [B, H] OUT (saved for backward) */
float* __restrict__ pred_out /* [B, 1] OUT */
float* __restrict__ logits_out, /* [B, K] OUT (saved for backward) */
float* __restrict__ softmax_out /* [B, K] OUT (saved for 3 consumers) */
) {
const int b = blockIdx.x;
if (b >= B) return;
const int tid = threadIdx.x;
const int H = AUX_HIDDEN_DIM;
extern __shared__ float smem[]; /* h cache [H] */
extern __shared__ float smem[]; /* [H + K] floats */
float* sh_h = smem; /* [H] hidden cache */
float* sh_logit = smem + H; /* [K] logits cache for softmax */
/* Step 1: per-hidden-unit Linear_1 + ELU. Each thread computes one or
* more hidden lanes via stride loop. SH2 is small enough that each
@@ -117,20 +148,43 @@ extern "C" __global__ void aux_next_bar_forward(
acc += w_row[j] * h_row[j];
}
const float h_post = aux_elu_fwd(acc);
smem[k] = h_post;
sh_h[k] = h_post;
hidden_out[(size_t)b * H + k] = h_post;
}
__syncthreads();
/* Step 2: scalar Linear_2. One thread does the H-element dot. H is
* small (32) — a tree reduce wastes more cycles in shfl bookkeeping
* than the serial pass. */
/* Step 2: K-way Linear_2. Thread 0 emits all K logits (K is small —
* 2 in B1.1a; mirrors the regime kernel's K=5 fanout shape). Stores
* each logit into the saved buffer AND a shmem scratch for Step 3's
* softmax. */
if (tid == 0) {
float acc = b2[0];
for (int k = 0; k < H; ++k) {
acc += w2[k] * smem[k];
for (int kc = 0; kc < K; ++kc) {
float acc = b2[kc];
const float* w2_row = w2 + (size_t)kc * H;
for (int k = 0; k < H; ++k) {
acc += w2_row[k] * sh_h[k];
}
logits_out[(size_t)b * K + kc] = acc;
sh_logit[kc] = acc;
}
}
__syncthreads();
/* Step 3: stable softmax on the K-vector — subtract max before exp.
* Thread 0 again (K=2; tree-reducing a K=2 vector is wasted bookkeeping). */
if (tid == 0) {
float lmax = sh_logit[0];
for (int kc = 1; kc < K; ++kc) {
if (sh_logit[kc] > lmax) lmax = sh_logit[kc];
}
float sum_e = 0.0f;
for (int kc = 0; kc < K; ++kc) {
sum_e += expf(sh_logit[kc] - lmax);
}
const float inv_sum = 1.0f / sum_e;
for (int kc = 0; kc < K; ++kc) {
softmax_out[(size_t)b * K + kc] = expf(sh_logit[kc] - lmax) * inv_sum;
}
pred_out[b] = acc;
}
}
@@ -192,48 +246,94 @@ extern "C" __global__ void aux_regime_forward(
}
/* =====================================================================
* aux_next_bar_loss_reduce — single-block shmem-reduce mean MSE.
* aux_next_bar_loss_reduce — single-block shmem-reduce mean cross-entropy
* over the K-class softmax tile (SP13 B1.1a; flipped from MSE).
*
* Computes loss = (1/B) * sum_b (pred[b] - label[b])^2 into loss_out[0].
* Single-block (AUX_BLOCK threads) shmem-tree reduction; no atomicAdd.
* For each sample b:
* if labels[b] == -1: contribute 0 to numerator AND 0 to valid count
* else : numerator += -log(softmax[b, labels[b]])
* valid_count += 1
*
* Block: AUX_BLOCK threads. Grid: (1, 1, 1). Shared mem: AUX_BLOCK floats.
* Reduces both the numerator and the valid count in lockstep via two
* parallel shmem-tree strips, then writes
* loss_out[0] = numerator / max(valid_count, 1.0f)
* valid_count_out[0] = valid_count (saved for backward)
*
* Mean-over-valid-bars semantics: an all-skip batch produces loss = 0
* (no NaN from `0 / max(0, 1)` since the divisor is masked to ≥1) and
* the saved valid_count = 0 makes the backward kernel emit zero
* gradients across the board.
*
* Single-block (AUX_BLOCK threads) shmem-tree reduction; no atomicAdd
* per `feedback_no_atomicadd.md`. Two parallel partial-reduction strips
* (loss numerator + valid count); shmem layout matches the K-class
* `aux_regime_loss_reduce` shape.
*
* The numerically-stable softmax was already materialised in the forward
* kernel — this reduce just reads `softmax[b, labels[b]]` directly,
* avoiding a second max-shift+exp pass per row.
*
* Block: AUX_BLOCK threads. Grid: (1, 1, 1).
* Shared memory: 2 * AUX_BLOCK floats (loss + valid-count partials).
* ===================================================================== */
extern "C" __global__ void aux_next_bar_loss_reduce(
const float* __restrict__ pred, /* [B] (== [B, 1] flat) */
const float* __restrict__ label, /* [B] */
const float* __restrict__ softmax, /* [B, K] row-major */
const int* __restrict__ labels, /* [B] i32 in {-1, 0, 1, …, K-1} */
int B,
float* __restrict__ loss_out /* [1] */
int K,
float* __restrict__ loss_out, /* [1] mean CE over valid rows */
float* __restrict__ valid_count_out /* [1] B_valid (read by backward) */
) {
if (blockIdx.x != 0) return;
extern __shared__ float smem[];
float* sh_loss = smem;
float* sh_valid = smem + blockDim.x;
const int tid = threadIdx.x;
const int block = blockDim.x;
/* 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;
/* Per-thread strided accumulation: each thread folds its slice of
* `[0, B)` into private (loss_numer, valid_count) accumulators.
* Mask-encoded rows (label == -1) contribute 0 to BOTH accumulators —
* no branch in the loss path; the divisor uses fmaxf(valid, 1) to
* prevent NaN when the entire batch is masked. */
float local_loss = 0.0f;
float local_valid = 0.0f;
for (int i = tid; i < B; i += block) {
const float d = pred[i] - label[i];
local += d * d;
const int tgt = labels[i];
if (tgt == -1) continue;
/* Defensive bounds: out-of-range labels degrade to skip rather
* than read past the end of the [B, K] softmax row. The producer
* (B1.1b) is responsible for emitting only -1 / [0, K) — this
* branch is dead under correct producer behavior, kept for kernel
* robustness against producer bugs. */
if (tgt < 0 || tgt >= K) continue;
const float p_tgt = softmax[(size_t)i * K + tgt];
/* Numerical floor: the forward softmax can produce p ≈ 0 for
* extreme negative logits; -log(0) = +inf would propagate to the
* EMA. Floor at FLT_MIN-ish (1e-30); equivalent to clipping the
* logit at log(1e-30) ≈ -69 which is well outside the trained
* regime. */
const float p_clipped = fmaxf(p_tgt, 1e-30f);
local_loss += -logf(p_clipped);
local_valid += 1.0f;
}
smem[tid] = local;
sh_loss[tid] = local_loss;
sh_valid[tid] = local_valid;
__syncthreads();
/* Standard log2(BLOCK) tree reduction over both arrays in lockstep. */
for (int s = block / 2; s > 0; s >>= 1) {
if (tid < s) smem[tid] += smem[tid + s];
if (tid < s) {
sh_loss[tid] += sh_loss[tid + s];
sh_valid[tid] += sh_valid[tid + s];
}
__syncthreads();
}
/* B > 0 invariant — caller guarantees positive batch. NaN propagation
* preferred over masking per the same rationale documented in
* h_s2_rms_ema_kernel.cu and vsn_mask_ema_kernel.cu. */
if (tid == 0) {
loss_out[0] = smem[0] / (float)B;
const float valid = sh_valid[0];
loss_out[0] = sh_loss[0] / fmaxf(valid, 1.0f);
valid_count_out[0] = valid;
}
}
@@ -382,18 +482,20 @@ extern "C" __global__ void aux_next_bar_backward(
/* Forward inputs (re-read for dW). */
const float* __restrict__ h_s2, /* [B, SH2] */
const float* __restrict__ w1, /* [H, SH2] */
const float* __restrict__ w2, /* [H] flat */
const float* __restrict__ w2, /* [K, H] row-major */
const float* __restrict__ hidden_post, /* [B, H] saved h_post */
/* Loss inputs. */
const float* __restrict__ pred, /* [B] */
const float* __restrict__ label, /* [B] */
/* Loss inputs (SP13 B1.1a softmax CE). */
const float* __restrict__ softmax, /* [B, K] saved softmax tile */
const int* __restrict__ labels, /* [B] i32 in {-1, 0..K-1} */
const float* __restrict__ valid_count, /* [1] B_valid from loss kernel */
int B,
int SH2,
int K,
/* Partial outputs (per-sample). Caller reduces along batch dim. */
float* __restrict__ dW1_partial, /* [B, H, SH2] flat */
float* __restrict__ db1_partial, /* [B, H] */
float* __restrict__ dW2_partial, /* [B, H] */
float* __restrict__ db2_partial, /* [B] */
float* __restrict__ dW2_partial, /* [B, K, H] flat */
float* __restrict__ db2_partial, /* [B, K] */
/* Trunk gradient (contiguous in B, no batch-reduce needed — caller
* accumulates into trunk's main-path d_h_s2 via SAXPY). */
float* __restrict__ dh_s2_out /* [B, SH2] */
@@ -404,40 +506,82 @@ extern "C" __global__ void aux_next_bar_backward(
const int H = AUX_HIDDEN_DIM;
extern __shared__ float smem[];
float* sh_dh_pre = smem; /* [H] */
float* sh_h_post = smem + H; /* [H] */
float* sh_dh_pre = smem; /* [H] */
float* sh_h_post = smem + H; /* [H] */
float* sh_dlogits = smem + 2 * H; /* [K] */
/* 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]);
/* SP13 B1.1a (2026-05-05): softmax CE backward. Mirrors
* `aux_regime_backward`'s K-fanout structure but reads the
* pre-computed softmax tile from forward (rather than re-softmaxing
* the logits) since the tile is already saved for three other
* consumers. Mean-over-valid-rows derivative:
* d_logits[b, kc] = (softmax[b, kc] - one_hot(labels[b], kc)) / B_valid
* for valid rows; masked rows (labels[b] == -1) zero out d_logits
* across the K-vector so the partial-grad write becomes a no-op.
* B_valid is read from the [1] scalar that the loss kernel wrote in
* the same step; floor at 1.0f to mirror the loss-kernel's
* fmaxf(valid, 1) divisor (an all-skip batch produces zero gradients
* across the board — d_logits = 0 for every row, no NaN). */
const int tgt = labels[b];
const float B_v = fmaxf(valid_count[0], 1.0f);
const float inv_B = 1.0f / B_v;
/* Step 1: cache h_post and compute d_h_pre per hidden unit. */
for (int k = tid; k < H; k += blockDim.x) {
const float h_post = hidden_post[(size_t)b * H + k];
sh_h_post[k] = h_post;
const float d_h_post = d_pred_b * w2[k];
sh_dh_pre[k] = d_h_post * aux_elu_bwd_from_post(h_post);
/* Step 0: compute d_logits[b, :] from softmax + label.
* Single thread — K is small (2 in B1.1a; mirrors regime kernel
* pattern at K=5). Masked rows zero the K-vector. */
if (tid == 0) {
if (tgt == -1 || tgt < 0 || tgt >= K) {
for (int kc = 0; kc < K; ++kc) {
sh_dlogits[kc] = 0.0f;
}
} else {
const float* p_row = softmax + (size_t)b * K;
for (int kc = 0; kc < K; ++kc) {
const float onehot = (kc == tgt) ? 1.0f : 0.0f;
sh_dlogits[kc] = inv_B * (p_row[kc] - onehot);
}
}
}
__syncthreads();
/* Step 2: per-sample partials for {db2, dW2, db1, dW1}.
* db2_partial[b] = d_pred_b
* dW2_partial[b, k] = d_pred_b * h_post[b, k]
* db1_partial[b, k] = sh_dh_pre[k]
* dW1_partial[b, k, j] = sh_dh_pre[k] * h_s2[b, j]
*/
/* Step 1: cache h_post and compute d_h_pre per hidden unit.
* d_h_post[b, k] = sum_kc d_logits[b, kc] * w2[kc, k]
* d_h_pre[b, k] = d_h_post[b, k] * elu_bwd(h_post[b, k]) */
for (int k = tid; k < H; k += blockDim.x) {
const float h_post = hidden_post[(size_t)b * H + k];
sh_h_post[k] = h_post;
float d_h_post = 0.0f;
for (int kc = 0; kc < K; ++kc) {
d_h_post += sh_dlogits[kc] * w2[(size_t)kc * H + k];
}
sh_dh_pre[k] = d_h_post * aux_elu_bwd_from_post(h_post);
}
__syncthreads();
/* Step 2: per-sample param-grad partials.
* db2_partial[b, kc] = sh_dlogits[kc]
* dW2_partial[b, kc, k] = sh_dlogits[kc] * h_post[b, k]
* db1_partial[b, k] = sh_dh_pre[k]
* dW1_partial[b, k, j] = sh_dh_pre[k] * h_s2[b, j] */
if (tid == 0) {
db2_partial[b] = d_pred_b;
for (int kc = 0; kc < K; ++kc) {
db2_partial[(size_t)b * K + kc] = sh_dlogits[kc];
}
}
/* dW2_partial[b, kc, k] = sh_dlogits[kc] * sh_h_post[k]. */
{
const int total = K * H;
for (int idx = tid; idx < total; idx += blockDim.x) {
const int kc = idx / H;
const int k = idx - kc * H;
dW2_partial[(size_t)b * K * H + idx] = sh_dlogits[kc] * sh_h_post[k];
}
}
for (int k = tid; k < H; k += blockDim.x) {
const float dh = sh_dh_pre[k];
dW2_partial[(size_t)b * H + k] = d_pred_b * sh_h_post[k];
db1_partial[(size_t)b * H + k] = dh;
db1_partial[(size_t)b * H + k] = sh_dh_pre[k];
}
/* dW1 partial — per-(k, j) write. Stride loop over (H * SH2). */
@@ -454,7 +598,9 @@ extern "C" __global__ void aux_next_bar_backward(
/* Step 3: dh_s2[b, j] = sum_k sh_dh_pre[k] * w1[k, j].
* Per-sample, per-feature serial reduction over H lanes — H is small
* (32) and SH2 is moderate; stride loop over j. */
* (32) and SH2 is moderate; stride loop over j. Masked rows propagate
* zero d_h_pre through this sum, so dh_s2 row is zero for skipped
* samples (the trunk SAXPY adds zero — no spurious gradient). */
{
float* dh_row = dh_s2_out + (size_t)b * SH2;
for (int j = tid; j < SH2; j += blockDim.x) {

87
crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu
View File

@@ -1,27 +1,37 @@
// crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu
//
// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction →
// ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar producer.
// ISV[AUX_DIR_PREDICTION_INDEX=375] bounded scalar producer.
// SP13 B1.1a (2026-05-05): rewrite — reads the K=2 softmax tile and
// computes `mean(softmax[:, 1] - softmax[:, 0])` ∈ [-1, +1] (structural
// bound, no `tanh()`). The kernel filename + ISV slot name are
// preserved across the rewrite — every consumer of slot 375 reads the
// same `[-1, +1]` bound it always did. The "tanh" in the filename is
// historical (regression-mode wrapped a raw scalar; B1.1a's softmax
// classifier emits structurally-bounded probabilities directly so no
// runtime squash is needed).
//
// Reads the aux next-bar regression head's `aux_pred [B]` tile (the
// per-bar return-prediction scalar produced inside the captured forward
// graph) and writes `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to the SHARED
// ISV scalar at slot 375. The tanh squash bounds the scalar so the
// downstream consumer (the direction Q-head input layer mirroring the
// SP13 spec's "feed aux prediction back into the policy" wiring) sees
// a well-conditioned signal rather than a raw regression scalar that
// could grow with `label_scale`.
// Architectural rationale (SP13 B1.1a per
// `pearl_bounded_modifier_outputs_require_structural_activation.md`):
// the previous regression-mode head emitted an unbounded scalar that
// needed a runtime `tanh()` to bound the ISV[375] consumer. The B1.1a
// head emits softmax probabilities directly: `softmax[1] - softmax[0]`
// is naturally bounded in `[-1, +1]` because the components are
// non-negative and sum to 1. The bound is now STRUCTURAL (a property
// of the model output's activation) rather than a runtime clamp on a
// raw scalar. Removing the tanh eliminates one nonlinearity from the
// ISV[375] producer chain.
//
// Why a batch mean instead of per-bar broadcast: ISV is a single
// `[ISV_TOTAL_DIM]` shared array (see `gpu_dqn_trainer.rs::isv_signals_pinned`
// allocation) — there is NO per-batch tile. Every per-step ISV producer
// reduces a per-bar tile to a single scalar before writing (cf.
// `h_s2_rms_ema_kernel.cu` reducing `save_h_s2 [B, SH2]` to a single
// `producer_step_scratch_buf` slot, or `aux_label_scale_ema_update`
// reducing `aux_nb_label_buf [B]` to one EMA slot). The mean of
// per-bar predictions is the natural batch-aggregate signal for a
// shared scalar — it sits at zero when the head has no consensus
// direction, drifts toward +1/-1 as the head develops conviction.
// `producer_step_scratch_buf` slot). The batch mean of the directional
// margin is the natural batch-aggregate signal for a shared scalar — it
// sits at zero when the head has no consensus direction (50/50 split),
// drifts toward +1 when the batch consensus is "up" (softmax[1]
// dominates), drifts toward -1 when the batch consensus is "down".
//
// Sentinel-friendly: per `pearl_first_observation_bootstrap`,
// `(prev == 0.0)` is the cold-start sentinel for slot 375 (constructor
@@ -31,36 +41,46 @@
// branch is needed here because the kernel ALWAYS overwrites the
// slot with the current step's batch mean — it's a per-step state,
// not an EMA, and the registry entry skips slot 375 explicitly per
// the SP13 P0a registry comments.
// the SP13 P0a registry comments. B1.1b will swap this slot's
// per-step write semantics to be aware of mask rows once the producer
// kernel for `aux_sign_labels` lands.
//
// Single block, BLOCK_SIZE=256, one shared-memory float array — no
// atomicAdd per `feedback_no_atomicadd.md`. Pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Stream-ordered with the
// producer that wrote `aux_pred`; same-stream barrier suffices
// producer that wrote the softmax tile; same-stream barrier suffices
// (cf. `aux_heads_loss_ema_update`'s same-stream contract with
// `aux_next_bar_loss_reduce`).
//
// Cost per launch: 256 threads × ceil(B/256) strided tanh + sum
// passes, plus log2(256)=8 reduce steps + 1 global write. B is
// O(1k-10k) in production training; total ~1-2µs per step, well
// under the per-step ISV-producer budget.
// Cost per launch: 256 threads × ceil(B/256) strided sum passes (two
// loads per row, no transcendentals — this is a bare additive reduce
// since the squash was retired), plus log2(256)=8 reduce steps + 1
// global write. B is O(1k-10k) in production training; total ~1-2µs
// per step, well under the per-step ISV-producer budget. Cheaper than
// the regression-mode kernel by exactly the cost of the tanh transcend.
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
extern "C" __global__ void aux_pred_to_isv_tanh_kernel(
/* Per-bar aux next-bar regression scalar [B] — produced inside
* the captured forward graph by `aux_next_bar_forward`, valid
* after the launcher's same-stream sync against the producer. */
const float* __restrict__ aux_pred,
/* Per-bar K=2 softmax tile [B, K] — produced inside the captured
* forward graph by `aux_next_bar_forward` (SP13 B1.1a). Valid after
* the launcher's same-stream sync against the producer. Layout
* row-major; the kernel reads `softmax[b * K + 0]` and
* `softmax[b * K + 1]` per-bar. */
const float* __restrict__ aux_softmax,
/* Number of bars in the batch. The strided loop covers any
* `batch_size > 0`; empty batch falls through to the sentinel
* write (see thread-0 finaliser). */
int batch_size,
/* Softmax K (= AUX_NEXT_BAR_K = 2 in B1.1a). Passed runtime so the
* kernel signature stays stable across head-dim flips; the only
* indices read are `b*K + 0` and `b*K + 1`. */
int K,
/* ISV slot index for the SHARED scalar — `AUX_DIR_PREDICTION_INDEX = 375`
* in the SP13 layout. The kernel writes
* isv[isv_slot_offset] = mean(tanh(aux_pred[i]))
* isv[isv_slot_offset] = mean(softmax[:, 1] - softmax[:, 0])
* after the tree-reduce, with `__threadfence_system()` to make the
* write visible to subsequent same-stream consumers. */
int isv_slot_offset,
@@ -75,15 +95,18 @@ extern "C" __global__ void aux_pred_to_isv_tanh_kernel(
const int tid = threadIdx.x;
const int bdim = blockDim.x;
/* Per-thread strided accumulation: each thread folds tanh(aux_pred[i])
* over its slice of `[0, batch_size)`. tanh squashes any unbounded
* regression scalar into [-1, +1] before the reduce so the batch
* mean is bounded by construction (the pre-squash mean could grow
* unbounded with `label_scale` and contaminate downstream
* consumers). */
/* Per-thread strided accumulation: each thread folds
* `(softmax[b, 1] - softmax[b, 0])` over its slice of `[0, batch_size)`.
* The pairwise difference is structurally bounded in [-1, +1] because
* each softmax component is in [0, 1] and they sum to 1; the batch
* mean is therefore also in [-1, +1] (the bound used to come from the
* tanh squash, but B1.1a's softmax classifier provides it
* structurally per `pearl_bounded_modifier_outputs_require_structural_activation`). */
float local_sum = 0.0f;
for (int i = tid; i < batch_size; i += bdim) {
local_sum += tanhf(aux_pred[i]);
const float p0 = aux_softmax[(size_t)i * K + 0];
const float p1 = aux_softmax[(size_t)i * K + 1];
local_sum += (p1 - p0);
}
sh_sum[tid] = local_sum;
__syncthreads();

107
crates/ml/src/cuda_pipeline/gpu_aux_heads.rs
View File

@@ -51,8 +51,14 @@ use super::gpu_dqn_trainer::{AUX_HEADS_CUBIN, AUX_HEADS_LOSS_EMA_CUBIN};
/// `aux_heads_kernel.cu`.
pub(crate) const AUX_HIDDEN_DIM: usize = 32;
/// Output cardinality of the next-bar regression head (scalar prediction).
pub(crate) const AUX_NEXT_BAR_K: usize = 1;
/// Output cardinality of the next-bar direction classification head
/// (SP13 B1.1a, 2026-05-05). Flipped 1 → 2: the head was a K=1 MSE
/// regression head; B1.1a flips it to a K=2 softmax classifier reading
/// i32 direction labels in {-1, 0, 1} where -1 = mask/skip. The
/// per-class softmax is structurally bounded in [0, 1] so downstream
/// ISV consumers (slot 375) see naturally-bounded inputs without a
/// runtime tanh squash per `pearl_bounded_modifier_outputs_require_structural_activation`.
pub(crate) const AUX_NEXT_BAR_K: usize = 2;
/// Output cardinality of the regime classification head (5-way softmax).
pub(crate) const AUX_REGIME_K: usize = 5;
@@ -125,17 +131,32 @@ impl AuxHeadsForwardOps {
})
}
/// Launch `aux_next_bar_forward`: `Linear(h_s2) → ELU → Linear → pred`.
/// Launch `aux_next_bar_forward`: `Linear(h_s2) → ELU → Linear → softmax`.
///
/// SP13 B1.1a (2026-05-05): K-flipped 1 → 2 with softmax tile output.
/// Forward writes BOTH the `[B, K]` logits buffer (saved for backward
/// re-read; the previous regression-mode design only saved the scalar
/// prediction) AND the `[B, K]` softmax tile (the new structurally-
/// bounded prediction surface read by 3 downstream consumers: loss
/// reduce, backward, dir-acc / isv-tanh producers).
///
/// Caller-owned buffers (raw `u64` device pointers for graph-capture
/// safety; mirrors `gpu_grn.rs::forward_raw`):
/// * `h_s2_ptr` — `[B, SH2]` row-major.
/// * `w1_ptr`/`b1_ptr` — `[H, SH2]` / `[H]`. Slice from `params_buf[119]`/`[120]`.
/// * `w2_ptr`/`b2_ptr` — `[1, H]` / `[1]`. Slice from `params_buf[121]`/`[122]`.
/// * `w2_ptr`/`b2_ptr` — `[K, H]` / `[K]`. Slice from `params_buf[121]`/`[122]`.
/// * `hidden_out_ptr` — `[B, H]` SAVED post-ELU buffer (consumed by backward).
/// * `pred_out_ptr` — `[B, 1]` scalar prediction per sample.
/// * `logits_out_ptr` — `[B, K]` SAVED logits (consumed by backward
/// only; the consumers below read the softmax tile).
/// * `softmax_out_ptr` — `[B, K]` SAVED softmax tile (consumed by
/// loss reduce + backward + dir-acc + isv-tanh).
///
/// Block: `AUX_BLOCK` threads. Grid: `B` blocks. Shared mem: `H` floats.
/// `k` is `AUX_NEXT_BAR_K = 2` in B1.1a; passed as runtime arg for
/// kernel-signature stability (mirrors `forward_regime`'s K-runtime
/// pattern — a future K bump won't require an ABI change).
///
/// Block: `AUX_BLOCK` threads. Grid: `B` blocks.
/// Shared mem: `(H + K) * sizeof(f32)` bytes (h cache + softmax temp).
#[allow(clippy::too_many_arguments)]
pub(crate) fn forward_next_bar(
&self,
@@ -147,12 +168,15 @@ impl AuxHeadsForwardOps {
b2_ptr: u64,
b: usize,
sh2: usize,
k: usize,
hidden_out_ptr: u64,
pred_out_ptr: u64,
logits_out_ptr: u64,
softmax_out_ptr: u64,
) -> Result<(), MLError> {
let b_i32 = b as i32;
let sh2_i32 = sh2 as i32;
let smem_bytes = (AUX_HIDDEN_DIM as u32) * std::mem::size_of::<f32>() as u32;
let k_i32 = k as i32;
let smem_bytes = (AUX_HIDDEN_DIM as u32 + k as u32) * std::mem::size_of::<f32>() as u32;
unsafe {
stream
.launch_builder(&self.next_bar_forward_kernel)
@@ -163,8 +187,10 @@ impl AuxHeadsForwardOps {
.arg(&b2_ptr)
.arg(&b_i32)
.arg(&sh2_i32)
.arg(&k_i32)
.arg(&hidden_out_ptr)
.arg(&pred_out_ptr)
.arg(&logits_out_ptr)
.arg(&softmax_out_ptr)
.launch(LaunchConfig {
grid_dim: (b as u32, 1, 1),
block_dim: (AUX_BLOCK, 1, 1),
@@ -221,29 +247,43 @@ impl AuxHeadsForwardOps {
Ok(())
}
/// Launch `aux_next_bar_loss_reduce`: scalar mean MSE into `loss_out[1]`.
/// Launch `aux_next_bar_loss_reduce`: scalar mean cross-entropy +
/// scalar valid-count.
///
/// 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.
/// SP13 B1.1a (2026-05-05): MSE → softmax CE flip. Reads the `[B, K]`
/// softmax tile (produced by `forward_next_bar`) and the `[B] i32`
/// labels buffer (`-1 / 0..K`; -1 is the producer's mask sentinel).
/// Writes mean CE = `numerator / max(B_valid, 1)` to `loss_out[1]`
/// AND the `B_valid` count to `valid_count_out[1]` for the backward
/// kernel to consume (so loss + backward share the same divisor —
/// derivatives of the same scalar function).
///
/// Two parallel partial-reduction strips (loss numerator + valid
/// count); shmem matches the K-class `regime_loss_reduce` shape.
pub(crate) fn next_bar_loss_reduce(
&self,
stream: &Arc<CudaStream>,
pred_ptr: u64,
label_ptr: u64,
softmax_ptr: u64,
labels_ptr: u64,
b: usize,
k: usize,
loss_out_ptr: u64,
valid_count_out_ptr: u64,
) -> Result<(), MLError> {
let b_i32 = b as i32;
let smem_bytes = AUX_BLOCK * std::mem::size_of::<f32>() as u32;
let k_i32 = k as i32;
// Two parallel partial-reduction strips: loss + valid-count. See
// `aux_next_bar_loss_reduce` shmem layout in the kernel.
let smem_bytes = 2 * AUX_BLOCK * std::mem::size_of::<f32>() as u32;
unsafe {
stream
.launch_builder(&self.next_bar_loss_reduce_kernel)
.arg(&pred_ptr)
.arg(&label_ptr)
.arg(&softmax_ptr)
.arg(&labels_ptr)
.arg(&b_i32)
.arg(&k_i32)
.arg(&loss_out_ptr)
.arg(&valid_count_out_ptr)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (AUX_BLOCK, 1, 1),
@@ -417,10 +457,14 @@ impl AuxHeadsBackwardOps {
/// `param_grad_reduce` for each of {`dW1`, `db1`, `dW2`, `db2`} to
/// collapse along the batch dim.
///
/// 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.
/// SP13 B1.1a (2026-05-05): MSE → softmax CE flip. Reads the saved
/// `[B, K]` softmax tile (produced by `forward_next_bar`) plus the
/// `[B] i32` labels buffer and the `[1]` `valid_count` scalar
/// (written by `next_bar_loss_reduce`). Mirrors `aux_regime_backward`'s
/// K-fanout structure but with mean-over-valid-rows derivative
/// `d_logits[b, kc] = (softmax[b, kc] - one_hot(labels[b], kc)) / B_valid`
/// for valid rows, zeroed across the K-vector for masked rows
/// (label == -1) so partial-grad writes are no-ops.
#[allow(clippy::too_many_arguments)]
pub(crate) fn backward_next_bar(
&self,
@@ -429,10 +473,12 @@ impl AuxHeadsBackwardOps {
w1_ptr: u64,
w2_ptr: u64,
hidden_post_ptr: u64,
pred_ptr: u64,
label_ptr: u64,
softmax_ptr: u64,
labels_ptr: u64,
valid_count_ptr: u64,
b: usize,
sh2: usize,
k: usize,
dw1_partial_ptr: u64,
db1_partial_ptr: u64,
dw2_partial_ptr: u64,
@@ -441,7 +487,10 @@ impl AuxHeadsBackwardOps {
) -> Result<(), MLError> {
let b_i32 = b as i32;
let sh2_i32 = sh2 as i32;
let smem_bytes = (2 * AUX_HIDDEN_DIM as u32) * std::mem::size_of::<f32>() as u32;
let k_i32 = k as i32;
// Shared mem: 2*H (d_h_pre + h_post caches) + K (d_logits cache).
let smem_bytes =
((2 * AUX_HIDDEN_DIM as u32) + (k as u32)) * std::mem::size_of::<f32>() as u32;
unsafe {
stream
.launch_builder(&self.next_bar_backward_kernel)
@@ -449,10 +498,12 @@ impl AuxHeadsBackwardOps {
.arg(&w1_ptr)
.arg(&w2_ptr)
.arg(&hidden_post_ptr)
.arg(&pred_ptr)
.arg(&label_ptr)
.arg(&softmax_ptr)
.arg(&labels_ptr)
.arg(&valid_count_ptr)
.arg(&b_i32)
.arg(&sh2_i32)
.arg(&k_i32)
.arg(&dw1_partial_ptr)
.arg(&db1_partial_ptr)
.arg(&dw2_partial_ptr)

437
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -641,12 +641,15 @@ pub(crate) static SP13_APPLY_FIXED_ALPHA_EMA_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/apply_fixed_alpha_ema_kernel.cubin"));
/// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction →
/// ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar producer.
/// Single-block tree-reduce reads the captured-graph `aux_nb_pred_buf [B]`
/// tile and writes `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to the SHARED
/// ISV slot 375 (per-step overwrite — not an EMA). tanh squash bounds
/// the scalar so the downstream Q-head consumer sees a well-conditioned
/// signal regardless of `label_scale` magnitude. Consumed by
/// ISV[AUX_DIR_PREDICTION_INDEX=375] bounded scalar producer.
/// SP13 B1.1a (2026-05-05) rewrite: reads the K=2 softmax tile
/// (`aux_nb_softmax_buf [B, 2]`) and writes
/// `mean(softmax[:, 1] - softmax[:, 0])` ∈ [-1, +1] to the SHARED
/// ISV slot 375 (per-step overwrite — not an EMA). The bound is
/// structural (softmax components in [0, 1] sum to 1 ⇒ pairwise diff
/// in [-1, +1]) per `pearl_bounded_modifier_outputs_require_structural_activation`;
/// the runtime tanh squash retired. Cubin filename and ISV slot name
/// preserved for layout stability. Consumed by
/// `GpuDqnTrainer::launch_aux_pred_to_isv_tanh`. See
/// `aux_pred_to_isv_tanh_kernel.cu` for kernel contract details.
static SP13_AUX_PRED_TO_ISV_TANH_CUBIN: &[u8] =
@@ -2141,7 +2144,7 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
PARAM_VSN_W1_G3=107;PARAM_VSN_B1_G3=108;PARAM_VSN_W2_G3=109;PARAM_VSN_B2_G3=110;\
PARAM_VSN_W1_G4=111;PARAM_VSN_B1_G4=112;PARAM_VSN_W2_G4=113;PARAM_VSN_B2_G4=114;\
PARAM_VSN_W1_G5=115;PARAM_VSN_B1_G5=116;PARAM_VSN_W2_G5=117;PARAM_VSN_B2_G5=118;\
PARAM_AUX_NB_W1=119;PARAM_AUX_NB_B1=120;PARAM_AUX_NB_W2=121;PARAM_AUX_NB_B2=122;\
PARAM_AUX_NB_W1=119;PARAM_AUX_NB_B1=120;PARAM_AUX_NB_W2_K2=121;PARAM_AUX_NB_B2_K2=122;\
PARAM_AUX_RG_W1=123;PARAM_AUX_RG_B1=124;PARAM_AUX_RG_W2=125;PARAM_AUX_RG_B2=126;\
PARAM_MOE_GATE_W1=127;PARAM_MOE_GATE_B1=128;PARAM_MOE_GATE_W2=129;PARAM_MOE_GATE_B2=130;\
PARAM_MOE_EXPERT_0_W1=131;PARAM_MOE_EXPERT_0_B1=132;PARAM_MOE_EXPERT_0_W2=133;PARAM_MOE_EXPERT_0_B2=134;\
@@ -2893,16 +2896,17 @@ pub(crate) fn compute_param_sizes(cfg: &GpuDqnTrainConfig) -> [usize; NUM_WEIGHT
// Two heads: next-bar regression (K=1) + regime classification (K=5).
// Each head is `Linear(SH2 → AUX_HIDDEN_DIM=32) → ELU → Linear(32 → K)`.
// Per-head total = 32*SH2 + 32 + K*32 + K. Sum over both heads:
// 2*32*SH2 + 2*32 + (1+5)*32 + (1+5) = 2*32*SH2 + 64 + 192 + 6
// = 2*32*SH2 + 262 floats.
// No production callers in this commit — params consumed by Commit B
// (forward orchestrator + training-loop loss accumulation + backward
// dispatcher).
// 2*32*SH2 + 2*32 + (2+5)*32 + (2+5) = 2*32*SH2 + 64 + 224 + 7
// = 2*32*SH2 + 295 floats.
// SP13 B1.1a (2026-05-05): next-bar head K_out flipped 1 → 2;
// [121] grew from `H` (= 32) to `K_NB * H` (= 64), [122] grew from `1`
// to `K_NB` (= 2). Mirrors the regime row's expression for K=5.
let sh2 = cfg.shared_h2;
sizes[119] = crate::cuda_pipeline::gpu_aux_heads::AUX_HIDDEN_DIM * sh2; // [119] aux_nb_w1 [32, SH2]
sizes[120] = crate::cuda_pipeline::gpu_aux_heads::AUX_HIDDEN_DIM; // [120] aux_nb_b1 [32]
sizes[121] = crate::cuda_pipeline::gpu_aux_heads::AUX_HIDDEN_DIM; // [121] aux_nb_w2 [1, 32]
sizes[122] = 1; // [122] aux_nb_b2 [1]
sizes[121] = crate::cuda_pipeline::gpu_aux_heads::AUX_NEXT_BAR_K
* crate::cuda_pipeline::gpu_aux_heads::AUX_HIDDEN_DIM; // [121] aux_nb_w2 [K_NB=2, 32]
sizes[122] = crate::cuda_pipeline::gpu_aux_heads::AUX_NEXT_BAR_K; // [122] aux_nb_b2 [K_NB=2]
sizes[123] = crate::cuda_pipeline::gpu_aux_heads::AUX_HIDDEN_DIM * sh2; // [123] aux_rg_w1 [32, SH2]
sizes[124] = crate::cuda_pipeline::gpu_aux_heads::AUX_HIDDEN_DIM; // [124] aux_rg_b1 [32]
sizes[125] = crate::cuda_pipeline::gpu_aux_heads::AUX_REGIME_K
@@ -4309,19 +4313,41 @@ pub struct GpuDqnTrainer {
/// `[B, AUX_HIDDEN_DIM]` saved post-ELU activation of the next-bar head's
/// Linear_1 (consumed by `aux_next_bar_backward` for the dW1 / dh_s2 chain).
aux_nb_hidden_buf: CudaSlice<f32>,
/// `[B, AUX_NEXT_BAR_K=1]` scalar prediction per sample. Read by
/// `aux_next_bar_loss_reduce` (vs `next_states[:, 0]` label) AND by
/// `aux_next_bar_backward` (for the `(2/B)*(pred-label)` derivative).
aux_nb_pred_buf: CudaSlice<f32>,
/// `[B]` per-sample next-bar return label, gathered from
/// `next_states_buf[:, 0]` via `strided_gather` at the start of every
/// forward call. DEDICATED buffer (no aliasing — the previous WIP
/// over-aliased `aux_partial_nb_b2` here and triggered a regression).
/// Consumed by both `aux_next_bar_loss_reduce` and `aux_next_bar_backward`.
aux_nb_label_buf: CudaSlice<f32>,
/// `[1]` scalar mean MSE — written by `aux_next_bar_loss_reduce`, read by
/// `aux_heads_loss_ema_update` (the ISV producer).
/// `[B, AUX_NEXT_BAR_K=2]` saved logits — written by
/// `aux_next_bar_forward` (B1.1a). Saved primarily for diagnostic
/// readback / future consumers; backward reads the softmax tile
/// instead of re-softmaxing the logits since the softmax is already
/// materialised. Field renamed from `aux_nb_pred_buf` (regression
/// scalar) to `aux_nb_logits_buf` (B1.1a logits) per
/// `feedback_no_legacy_aliases.md`.
aux_nb_logits_buf: CudaSlice<f32>,
/// `[B, AUX_NEXT_BAR_K=2]` saved softmax tile — written by
/// `aux_next_bar_forward` (B1.1a). Three downstream consumers read
/// this tile:
/// - `aux_next_bar_loss_reduce` (CE numerator)
/// - `aux_next_bar_backward` (`d_logits = softmax - one_hot(label)`)
/// - `aux_dir_acc_reduce_kernel` (argmax-vs-label compare)
/// - `aux_pred_to_isv_tanh_kernel` (mean(softmax[1] - softmax[0]) → ISV[375])
aux_nb_softmax_buf: CudaSlice<f32>,
/// `[B] i32` per-sample direction label in `{-1, 0, 1}` (B1.1a):
/// `-1` = mask/skip, `0` = down, `1` = up. **B1.1a is producer-less**
/// — the buffer is `alloc_zeros`-initialised (every sample receives
/// label 0 → "down") and stays at zero until B1.1b lands the
/// producer kernel that computes labels from the price trajectory.
/// "Model learns to predict class 0 (down) everywhere" is the known
/// degraded training behavior between B1.1a and B1.1b.
/// Consumed by `aux_next_bar_loss_reduce`, `aux_next_bar_backward`,
/// AND `aux_dir_acc_reduce_kernel`.
aux_nb_label_buf: CudaSlice<i32>,
/// `[1]` scalar mean cross-entropy — written by
/// `aux_next_bar_loss_reduce`, read by `aux_heads_loss_ema_update`
/// (the ISV producer).
aux_nb_loss_scalar_buf: CudaSlice<f32>,
/// `[1]` scalar `B_valid` count (number of non-mask labels in the
/// batch) — written by `aux_next_bar_loss_reduce`, read by
/// `aux_next_bar_backward` so loss + backward share the same
/// `1/B_valid` divisor (derivatives of the same scalar function).
aux_nb_valid_count_buf: CudaSlice<f32>,
/// `[B, AUX_HIDDEN_DIM]` saved post-ELU activation of the regime head's
/// Linear_1 (consumed by `aux_regime_backward`).
aux_rg_hidden_buf: CudaSlice<f32>,
@@ -4353,11 +4379,14 @@ pub struct GpuDqnTrainer {
aux_partial_nb_w1: CudaSlice<f32>,
/// `[B, AUX_HIDDEN_DIM]` per-sample db1 partial — next-bar head.
aux_partial_nb_b1: CudaSlice<f32>,
/// `[B, AUX_HIDDEN_DIM]` per-sample dW2 partial — next-bar head (W2 is
/// `[1, H]` flat-stored).
/// `[B, AUX_NEXT_BAR_K, AUX_HIDDEN_DIM]` per-sample dW2 partial —
/// next-bar head. SP13 B1.1a: K_NB flipped 1 → 2, so the partial
/// alloc grew from `[B, H]` to `[B, K, H]` (mirrors regime head's
/// `[B, K, H]` shape with K=5).
aux_partial_nb_w2: CudaSlice<f32>,
/// `[B]` per-sample db2 partial — next-bar head (b2 is scalar).
/// DEDICATED — does NOT alias `aux_nb_label_buf`.
/// `[B, AUX_NEXT_BAR_K]` per-sample db2 partial — next-bar head.
/// SP13 B1.1a: K_NB flipped 1 → 2, so the partial alloc grew from
/// `[B]` to `[B, K]`. DEDICATED — does NOT alias `aux_nb_label_buf`.
aux_partial_nb_b2: CudaSlice<f32>,
/// `[B, AUX_HIDDEN_DIM, SH2]` per-sample dW1 partial — regime head.
aux_partial_rg_w1: CudaSlice<f32>,
@@ -4985,26 +5014,28 @@ pub struct GpuDqnTrainer {
pub(crate) sp11_popart_component_len: usize,
/// SP13 Phase 0a (2026-05-04): aux-head directional-accuracy
/// reducer kernel. Single-block tree-reduce loaded from
/// `aux_dir_acc_reduce_kernel.cubin`. Reads the aux next-bar
/// regression-head prediction buffer + a per-bar sign-encoded
/// `next_bar_label` buffer and writes
/// `(dir_acc, pos_pred_frac, pos_label_frac)` to
/// `aux_dir_acc_buf` (mapped-pinned, 3 floats). Producer is wired
/// by P0a.T4; the kernel handle here lands the constructor edge
/// per `feedback_no_partial_refactor.md` (P0a is one atomic
/// commit). See the kernel source for the per-bar prediction sign
/// convention (`(p > 0.0f) ? 1 : 0`) and the all-zero-label
/// sentinel (0.5, matching `DIR_ACC_EMA_SENTINEL`).
/// `aux_dir_acc_reduce_kernel.cubin`. SP13 B1.1a (2026-05-05) ABI
/// flip: reads the aux next-bar K=2 softmax tile + i32 labels
/// (replacing the regression-mode `(aux_pred f32, label_sign f32)`
/// pair) and writes
/// `(dir_acc, pos_pred_frac, pos_label_frac, n_down, n_up, n_skip)`
/// to `aux_dir_acc_buf` (mapped-pinned, 6 floats — was 3). Producer
/// is wired by P0a.T4; the kernel handle here lands the constructor
/// edge per `feedback_no_partial_refactor.md` (P0a is one atomic
/// commit). See the kernel source for the argmax prediction
/// convention (`softmax[1] > softmax[0] ? 1 : 0`) and the
/// all-mask-label sentinel (0.5, matching `DIR_ACC_EMA_SENTINEL`).
aux_dir_acc_reduce: CudaFunction,
/// SP13 Phase 0a (2026-05-04): mapped-pinned 3-element output
/// buffer for `aux_dir_acc_reduce`. Layout `[dir_acc,
/// pos_pred_frac, pos_label_frac]` per the kernel contract. Per
/// `feedback_no_htod_htoh_only_mapped_pinned.md`: every CPU↔GPU
/// scalar read-back goes through `MappedF32Buffer` (host writes
/// the launcher reads via `read_all()`'s volatile read after stream
/// sync). Constructor-zero-initialised; the kernel overwrites all
/// three slots on every launch (or all three to the sentinel 0.5
/// when the batch is empty / all-zero-label).
/// SP13 Phase 0a (2026-05-04): mapped-pinned 6-element output
/// buffer for `aux_dir_acc_reduce` (SP13 B1.1a grew 3 → 6 to
/// expose per-class counts for HEALTH_DIAG). Layout
/// `[dir_acc, pos_pred_frac, pos_label_frac, n_down, n_up, n_skip]`
/// per the kernel contract. Per `feedback_no_htod_htoh_only_mapped_pinned.md`:
/// every CPU↔GPU scalar read-back goes through `MappedF32Buffer`
/// (host writes the launcher reads via `read_all()`'s volatile
/// read after stream sync). Constructor-zero-initialised; the
/// kernel overwrites all six slots on every launch (or first three
/// to the sentinel 0.5 when the batch is empty / all-mask-label).
pub(crate) aux_dir_acc_buf: super::mapped_pinned::MappedF32Buffer,
/// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-α EMA applicator
/// kernel handle. Loaded from `apply_fixed_alpha_ema_kernel.cubin`.
@@ -5017,12 +5048,13 @@ pub struct GpuDqnTrainer {
/// defeating the stagnation detector that compares slot 373 vs 374.
apply_fixed_alpha_ema_kernel: CudaFunction,
/// SP13 Phase 0a P0a.T4 (2026-05-04): aux-head per-bar prediction
/// → ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-bounded scalar producer
/// → ISV[AUX_DIR_PREDICTION_INDEX=375] bounded scalar producer
/// kernel handle. Loaded from `aux_pred_to_isv_tanh_kernel.cubin`.
/// Single-block tree-reduce reads `aux_nb_pred_buf [B]` and writes
/// `mean(tanh(aux_pred[i]))` ∈ [-1, +1] to ISV[375]. Per-step state
/// (not an EMA — slot is overwritten each launch); no FoldReset
/// registry entry. Consumed by `launch_aux_pred_to_isv_tanh`.
/// SP13 B1.1a (2026-05-05) rewrite: reads `aux_nb_softmax_buf [B, K=2]`
/// and writes `mean(softmax[:, 1] - softmax[:, 0])` ∈ [-1, +1] to
/// ISV[375]. Per-step state (not an EMA — slot is overwritten each
/// launch); no FoldReset registry entry. Bound is structural (no
/// runtime tanh). Consumed by `launch_aux_pred_to_isv_tanh`.
aux_pred_to_isv_tanh_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller
/// kernel. Single-block, 10-thread producer reading 5 canary ISV slots
@@ -13148,50 +13180,55 @@ impl GpuDqnTrainer {
/// SP13 Phase 0a (2026-05-04): launch the aux-head directional-
/// accuracy reducer.
///
/// Reads `aux_pred [B]` (the per-bar aux next-bar regression
/// prediction tile, P0a regression mode) and `next_bar_label [B]`
/// (per-bar sign-encoded label: +1 / 0 / -1) and writes
/// `(dir_acc, pos_pred_frac, pos_label_frac)` into the 3-element
/// output buffer at `out_3_dev`. Both predictions and labels live
/// on the GPU; the 3-element output is the trainer's mapped-pinned
/// `aux_dir_acc_buf` in production (the launcher is invoked with
/// `aux_dir_acc_buf.dev_ptr` from `training_loop.rs` after the
/// captured forward graph populates `aux_nb_pred_buf` /
/// `aux_nb_label_buf`). The launcher takes raw `u64` device
/// pointers per Decision D / `feedback_no_htod_htoh_only_mapped_pinned.md`
/// because the production source for `out_3_dev` is a
/// `MappedF32Buffer` (which exposes only `dev_ptr: u64`).
/// SP13 B1.1a (2026-05-05): reads `aux_softmax [B, K]` (the K=2
/// softmax tile produced inside the captured forward graph by
/// `aux_next_bar_forward`) and `labels [B] i32` (`-1 / 0 / 1` —
/// `-1` is producer mask sentinel; producer wires in B1.1b) and
/// writes
/// `(dir_acc, pos_pred_frac, pos_label_frac, n_down, n_up, n_skip)`
/// into the 6-element output buffer at `out_6_dev`. Both softmax
/// and labels live on the GPU; the 6-element output is the
/// trainer's mapped-pinned `aux_dir_acc_buf` in production (the
/// launcher is invoked with `aux_dir_acc_buf.dev_ptr` from
/// `training_loop.rs` after the captured forward graph populates
/// `aux_nb_softmax_buf` / `aux_nb_label_buf`). The launcher takes
/// raw `u64` device pointers per Decision D /
/// `feedback_no_htod_htoh_only_mapped_pinned.md` because the
/// production source for `out_6_dev` is a `MappedF32Buffer` (which
/// exposes only `dev_ptr: u64`).
///
/// Single-block tree-reduce (BLOCK_SIZE=256) — four shared-memory
/// int arrays (correct/pos_pred/pos_label/valid) reduce in lockstep
/// per `feedback_no_atomicadd.md`. The launcher computes
/// `shared_mem_bytes = 4 × bdim × sizeof(i32)` to match the kernel's
/// `extern __shared__ int shared[]` declaration; passing a smaller
/// value would make the kernel read past the dynamic shared-memory
/// region and corrupt other data.
/// Single-block tree-reduce (BLOCK_SIZE=256) — six shared-memory
/// int arrays (correct/pos_pred/pos_label/n_down/n_up/n_skip)
/// reduce in lockstep per `feedback_no_atomicadd.md`. The launcher
/// computes `shared_mem_bytes = 6 × bdim × sizeof(i32)` to match
/// the kernel's `extern __shared__ int shared[]` declaration; passing
/// a smaller value would make the kernel read past the dynamic
/// shared-memory region and corrupt other data.
///
/// Per `feedback_cudarc_f64_f32_abi.md`: `batch_size` is `i32`
/// (matches the kernel signature exactly — no implicit cast).
/// Per `feedback_cudarc_f64_f32_abi.md`: `batch_size` and `k` are
/// `i32` (match the kernel signature exactly — no implicit cast).
pub(crate) fn launch_aux_dir_acc_reduce(
&self,
aux_pred_dev: u64,
next_bar_label_dev: u64,
aux_softmax_dev: u64,
labels_dev: u64,
batch_size: i32,
out_3_dev: u64,
k: i32,
out_6_dev: u64,
) -> Result<(), MLError> {
let bdim: u32 = 256;
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: 4 * bdim * std::mem::size_of::<i32>() as u32,
shared_mem_bytes: 6 * bdim * std::mem::size_of::<i32>() as u32,
};
unsafe {
self.stream
.launch_builder(&self.aux_dir_acc_reduce)
.arg(&aux_pred_dev)
.arg(&next_bar_label_dev)
.arg(&aux_softmax_dev)
.arg(&labels_dev)
.arg(&batch_size)
.arg(&out_3_dev)
.arg(&k)
.arg(&out_6_dev)
.launch(cfg)
.map_err(|e| MLError::ModelError(format!("sp13 aux_dir_acc_reduce launch: {e}")))?;
}
@@ -13265,26 +13302,31 @@ impl GpuDqnTrainer {
}
/// SP13 Phase 0a P0a.T4 (2026-05-04): launch the aux-head
/// per-bar prediction → ISV[AUX_DIR_PREDICTION_INDEX=375] tanh-
/// bounded scalar producer.
/// per-bar prediction → ISV[AUX_DIR_PREDICTION_INDEX=375] bounded
/// scalar producer.
///
/// Reads `aux_pred [B]` (production source: `aux_nb_pred_buf`,
/// populated by the captured forward graph's
/// `aux_next_bar_forward`) and writes `mean(tanh(aux_pred[i]))`
/// ∈ [-1, +1] to the SHARED ISV slot at `isv_slot_offset`. Per-step
/// state, not an EMA — slot is overwritten each launch (no
/// FoldReset registry entry, per the SP13 P0a registry comments).
/// SP13 B1.1a (2026-05-05): reads `aux_softmax [B, K]` (production
/// source: `aux_nb_softmax_buf`, populated by the captured forward
/// graph's `aux_next_bar_forward`) and writes
/// `mean(softmax[:, 1] - softmax[:, 0])` ∈ [-1, +1] to the SHARED
/// ISV slot at `isv_slot_offset`. Per-step state, not an EMA —
/// slot is overwritten each launch (no FoldReset registry entry,
/// per the SP13 P0a registry comments). The bound is now structural
/// (softmax components in [0, 1] sum to 1 ⇒ pairwise diff in
/// [-1, +1]) per `pearl_bounded_modifier_outputs_require_structural_activation`;
/// the runtime tanh squash is retired.
///
/// Single-block tree-reduce (BLOCK_SIZE=256) — one shared-memory
/// float array reduces in lockstep per `feedback_no_atomicadd.md`.
/// `shared_mem_bytes = bdim × sizeof(f32)`.
///
/// Stream-ordered with the producer that wrote `aux_pred_dev`;
/// Stream-ordered with the producer that wrote `aux_softmax_dev`;
/// same-stream contract mirrors `launch_aux_heads_loss_ema`.
pub(crate) fn launch_aux_pred_to_isv_tanh(
&self,
aux_pred_dev: u64,
aux_softmax_dev: u64,
batch_size: i32,
k: i32,
isv_slot_offset: i32,
) -> Result<(), MLError> {
debug_assert!(self.isv_signals_dev_ptr != 0,
@@ -13300,8 +13342,9 @@ impl GpuDqnTrainer {
unsafe {
self.stream
.launch_builder(&self.aux_pred_to_isv_tanh_kernel)
.arg(&aux_pred_dev)
.arg(&aux_softmax_dev)
.arg(&batch_size)
.arg(&k)
.arg(&isv_slot_offset)
.arg(&isv_dev)
.launch(cfg)
@@ -13313,16 +13356,21 @@ impl GpuDqnTrainer {
/// SP13 Phase 0a P0a.T4 (2026-05-04): orchestrator for the per-step
/// aux directional-accuracy producer chain.
///
/// SP13 B1.1a (2026-05-05): reads softmax tile + i32 labels (replaces
/// the regression-mode pred/label_f32 pair); dir-acc kernel emits
/// 6 floats now (was 3 — added `n_down`/`n_up`/`n_skip` for HEALTH_DIAG
/// observability of the producer-less B1.1a label distribution).
///
/// Single-call API for `training_loop.rs` to fire the SP13 P0a
/// per-step metrics:
/// 1. `launch_aux_dir_acc_reduce` — reduce aux_nb_pred_buf +
/// aux_nb_label_buf into `aux_dir_acc_buf [3]`
/// (dir_acc, pos_pred_frac, pos_label_frac).
/// 1. `launch_aux_dir_acc_reduce` — reduce `aux_nb_softmax_buf` +
/// `aux_nb_label_buf` into `aux_dir_acc_buf [6]`
/// `(dir_acc, pos_pred_frac, pos_label_frac, n_down, n_up, n_skip)`.
/// 2. `launch_apply_fixed_alpha_ema` — short EMA (α=0.3) into
/// ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373], sentinel 0.5.
/// 3. `launch_apply_fixed_alpha_ema` — slow EMA (α=0.05) into
/// ISV[AUX_DIR_ACC_LONG_EMA_INDEX=374], sentinel 0.5.
/// 4. `launch_aux_pred_to_isv_tanh` — mean(tanh(aux_pred)) →
/// 4. `launch_aux_pred_to_isv_tanh` — mean(softmax[1]-softmax[0]) →
/// ISV[AUX_DIR_PREDICTION_INDEX=375].
///
/// All launches are stream-ordered; the producer's
@@ -13333,7 +13381,7 @@ impl GpuDqnTrainer {
/// invariant from commit `a5f23b28f`).
///
/// Pre-condition: the captured forward graph has run this step,
/// so `aux_nb_pred_buf` and `aux_nb_label_buf` are populated.
/// so `aux_nb_softmax_buf` and `aux_nb_label_buf` are populated.
/// Producer-only — no consumer kernel reads slots 373/374/375 in
/// this commit; Phase 0b wires the controller and the Q-head
/// input layer that reads slot 375.
@@ -13346,21 +13394,23 @@ impl GpuDqnTrainer {
DIR_ACC_EMA_SENTINEL,
};
let aux_pred_dev = self.aux_nb_pred_buf.raw_ptr();
let aux_label_dev = self.aux_nb_label_buf.raw_ptr();
let out_3_dev = self.aux_dir_acc_buf.dev_ptr;
let batch_size_i32 = self.config.batch_size as i32;
let aux_softmax_dev = self.aux_nb_softmax_buf.raw_ptr();
let aux_label_dev = self.aux_nb_label_buf.raw_ptr();
let out_6_dev = self.aux_dir_acc_buf.dev_ptr;
let batch_size_i32 = self.config.batch_size as i32;
let k_i32 = AUX_NEXT_BAR_K as i32;
// 1. Reduce → mapped-pinned aux_dir_acc_buf [3]
// 1. Reduce → mapped-pinned aux_dir_acc_buf [6]
self.launch_aux_dir_acc_reduce(
aux_pred_dev, aux_label_dev, batch_size_i32, out_3_dev,
aux_softmax_dev, aux_label_dev, batch_size_i32, k_i32, out_6_dev,
)?;
// 2. Short EMA (α=0.3) — fast tracker for aux-w controller deficit term.
// n=1: only out_3[0] = dir_acc feeds the EMA; pos_pred/pos_label slots
// are diagnostic-only and don't have ISV destinations.
// n=1: only out_6[0] = dir_acc feeds the EMA; pos_pred/pos_label/
// n_down/n_up/n_skip slots are diagnostic-only and don't have ISV
// destinations.
self.launch_apply_fixed_alpha_ema(
out_3_dev,
out_6_dev,
1,
AUX_DIR_ACC_SHORT_EMA_INDEX as i32,
0.3,
@@ -13369,18 +13419,20 @@ impl GpuDqnTrainer {
// 3. Slow EMA (α=0.05) — stagnation detector comparator.
self.launch_apply_fixed_alpha_ema(
out_3_dev,
out_6_dev,
1,
AUX_DIR_ACC_LONG_EMA_INDEX as i32,
0.05,
DIR_ACC_EMA_SENTINEL,
)?;
// 4. mean(tanh(aux_pred)) → ISV[375]. SHARED scalar (per-step
// overwrite, no EMA — see kernel header for layout rationale).
// 4. mean(softmax[1] - softmax[0]) → ISV[375]. SHARED scalar
// (per-step overwrite, no EMA — see kernel header for layout
// rationale). Structurally bounded in [-1, +1] (no tanh).
self.launch_aux_pred_to_isv_tanh(
aux_pred_dev,
aux_softmax_dev,
batch_size_i32,
k_i32,
AUX_DIR_PREDICTION_INDEX as i32,
)?;
@@ -15195,18 +15247,24 @@ impl GpuDqnTrainer {
/// off the trunk's just-produced `save_h_s2 [B, SH2]`. Sequence:
/// 1. regime label builder: `states[:, OFI_START + 29]` → uint8 [B]
/// 2. next-bar label gather: `next_states[:, 0]` → f32 [B] dedicated buf
/// 3. next-bar forward: Linear → ELU → Linear → pred[B, 1]
/// 3. next-bar forward: Linear → ELU → Linear → softmax[B, K=2]
/// 4. regime forward: Linear → ELU → Linear → logits[B, 5]
/// 5. next-bar MSE reduce → loss_scalar[1]
/// 5. next-bar softmax CE reduce → loss_scalar[1] + valid_count[1]
/// 6. regime CE reduce → loss_scalar[1] + correct_scalar[1]
///
/// SP13 B1.1a (2026-05-05): the next-bar head flipped from K=1 MSE
/// regression to K=2 softmax CE classification. The pre-B1.1a step 2
/// (strided_gather of `next_states[:, 0]` into a f32 label buffer)
/// is retired; the producer kernel for the new i32 labels lands in
/// B1.1b, so until then `aux_nb_label_buf` stays at its
/// `alloc_zeros` value.
///
/// Online-only — aux heads NEVER enter Bellman bootstrapping.
pub(crate) fn aux_heads_forward(&self) -> Result<(), MLError> {
let b = self.config.batch_size;
let sh2 = self.config.shared_h2;
let h_s2_ptr = self.ptrs.save_h_s2;
let states_ptr = self.ptrs.states_buf;
let next_states_ptr = self.ptrs.next_states_buf;
let param_sizes = compute_param_sizes(&self.config);
let on_w_ptrs = f32_weight_ptrs_from_base(self.ptrs.params_ptr, &param_sizes);
@@ -15237,61 +15295,28 @@ impl GpuDqnTrainer {
self.aux_rg_label_buf.raw_ptr(),
)?;
// Step 2: next-bar regression label = column 0 of next_states_buf
// (= `log_return` per `pipeline.rs`; first feature in `MARKET` group).
// Use `strided_gather` to extract into the DEDICATED `aux_nb_label_buf`
// (NOT aliased over a backward partial — the previous WIP aliased
// `aux_partial_nb_b2` here and triggered the F1/F2 regression).
// strided_gather signature: (src, dst, src_lda, col, dst_stride, total).
{
let label_dst_ptr = self.aux_nb_label_buf.raw_ptr();
let src_lda: i32 = state_dim_padded as i32;
let col: i32 = 0;
let dst_stride: i32 = 1;
let total: i32 = b as i32;
let blocks = ((b as u32 + 255) / 256).max(1);
unsafe {
self.stream
.launch_builder(&self.vsn_logit_gather_kernel)
.arg(&next_states_ptr)
.arg(&label_dst_ptr)
.arg(&src_lda)
.arg(&col)
.arg(&dst_stride)
.arg(&total)
.launch(LaunchConfig {
grid_dim: (blocks, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"aux_heads next-bar label gather (col 0 of next_states): {e}"
)))?;
}
}
// SP13 B1.1a (2026-05-05): the previous regression-mode head
// gathered `next_states[:, 0]` (= `log_return`) into a f32 label
// buffer here. The B1.1a head is a 2-class softmax classifier
// reading i32 direction labels in `{-1, 0, 1}`, and the producer
// kernel that fills those labels lands in B1.1b. Until then,
// `aux_nb_label_buf` stays at its `alloc_zeros` value (every
// sample receives label 0 = "down") — the model converges on
// "predict class 0 everywhere" until B1.1b lands the producer.
// This is intentional, known, and documented per the B1.1a brief.
// The strided_gather block + ISV[117] retirement comments are
// both retired in this commit.
// 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.
// Step 3: next-bar direction-classification forward (SP13 B1.1a:
// K=2 softmax). Writes hidden tile + logits tile + softmax tile.
self.aux_heads_fwd.forward_next_bar(
&self.stream,
h_s2_ptr,
nb_w1, nb_b1, nb_w2, nb_b2,
b, sh2,
b, sh2, AUX_NEXT_BAR_K,
self.aux_nb_hidden_buf.raw_ptr(),
self.aux_nb_pred_buf.raw_ptr(),
self.aux_nb_logits_buf.raw_ptr(),
self.aux_nb_softmax_buf.raw_ptr(),
)?;
// Step 4: regime classification forward.
@@ -15304,16 +15329,16 @@ impl GpuDqnTrainer {
self.aux_rg_logits_buf.raw_ptr(),
)?;
// Step 5: next-bar MSE reduce against the dedicated label buffer.
// 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.
// Step 5: next-bar softmax CE reduce (SP13 B1.1a). Reads the
// softmax tile + i32 labels; writes loss + B_valid scalars.
self.aux_heads_fwd.next_bar_loss_reduce(
&self.stream,
self.aux_nb_pred_buf.raw_ptr(),
self.aux_nb_softmax_buf.raw_ptr(),
self.aux_nb_label_buf.raw_ptr(),
b,
AUX_NEXT_BAR_K,
self.aux_nb_loss_scalar_buf.raw_ptr(),
self.aux_nb_valid_count_buf.raw_ptr(),
)?;
// Step 6: regime CE reduce.
@@ -15355,6 +15380,7 @@ impl GpuDqnTrainer {
let aux_w = self.aux_weight;
let aux_h = AUX_HIDDEN_DIM;
let aux_kr = AUX_REGIME_K;
let aux_knb = AUX_NEXT_BAR_K;
let param_sizes = compute_param_sizes(&self.config);
let on_w_ptrs = f32_weight_ptrs_from_base(self.ptrs.params_ptr, &param_sizes);
@@ -15366,19 +15392,18 @@ impl GpuDqnTrainer {
// Run BOTH backward kernels. They WRITE to disjoint output buffers
// (per-head partials + per-head dh_s2), so no inter-dependency beyond
// the just-completed forward.
// SP13 B1.1a (2026-05-05): softmax CE backward. Reads the saved
// softmax tile + i32 labels + B_valid scalar (written by loss
// reduce) so loss + gradient share the same `1/B_valid` divisor.
self.aux_heads_bwd.backward_next_bar(
&self.stream,
h_s2_ptr,
nb_w1, nb_w2,
self.aux_nb_hidden_buf.raw_ptr(),
self.aux_nb_pred_buf.raw_ptr(),
self.aux_nb_softmax_buf.raw_ptr(),
self.aux_nb_label_buf.raw_ptr(),
b, sh2,
// 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_nb_valid_count_buf.raw_ptr(),
b, sh2, aux_knb,
self.aux_partial_nb_w1.raw_ptr(),
self.aux_partial_nb_b1.raw_ptr(),
self.aux_partial_nb_w2.raw_ptr(),
@@ -15400,24 +15425,25 @@ impl GpuDqnTrainer {
self.aux_dh_s2_rg_buf.raw_ptr(),
)?;
// Reduce + SAXPY each of the 8 aux param tensors. Tensor layout (Commit A):
// [119] aux_nb_w1 [H, SH2] ← partial [B, H*SH2]
// [120] aux_nb_b1 [H] ← partial [B, H]
// [121] aux_nb_w2 [1, H] ← partial [B, H]
// [122] aux_nb_b2 [1] ← partial [B]
// Reduce + SAXPY each of the 8 aux param tensors. Tensor layout
// (SP13 B1.1a; next-bar K_NB flipped 1 → 2):
// [119] aux_nb_w1 [H, SH2] ← partial [B, H*SH2]
// [120] aux_nb_b1 [H] ← partial [B, H]
// [121] aux_nb_w2 [K_NB=2, H] ← partial [B, K_NB*H]
// [122] aux_nb_b2 [K_NB=2] ← partial [B, K_NB]
// [123] aux_rg_w1 [H, SH2]
// [124] aux_rg_b1 [H]
// [125] aux_rg_w2 [K, H]
// [126] aux_rg_b2 [K]
// Note: kernels emit (1/B)*sum_b grad — `aux_param_grad_reduce`
// [125] aux_rg_w2 [K_RG=5, H]
// [126] aux_rg_b2 [K_RG=5]
// Note: kernels emit (1/B_valid)*sum_b grad — `aux_param_grad_reduce`
// sums per-sample partials, no extra division. SAXPY alpha is
// `aux_weight` (small positive, [0.05, 0.3]).
let final_ptr = self.aux_param_grad_final_buf.raw_ptr();
let aux_param_specs: [(usize, u64, usize); 8] = [
(119, self.aux_partial_nb_w1.raw_ptr(), aux_h * sh2),
(120, self.aux_partial_nb_b1.raw_ptr(), aux_h),
(121, self.aux_partial_nb_w2.raw_ptr(), aux_h),
(122, self.aux_partial_nb_b2.raw_ptr(), 1),
(121, self.aux_partial_nb_w2.raw_ptr(), aux_knb * aux_h),
(122, self.aux_partial_nb_b2.raw_ptr(), aux_knb),
(123, self.aux_partial_rg_w1.raw_ptr(), aux_h * sh2),
(124, self.aux_partial_rg_b1.raw_ptr(), aux_h),
(125, self.aux_partial_rg_w2.raw_ptr(), aux_kr * aux_h),
@@ -16799,10 +16825,15 @@ impl GpuDqnTrainer {
module.load_function("aux_dir_acc_reduce_kernel")
.map_err(|e| MLError::ModelError(format!("sp13 aux_dir_acc_reduce load: {e}")))?
};
// SP13 B1.1a (2026-05-05): output grew 3 → 6 (added n_down/n_up/
// n_skip for HEALTH_DIAG observability of the B1.1a producer-less
// label distribution). Layout:
// [0] dir_acc, [1] pos_pred_frac, [2] pos_label_frac,
// [3] n_down, [4] n_up, [5] n_skip
let aux_dir_acc_buf = unsafe {
super::mapped_pinned::MappedF32Buffer::new(3)
super::mapped_pinned::MappedF32Buffer::new(6)
}.map_err(|e| MLError::ModelError(
format!("SP13 aux_dir_acc_buf alloc (3 f32): {e}")
format!("SP13 aux_dir_acc_buf alloc (6 f32): {e}")
))?;
// SP13 v3 P0a.T3 (2026-05-04): the per-step Hold-rate observer
// chain lives on the `GpuExperienceCollector` stream — the
@@ -18424,9 +18455,22 @@ impl GpuDqnTrainer {
let aux_kr = AUX_REGIME_K;
let aux_knb = AUX_NEXT_BAR_K;
let aux_nb_hidden_buf = alloc_f32(&stream, aux_b * aux_h, "aux_nb_hidden_buf")?;
let aux_nb_pred_buf = alloc_f32(&stream, aux_b * aux_knb, "aux_nb_pred_buf")?;
let aux_nb_label_buf = alloc_f32(&stream, aux_b, "aux_nb_label_buf")?;
// SP13 B1.1a (2026-05-05): next-bar head K_out flipped 1 → 2.
// - `aux_nb_logits_buf` (renamed from `aux_nb_pred_buf`) holds the
// `[B, K]` logits saved by forward.
// - `aux_nb_softmax_buf` (NEW) holds the `[B, K]` softmax tile
// read by 3 consumers (loss, backward, dir-acc, isv-tanh).
// - `aux_nb_label_buf` flipped f32 → i32; producer wires in B1.1b
// so the buffer stays at `alloc_zeros` value (label 0 = "down")
// until then.
// - `aux_nb_valid_count_buf` (NEW) carries `B_valid` from loss
// reduce to backward so they share the same `1/B_valid` divisor.
let aux_nb_logits_buf = alloc_f32(&stream, aux_b * aux_knb, "aux_nb_logits_buf")?;
let aux_nb_softmax_buf = alloc_f32(&stream, aux_b * aux_knb, "aux_nb_softmax_buf")?;
let aux_nb_label_buf = stream.alloc_zeros::<i32>(aux_b)
.map_err(|e| MLError::ModelError(format!("alloc aux_nb_label_buf (B1.1a i32 flip): {e}")))?;
let aux_nb_loss_scalar_buf = alloc_f32(&stream, 1, "aux_nb_loss_scalar_buf")?;
let aux_nb_valid_count_buf = alloc_f32(&stream, 1, "aux_nb_valid_count_buf")?;
let aux_rg_hidden_buf = alloc_f32(&stream, aux_b * aux_h, "aux_rg_hidden_buf")?;
let aux_rg_logits_buf = alloc_f32(&stream, aux_b * aux_kr, "aux_rg_logits_buf")?;
let aux_rg_label_buf = stream.alloc_zeros::<u8>(aux_b)
@@ -18437,14 +18481,19 @@ impl GpuDqnTrainer {
let aux_dh_s2_rg_buf = alloc_f32(&stream, aux_b * aux_sh2, "aux_dh_s2_rg_buf")?;
let aux_partial_nb_w1 = alloc_f32(&stream, aux_b * aux_h * aux_sh2, "aux_partial_nb_w1")?;
let aux_partial_nb_b1 = alloc_f32(&stream, aux_b * aux_h, "aux_partial_nb_b1")?;
let aux_partial_nb_w2 = alloc_f32(&stream, aux_b * aux_h, "aux_partial_nb_w2")?;
let aux_partial_nb_b2 = alloc_f32(&stream, aux_b, "aux_partial_nb_b2")?;
// SP13 B1.1a: nb_w2 partial grew [B, H] → [B, K, H]; nb_b2 partial
// grew [B] → [B, K]. Mirrors regime-head shapes with K_NB=2.
let aux_partial_nb_w2 = alloc_f32(&stream, aux_b * aux_knb * aux_h, "aux_partial_nb_w2")?;
let aux_partial_nb_b2 = alloc_f32(&stream, aux_b * aux_knb, "aux_partial_nb_b2")?;
let aux_partial_rg_w1 = alloc_f32(&stream, aux_b * aux_h * aux_sh2, "aux_partial_rg_w1")?;
let aux_partial_rg_b1 = alloc_f32(&stream, aux_b * aux_h, "aux_partial_rg_b1")?;
let aux_partial_rg_w2 = alloc_f32(&stream, aux_b * aux_kr * aux_h, "aux_partial_rg_w2")?;
let aux_partial_rg_b2 = alloc_f32(&stream, aux_b * aux_kr, "aux_partial_rg_b2")?;
// SP13 B1.1a: include the K_NB-scaled tensor sizes in the max.
// Both `aux_knb * aux_h` and `aux_knb` need coverage now.
let max_aux_tensor_len = (aux_h * aux_sh2)
.max(aux_kr * aux_h)
.max(aux_knb * aux_h)
.max(aux_h)
.max(aux_kr)
.max(aux_knb);
@@ -20220,9 +20269,11 @@ impl GpuDqnTrainer {
aux_heads_fwd,
aux_heads_bwd,
aux_nb_hidden_buf,
aux_nb_pred_buf,
aux_nb_logits_buf,
aux_nb_softmax_buf,
aux_nb_label_buf,
aux_nb_loss_scalar_buf,
aux_nb_valid_count_buf,
aux_rg_hidden_buf,
aux_rg_logits_buf,
aux_rg_label_buf,

58
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -4014,15 +4014,20 @@ impl DQNTrainer {
}
// SP13 Phase 0a P0a.T4 (2026-05-04): per-step aux
// directional-accuracy producer chain. The orchestrator
// reads the captured-forward-graph `aux_nb_pred_buf` +
// `aux_nb_label_buf` (just populated above by the same
// per-step block that fires `launch_aux_heads_loss_ema`),
// reduces them into the trainer's mapped-pinned
// `aux_dir_acc_buf [3]`, then chains:
// directional-accuracy producer chain. SP13 B1.1a
// (2026-05-05) flipped the upstream head from K=1 MSE
// regression to K=2 softmax CE classification — this
// orchestrator now reads the captured-forward-graph
// `aux_nb_softmax_buf [B, K]` + `aux_nb_label_buf [B] i32`
// (just populated above by the same per-step block that
// fires `launch_aux_heads_loss_ema`), reduces them into
// the trainer's mapped-pinned `aux_dir_acc_buf [6]`
// (was [3]; B1.1a added n_down/n_up/n_skip for
// HEALTH_DIAG observability), then chains:
// - fixed-α EMA (α=0.3, sentinel 0.5) → ISV[373]
// - fixed-α EMA (α=0.05, sentinel 0.5) → ISV[374]
// - mean(tanh(aux_pred)) → ISV[375]
// - mean(softmax[1]-softmax[0]) → ISV[375]
// (structural [-1, +1] bound; no runtime tanh)
// All on the same stream as the producer, so no host
// sync is required between captured graph and consumer.
// Producer-only — no consumer kernel reads slots 373/374/
@@ -4031,6 +4036,16 @@ impl DQNTrainer {
// post-cascade-fix invariant as `launch_aux_heads_loss_ema`:
// ISV producer launches MUST run BEFORE the HEALTH_DIAG
// line emit further down.
//
// B1.1a known degraded state: labels stay zero-init
// (`aux_nb_label_buf` is `alloc_zeros`-i32 every step;
// every sample receives label 0 = "down") until B1.1b
// lands the producer kernel that fills real -1/0/1 from
// the price trajectory. Until then, the model converges
// on "predict class 0 (down) everywhere" — the cascade
// is internally consistent (every consumer migrated
// atomically per `feedback_no_partial_refactor`); the
// labels are placeholder.
if let Some(ref fused) = self.fused_ctx {
if let Err(e) = fused.trainer().launch_sp13_aux_dir_metrics() {
tracing::warn!("SP13 P0a.T4 aux_dir_metrics launch failed: {e}");
@@ -4888,6 +4903,35 @@ impl DQNTrainer {
"HEALTH_DIAG[{}]: hold_pricing observed_rate={:.4} target={:.4} cost={:.6}",
epoch, hold_observed, hold_target, hold_cost,
);
// SP13 B1.1a (2026-05-05): per-class label distribution
// diagnostic. The 6-element `aux_dir_acc_buf` exposes
// `n_down / n_up / n_skip` directly from the GPU dir-acc
// reduce; B1.1a labels are zero-init (producer wires in
// B1.1b), so this line shows `n_down=B, n_up=0, n_skip=0`
// until B1.1b lands. Once B1.1b's producer fills real
// -1/0/1 labels, the distribution becomes diagnostic for
// the price-trajectory binarisation. `mask_frac = n_skip
// / B` is the fraction of bars excluded from the CE
// numerator. The `aux_dir_acc` short/long EMAs are read
// above; the per-class counts here add the missing
// observability that B1.1b's producer + smoke A will
// need.
let (n_down, n_up, n_skip, mask_frac) = if let Some(ref fused) = self.fused_ctx {
let dist = fused.trainer().aux_dir_acc_buf.read_all();
// Layout: [dir_acc, pos_pred_frac, pos_label_frac, n_down, n_up, n_skip]
let n_down = dist.get(3).copied().unwrap_or(0.0);
let n_up = dist.get(4).copied().unwrap_or(0.0);
let n_skip = dist.get(5).copied().unwrap_or(0.0);
let total = (n_down + n_up + n_skip).max(1.0);
(n_down, n_up, n_skip, n_skip / total)
} else {
(0.0, 0.0, 0.0, 0.0)
};
tracing::info!(
"HEALTH_DIAG[{}]: aux_b1_diag dir_acc_short={:.4} dir_acc_long={:.4} n_down={:.0} n_up={:.0} n_skip={:.0} mask_frac={:.4}",
epoch, dir_short, dir_long, n_down, n_up, n_skip, mask_frac,
);
}
// Phase 3 T3.5: MoE expert-utilisation + gate entropy HEALTH_DIAG.

821
crates/ml/tests/sp13_layer_b_oracle_tests.rs Normal file
View File

@@ -0,0 +1,821 @@
#![allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
//! SP13 Layer B Commit B1.1a (2026-05-05) GPU oracle tests.
//!
//! Validates the kernel rewrites that take the aux next-bar head from
//! K=1 MSE regression to K=2 softmax CE classification:
//! - `aux_next_bar_loss_reduce` (softmax + i32 labels → mean CE +
//! `B_valid` count; `-1` label = mask/skip)
//! - `aux_next_bar_backward` (softmax + i32 labels +
//! `B_valid` → per-sample dW/dh partials; masked rows zero)
//! - `aux_dir_acc_reduce_kernel` argmax over softmax-vs-i32-labels
//! (B1.1a-flipped — was f32 pred + f32 sign label)
//! - `aux_pred_to_isv_tanh_kernel` `mean(softmax[1] - softmax[0])`
//! into the SHARED ISV slot 375 (B1.1a-flipped — runtime tanh
//! retired; bound is structural)
//! - `LAYOUT_FINGERPRINT_CURRENT` regression — verifies the
//! fingerprint bumped from the pre-B1.1a value because the seed
//! names `PARAM_AUX_NB_W2 → PARAM_AUX_NB_W2_K2` (and `_B2`
//! equivalent) flipped.
//!
//! Per the B1.1a brief, this file covers tests 7-11 (CE loss
//! correctness), 12-13 (dir_acc B1.1a additions), 14-15 (isv_tanh
//! B1.1a additions), and 16+18 (fingerprint regression + HEALTH_DIAG
//! snap stability). Tests 1-6 (producer kernel) and 17 (end-to-end
//! round-trip) defer to B1.1b — they require the producer kernel +
//! replay direct path which aren't wired in B1.1a.
//!
//! Per `feedback_no_cpu_test_fallbacks.md`: GPU oracle only, no CPU
//! reference impl. Per `feedback_no_htod_htoh_only_mapped_pinned.md`:
//! every CPU↔GPU buffer is a `MappedF32Buffer` / `MappedI32Buffer`
//! (cuMemHostAlloc with DEVICEMAP|PORTABLE); zero `htod_copy`, zero
//! `dtoh_sync_copy`. Tests are NOT exempt from this rule.
//!
//! All GPU tests are `#[ignore = "requires GPU"]`-gated to match every
//! other GPU oracle test in this crate (sp4/sp5/sp11/sp12/sp13_phase0).
//! Run on a GPU host:
//!
//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//! cargo test -p ml --test sp13_layer_b_oracle_tests --features cuda \
//! -- --ignored --nocapture
//!
//! The fingerprint regression test (test 16) is a pure-Rust const-eval
//! test — does NOT require a GPU and runs without `--ignored`.
#![cfg(feature = "cuda")]
use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
// ── Test-only cubin handles ───────────────────────────────────────────────
//
// All three kernels live in the same `aux_heads_kernel.cubin` (loss +
// backward) plus the dedicated `aux_dir_acc_reduce_kernel.cubin` and
// `aux_pred_to_isv_tanh_kernel.cubin`. The cubins are emitted by
// `crates/ml/build.rs` into `OUT_DIR`; including them here mirrors the
// production loaders' include_bytes! pattern.
const SP13_AUX_HEADS_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_heads_kernel.cubin"));
const SP13_AUX_DIR_ACC_REDUCE_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_dir_acc_reduce_kernel.cubin"));
const SP13_AUX_PRED_TO_ISV_TANH_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_pred_to_isv_tanh_kernel.cubin"));
// ── Stream + kernel loaders ──────────────────────────────────────────────
fn make_test_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_loss_reduce(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_AUX_HEADS_CUBIN.to_vec())
.expect("load aux_heads cubin");
module
.load_function("aux_next_bar_loss_reduce")
.expect("load aux_next_bar_loss_reduce function")
}
fn load_backward(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_AUX_HEADS_CUBIN.to_vec())
.expect("load aux_heads cubin");
module
.load_function("aux_next_bar_backward")
.expect("load aux_next_bar_backward function")
}
fn load_dir_acc(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_AUX_DIR_ACC_REDUCE_CUBIN.to_vec())
.expect("load aux_dir_acc_reduce cubin");
module
.load_function("aux_dir_acc_reduce_kernel")
.expect("load aux_dir_acc_reduce_kernel function")
}
fn load_isv_tanh(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP13_AUX_PRED_TO_ISV_TANH_CUBIN.to_vec())
.expect("load aux_pred_to_isv_tanh cubin");
module
.load_function("aux_pred_to_isv_tanh_kernel")
.expect("load aux_pred_to_isv_tanh_kernel function")
}
// ── CE loss reduce scaffold ──────────────────────────────────────────────
const AUX_BLOCK: u32 = 256;
/// Drive one launch of `aux_next_bar_loss_reduce` with the given softmax
/// tile + i32 labels. Returns `(mean_ce, b_valid)`.
fn run_loss_reduce(
stream: &Arc<CudaStream>,
f: &CudaFunction,
softmax: &[f32],
labels: &[i32],
k: i32,
) -> (f32, f32) {
let n = labels.len();
assert_eq!(
softmax.len(),
n * k as usize,
"softmax must be [B, K]; got {} for B={} K={}",
softmax.len(),
n,
k
);
let alloc_smx = (n * k as usize).max(1);
let alloc_lbl = n.max(1);
let softmax_buf = unsafe { MappedF32Buffer::new(alloc_smx) }
.expect("alloc softmax buffer (mapped-pinned)");
let labels_buf = unsafe { MappedI32Buffer::new(alloc_lbl) }
.expect("alloc labels buffer (mapped-pinned)");
if n > 0 {
softmax_buf.write_from_slice(softmax);
labels_buf.write_from_slice(labels);
}
let loss_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc loss buffer (mapped-pinned)");
let valid_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc valid_count buffer (mapped-pinned)");
let softmax_dev = softmax_buf.dev_ptr;
let labels_dev = labels_buf.dev_ptr;
let loss_dev = loss_buf.dev_ptr;
let valid_dev = valid_buf.dev_ptr;
let b_i32: i32 = n as i32;
unsafe {
stream
.launch_builder(f)
.arg(&softmax_dev)
.arg(&labels_dev)
.arg(&b_i32)
.arg(&k)
.arg(&loss_dev)
.arg(&valid_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (AUX_BLOCK, 1, 1),
shared_mem_bytes: 2 * AUX_BLOCK * std::mem::size_of::<f32>() as u32,
})
.expect("launch aux_next_bar_loss_reduce");
}
stream.synchronize().expect("sync after aux_next_bar_loss_reduce");
(loss_buf.read_all()[0], valid_buf.read_all()[0])
}
// ── CE backward scaffold (test 10/11) ────────────────────────────────────
/// Drive one launch of `aux_next_bar_backward` with hand-crafted forward
/// state (uniform identity-ish weights so the test focuses on the
/// CE backward arithmetic, not the matmul). Returns the per-sample
/// `db2_partial [B, K]` slice (which equals `d_logits[b, kc]` directly
/// per the kernel's `db2_partial[b, kc] = sh_dlogits[kc]` write).
///
/// Hand-crafted forward state:
/// * h_s2[B, SH2] — zeros (so dh_s2 outputs zero; not under test)
/// * w1[H, SH2] — zeros
/// * w2[K, H] — zeros (so d_h_post = 0 → dW1/db1 = 0)
/// * hidden_post[B, H] — zeros (so ELU-bwd factor is 1.0)
///
/// With all-zero h_post and w2, the backward kernel's d_h_pre is zero,
/// so dW1, db1, dW2, dh_s2 are all zero. The interesting output is
/// `db2_partial[b, kc] = sh_dlogits[kc] = (softmax[b,kc] - one_hot)/B_valid`
/// for valid rows, zero for masked rows.
#[allow(clippy::too_many_arguments)]
fn run_backward_db2_only(
stream: &Arc<CudaStream>,
f: &CudaFunction,
softmax: &[f32],
labels: &[i32],
valid_count: f32,
b: usize,
sh2: usize,
k: i32,
) -> Vec<f32> {
let h: usize = 32; // AUX_HIDDEN_DIM
let kk = k as usize;
// Allocate all forward + partial buffers (mapped-pinned for parity
// with the production launcher; the kernel only reads dev_ptr).
let h_s2_buf = unsafe { MappedF32Buffer::new(b * sh2) }.unwrap();
let w1_buf = unsafe { MappedF32Buffer::new(h * sh2) }.unwrap();
let w2_buf = unsafe { MappedF32Buffer::new(kk * h) }.unwrap();
let hidden_post_buf = unsafe { MappedF32Buffer::new(b * h) }.unwrap();
let softmax_buf = unsafe { MappedF32Buffer::new(b * kk) }.unwrap();
let labels_buf = unsafe { MappedI32Buffer::new(b) }.unwrap();
let valid_buf = unsafe { MappedF32Buffer::new(1) }.unwrap();
let dw1_buf = unsafe { MappedF32Buffer::new(b * h * sh2) }.unwrap();
let db1_buf = unsafe { MappedF32Buffer::new(b * h) }.unwrap();
let dw2_buf = unsafe { MappedF32Buffer::new(b * kk * h) }.unwrap();
let db2_buf = unsafe { MappedF32Buffer::new(b * kk) }.unwrap();
let dh_s2_buf = unsafe { MappedF32Buffer::new(b * sh2) }.unwrap();
softmax_buf.write_from_slice(softmax);
labels_buf.write_from_slice(labels);
valid_buf.write_from_slice(&[valid_count]);
let h_s2_dev = h_s2_buf.dev_ptr;
let w1_dev = w1_buf.dev_ptr;
let w2_dev = w2_buf.dev_ptr;
let hpost_dev = hidden_post_buf.dev_ptr;
let softmax_dev = softmax_buf.dev_ptr;
let labels_dev = labels_buf.dev_ptr;
let valid_dev = valid_buf.dev_ptr;
let dw1_dev = dw1_buf.dev_ptr;
let db1_dev = db1_buf.dev_ptr;
let dw2_dev = dw2_buf.dev_ptr;
let db2_dev = db2_buf.dev_ptr;
let dhs2_dev = dh_s2_buf.dev_ptr;
let b_i32: i32 = b as i32;
let sh2_i32: i32 = sh2 as i32;
let smem_bytes = ((2 * h as u32) + (k as u32)) * std::mem::size_of::<f32>() as u32;
unsafe {
stream
.launch_builder(f)
.arg(&h_s2_dev)
.arg(&w1_dev)
.arg(&w2_dev)
.arg(&hpost_dev)
.arg(&softmax_dev)
.arg(&labels_dev)
.arg(&valid_dev)
.arg(&b_i32)
.arg(&sh2_i32)
.arg(&k)
.arg(&dw1_dev)
.arg(&db1_dev)
.arg(&dw2_dev)
.arg(&db2_dev)
.arg(&dhs2_dev)
.launch(LaunchConfig {
grid_dim: (b as u32, 1, 1),
block_dim: (AUX_BLOCK, 1, 1),
shared_mem_bytes: smem_bytes,
})
.expect("launch aux_next_bar_backward");
}
stream.synchronize().expect("sync after aux_next_bar_backward");
db2_buf.read_all()
}
// ── dir_acc + isv_tanh helpers (mirror sp13_phase0 launchers) ───────────
fn run_dir_acc(
stream: &Arc<CudaStream>,
f: &CudaFunction,
softmax: &[f32],
labels: &[i32],
k: i32,
) -> [f32; 6] {
let n = labels.len();
let alloc_smx = (n * k as usize).max(1);
let alloc_lbl = n.max(1);
let softmax_buf = unsafe { MappedF32Buffer::new(alloc_smx) }.unwrap();
let labels_buf = unsafe { MappedI32Buffer::new(alloc_lbl) }.unwrap();
if n > 0 {
softmax_buf.write_from_slice(softmax);
labels_buf.write_from_slice(labels);
}
let out_buf = unsafe { MappedF32Buffer::new(6) }.unwrap();
let bdim: u32 = 256;
let b_i32: i32 = n as i32;
unsafe {
stream
.launch_builder(f)
.arg(&softmax_buf.dev_ptr)
.arg(&labels_buf.dev_ptr)
.arg(&b_i32)
.arg(&k)
.arg(&out_buf.dev_ptr)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: 6 * bdim * std::mem::size_of::<i32>() as u32,
})
.expect("launch aux_dir_acc_reduce_kernel");
}
stream.synchronize().expect("sync after dir_acc");
let v = out_buf.read_all();
[v[0], v[1], v[2], v[3], v[4], v[5]]
}
fn run_isv_tanh(
stream: &Arc<CudaStream>,
f: &CudaFunction,
softmax: &[f32],
k: i32,
) -> f32 {
let n = if k > 0 { softmax.len() / k as usize } else { 0 };
let alloc = (n * k as usize).max(1);
let smx_buf = unsafe { MappedF32Buffer::new(alloc) }.unwrap();
if !softmax.is_empty() {
smx_buf.write_from_slice(softmax);
}
let isv_buf = unsafe { MappedF32Buffer::new(1) }.unwrap();
let bdim: u32 = 256;
let b_i32: i32 = n as i32;
let isv_off: i32 = 0;
unsafe {
stream
.launch_builder(f)
.arg(&smx_buf.dev_ptr)
.arg(&b_i32)
.arg(&k)
.arg(&isv_off)
.arg(&isv_buf.dev_ptr)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: bdim * std::mem::size_of::<f32>() as u32,
})
.expect("launch aux_pred_to_isv_tanh_kernel");
}
stream.synchronize().expect("sync after isv_tanh");
isv_buf.read_all()[0]
}
// ─── Test 7: CE loss single-row hand-crafted ──────────────────────────────
/// Single sample, K=2. softmax = [0.25, 0.75], label = 1 (up).
/// Expected loss = -log(0.75) ≈ 0.2877. valid_count = 1.
#[test]
#[ignore = "requires GPU"]
fn ce_loss_single_row_handcrafted() {
let stream = make_test_stream();
let f = load_loss_reduce(&stream);
let softmax = vec![0.25_f32, 0.75];
let labels = vec![1_i32];
let (loss, valid) = run_loss_reduce(&stream, &f, &softmax, &labels, 2);
let expected = (-0.75_f32.ln()) / 1.0;
assert!(
(loss - expected).abs() < 1e-5,
"expected loss = {expected}, got {loss}"
);
assert!(
(valid - 1.0).abs() < 1e-6,
"expected valid_count = 1.0, got {valid}"
);
}
// ─── Test 8: CE loss batch of 4 mixed ─────────────────────────────────────
/// 4 rows: 3 valid + 1 mask.
/// rows 0,1: softmax = [0.2, 0.8], label = 1 → CE = -log(0.8) each.
/// row 2: softmax = [0.9, 0.1], label = 0 → CE = -log(0.9).
/// row 3: softmax = [0.5, 0.5], label = -1 → masked, contributes 0.
/// Expected: B_valid = 3; mean = (2*-ln(0.8) + -ln(0.9)) / 3.
#[test]
#[ignore = "requires GPU"]
fn ce_loss_batch_mixed() {
let stream = make_test_stream();
let f = load_loss_reduce(&stream);
let softmax = vec![
0.2_f32, 0.8,
0.2, 0.8,
0.9, 0.1,
0.5, 0.5,
];
let labels = vec![1_i32, 1, 0, -1];
let (loss, valid) = run_loss_reduce(&stream, &f, &softmax, &labels, 2);
let numer = 2.0 * (-0.8_f32.ln()) + (-0.9_f32.ln());
let expected = numer / 3.0;
assert!(
(loss - expected).abs() < 1e-5,
"expected mean CE = {expected}, got {loss}"
);
assert!(
(valid - 3.0).abs() < 1e-6,
"expected valid_count = 3.0, got {valid}"
);
}
// ─── Test 9: CE loss all-skip batch ───────────────────────────────────────
/// 4 rows all masked → B_valid = 0; loss = 0/max(0,1) = 0 (no NaN).
/// `valid_count_out` reports the actual count (0.0).
#[test]
#[ignore = "requires GPU"]
fn ce_loss_all_skip_returns_zero_no_nan() {
let stream = make_test_stream();
let f = load_loss_reduce(&stream);
let softmax = vec![0.5_f32; 8]; // 4 rows, all balanced
let labels = vec![-1_i32; 4];
let (loss, valid) = run_loss_reduce(&stream, &f, &softmax, &labels, 2);
assert!(
loss.is_finite(),
"expected finite loss on all-mask batch, got {loss}"
);
assert!(
loss.abs() < 1e-6,
"expected loss = 0.0 on all-mask batch, got {loss}"
);
assert!(
valid.abs() < 1e-6,
"expected valid_count = 0.0 on all-mask batch, got {valid}"
);
}
// ─── Test 10: CE backward single-row ──────────────────────────────────────
/// Single row, K=2. softmax = [0.3, 0.7], label = 1 (up).
/// Expected `db2_partial[0, :] = (softmax - one_hot(1)) / B_valid`
/// = ([0.3, 0.7] - [0, 1]) / 1
/// = [0.3, -0.3].
#[test]
#[ignore = "requires GPU"]
fn ce_backward_single_row_handcrafted() {
let stream = make_test_stream();
let f = load_backward(&stream);
let softmax = vec![0.3_f32, 0.7];
let labels = vec![1_i32];
let valid = 1.0_f32;
let db2 = run_backward_db2_only(&stream, &f, &softmax, &labels, valid, 1, 4, 2);
assert!(
(db2[0] - 0.3).abs() < 1e-5,
"expected db2[0] = 0.3, got {}",
db2[0]
);
assert!(
(db2[1] - (-0.3)).abs() < 1e-5,
"expected db2[1] = -0.3, got {}",
db2[1]
);
}
// ─── Test 11: CE backward batch of 4 mixed ────────────────────────────────
/// 4 rows: 3 valid + 1 mask. Batch_valid = 3.
/// row 0: softmax = [0.2, 0.8], label = 1 → d_logits = [0.2, -0.2] / 3
/// row 1: softmax = [0.9, 0.1], label = 0 → d_logits = [-0.1, 0.1] / 3
/// row 2: softmax = [0.4, 0.6], label = 1 → d_logits = [0.4, -0.4] / 3
/// row 3: softmax = [0.5, 0.5], label = -1 → masked → zero across K.
#[test]
#[ignore = "requires GPU"]
fn ce_backward_batch_mixed() {
let stream = make_test_stream();
let f = load_backward(&stream);
let softmax = vec![
0.2_f32, 0.8,
0.9, 0.1,
0.4, 0.6,
0.5, 0.5,
];
let labels = vec![1_i32, 0, 1, -1];
let valid = 3.0_f32;
let db2 = run_backward_db2_only(&stream, &f, &softmax, &labels, valid, 4, 4, 2);
// Row 0
assert!((db2[0] - 0.2 / 3.0).abs() < 1e-5, "db2[0,0] mismatch: {}", db2[0]);
assert!((db2[1] - (-0.2) / 3.0).abs() < 1e-5, "db2[0,1] mismatch: {}", db2[1]);
// Row 1
assert!((db2[2] - (-0.1) / 3.0).abs() < 1e-5, "db2[1,0] mismatch: {}", db2[2]);
assert!((db2[3] - 0.1 / 3.0).abs() < 1e-5, "db2[1,1] mismatch: {}", db2[3]);
// Row 2
assert!((db2[4] - 0.4 / 3.0).abs() < 1e-5, "db2[2,0] mismatch: {}", db2[4]);
assert!((db2[5] - (-0.4) / 3.0).abs() < 1e-5, "db2[2,1] mismatch: {}", db2[5]);
// Row 3 (masked) → zero
assert!(db2[6].abs() < 1e-6, "db2[3,0] expected 0 (masked), got {}", db2[6]);
assert!(db2[7].abs() < 1e-6, "db2[3,1] expected 0 (masked), got {}", db2[7]);
}
// ─── Test 12: dir_acc argmax correctness on hand-crafted ─────────────────
/// 4 rows: argmax-vs-label cases:
/// row 0: softmax [0.1, 0.9], label 1 → correct
/// row 1: softmax [0.1, 0.9], label 1 → correct
/// row 2: softmax [0.1, 0.9], label 0 → wrong
/// row 3: softmax [0.5, 0.5], label -1 → masked
/// Valid bars = 3 (2 correct + 1 wrong). dir_acc = 2/3.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_handcrafted_argmax() {
let stream = make_test_stream();
let f = load_dir_acc(&stream);
let softmax = vec![
0.1_f32, 0.9,
0.1, 0.9,
0.1, 0.9,
0.5, 0.5,
];
let labels = vec![1_i32, 1, 0, -1];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
assert!(
(out[0] - (2.0 / 3.0)).abs() < 1e-5,
"expected dir_acc = 2/3, got {}",
out[0]
);
assert!((out[3] - 1.0).abs() < 1e-6, "n_down expected 1, got {}", out[3]);
assert!((out[4] - 2.0).abs() < 1e-6, "n_up expected 2, got {}", out[4]);
assert!((out[5] - 1.0).abs() < 1e-6, "n_skip expected 1, got {}", out[5]);
}
// ─── Test 13: dir_acc all-skip (NaN-safe) ────────────────────────────────
/// All-mask batch → sentinel 0.5 in slots 0..3, raw counts 0/0/B in
/// slots 3..6. NO NaN anywhere.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_all_skip_nan_safe() {
let stream = make_test_stream();
let f = load_dir_acc(&stream);
let softmax = vec![0.5_f32; 8];
let labels = vec![-1_i32; 4];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
for (i, v) in out.iter().enumerate() {
assert!(v.is_finite(), "slot {i} non-finite: {v}");
}
assert!(
(out[0] - 0.5).abs() < 1e-6,
"expected dir_acc sentinel 0.5, got {}",
out[0]
);
assert!((out[3]).abs() < 1e-6, "n_down expected 0, got {}", out[3]);
assert!((out[4]).abs() < 1e-6, "n_up expected 0, got {}", out[4]);
assert!((out[5] - 4.0).abs() < 1e-6, "n_skip expected 4, got {}", out[5]);
}
// ─── Test 14: aux_pred_to_isv_tanh bounded in [-1, +1] (3 fuzz) ──────────
/// Three randomly-shaped softmax tiles; each must produce mean diff
/// inside `[-1, +1]` (structural bound). Uses a simple PRNG so the
/// test is deterministic across runs.
#[test]
#[ignore = "requires GPU"]
fn isv_tanh_bounded_fuzz() {
let stream = make_test_stream();
let f = load_isv_tanh(&stream);
// Three deterministic fuzz inputs: skewed up, skewed down, mixed.
let cases: [Vec<f32>; 3] = [
// 8 rows skewed up: each row ≈ [0.2, 0.8].
vec![
0.2, 0.8, 0.15, 0.85, 0.25, 0.75, 0.3, 0.7,
0.18, 0.82, 0.22, 0.78, 0.21, 0.79, 0.27, 0.73,
],
// 8 rows skewed down: each row ≈ [0.8, 0.2].
vec![
0.8, 0.2, 0.85, 0.15, 0.75, 0.25, 0.7, 0.3,
0.82, 0.18, 0.78, 0.22, 0.79, 0.21, 0.73, 0.27,
],
// 4 mixed rows.
vec![0.1, 0.9, 0.9, 0.1, 0.4, 0.6, 0.6, 0.4],
];
for (i, softmax) in cases.iter().enumerate() {
let out = run_isv_tanh(&stream, &f, softmax, 2);
assert!(
out >= -1.0 && out <= 1.0,
"case {i}: out {out} outside [-1, +1] structural bound"
);
assert!(out.is_finite(), "case {i}: non-finite out {out}");
}
}
// ─── Test 15: aux_pred_to_isv_tanh mean correctness on hand-crafted ──────
/// 3 rows: [0.2, 0.8], [0.6, 0.4], [0.5, 0.5].
/// Diffs: +0.6, -0.2, 0.0
/// Mean: (0.6 - 0.2 + 0.0) / 3 = 0.4 / 3 ≈ 0.1333.
#[test]
#[ignore = "requires GPU"]
fn isv_tanh_mean_handcrafted() {
let stream = make_test_stream();
let f = load_isv_tanh(&stream);
let softmax = vec![0.2_f32, 0.8, 0.6, 0.4, 0.5, 0.5];
let out = run_isv_tanh(&stream, &f, &softmax, 2);
let expected = (0.6 - 0.2 + 0.0) / 3.0;
assert!(
(out - expected).abs() < 1e-5,
"expected mean diff = {expected}, got {out}"
);
}
// ─── Test 16: Fingerprint regression (CPU-only; no GPU required) ─────────
/// Inline FNV-1a 64-bit hash for the regression test. Mirrors the
/// `const fn fnv1a_64` in `gpu_dqn_trainer.rs` so the test computes
/// the same hash the production fingerprint uses.
const fn fnv1a_64(bytes: &[u8]) -> u64 {
const OFFSET: u64 = 0xcbf29ce484222325;
const PRIME: u64 = 0x00000100000001b3;
let mut h: u64 = OFFSET;
let mut i = 0;
while i < bytes.len() {
h ^= bytes[i] as u64;
h = h.wrapping_mul(PRIME);
i += 1;
}
h
}
/// Pre-B1.1a seed — the exact same seed string as the post-B1.0 layout
/// fingerprint EXCEPT the next-bar tensor names are
/// `PARAM_AUX_NB_W2` / `PARAM_AUX_NB_B2` (regression-mode names) rather
/// than the B1.1a `_K2` suffix variants. Hashing this constant gives
/// the pre-B1.1a fingerprint that the bump must differ from. Any
/// silent revert of the rename would make this test fail.
///
/// Layout: this is the verbatim seed string from `gpu_dqn_trainer.rs`'s
/// `layout_fingerprint_seed()` at HEAD `75e94858c` (B1.0). The only
/// section that changed in B1.1a is the `PARAM_AUX_NB_W2/B2 →
/// PARAM_AUX_NB_W2_K2/B2_K2` renames.
const PRE_B1_1A_SEED: &[u8] = b"SLOT_0_Q_DRIFT=0;\
SLOT_1_GRAD_NORM_EMA=1;\
SLOT_2_TD_ERR_EMA=2;\
SLOT_3_ENS_VAR_EMA=3;\
SLOT_4_ENS_VAR_VEL=4;\
SLOT_5_REWARD_EMA=5;\
SLOT_6_ATOM_UTIL_EMA=6;\
SLOT_7_LOSS_EMA=7;\
SLOT_8_ADX_EMA=8;\
SLOT_9_REGIME_DISAGREE=9;\
SLOT_10_REGIME_VEL_EMA=10;\
SLOT_11_REGIME_STABILITY=11;\
LEARNING_HEALTH=12;\
Q_MAG_MEAN_QUARTER=13;Q_MAG_MEAN_HALF=14;Q_MAG_MEAN_FULL=15;Q_ABS_REF=16;\
Q_DIR_MEAN_SHORT=17;Q_DIR_MEAN_HOLD=18;Q_DIR_MEAN_LONG=19;Q_DIR_MEAN_FLAT=20;Q_DIR_ABS_REF=21;\
SHARPE_EMA=22;\
V_CENTER_DIR=23;V_HALF_DIR=24;V_CENTER_MAG=25;V_HALF_MAG=26;\
V_CENTER_ORD=27;V_HALF_ORD=28;V_CENTER_URG=29;V_HALF_URG=30;\
GRAD_NORM_TARGET_DIR=31;GRAD_NORM_TARGET_MAG=32;GRAD_NORM_TARGET_ORD=33;GRAD_NORM_TARGET_URG=34;\
GRAD_SCALE_LIMIT=35;IQL_BRANCH_SCALE_FLOOR=36;\
EPOCH_IDX=39;TOTAL_EPOCHS=40;\
EPSILON_EFF=41;TAU_EFF=42;\
GAMMA_DIR_EFF=43;GAMMA_MAG_EFF=44;GAMMA_ORD_EFF=45;GAMMA_URG_EFF=46;\
KELLY_CAP_EFF=47;CQL_ALPHA=48;PLAN_THRESHOLD=49;\
Q_P05_DIR=50;Q_P05_MAG=51;Q_P05_ORD=52;Q_P05_URG=53;\
Q_P95_DIR=54;Q_P95_MAG=55;Q_P95_ORD=56;Q_P95_URG=57;\
TLOB_REGIME_FOCUS_EMA=60;\
REWARD_POPART_EMA=63;REWARD_CF_EMA=64;REWARD_TRAIL_EMA=65;\
REWARD_MICRO_EMA=66;REWARD_OPP_COST_EMA=67;REWARD_BONUS_EMA=68;\
TRADE_ATTEMPT_RATE_EMA=71;TRADE_TARGET_RATE=72;\
READINESS_EMA=75;\
STATE_KL_EMA=78;STATE_KL_AMP=79;\
SEED_STEPS_TARGET=82;SEED_STEPS_DONE=83;SEED_FRAC_EMA=84;\
VSN_MAG_EMA=87;VSN_DIR_EMA=88;MAMBA2_RETENTION_EMA=89;\
TARGET_DRIFT_MAG_EMA=92;TARGET_DRIFT_DIR_EMA=93;\
H_S2_RMS_EMA=96;\
IQN_Q_P05_EMA=99;IQN_Q_P25_EMA=100;IQN_Q_P75_EMA=101;IQN_Q_P95_EMA=102;\
VSN_MASK_GROUP_0_EMA=105;VSN_MASK_GROUP_1_EMA=106;VSN_MASK_GROUP_2_EMA=107;\
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;\
MOE_EXPERT_UTIL_EMA_BASE=118;MOE_EXPERT_UTIL_EMA_COUNT=8;\
MOE_GATE_ENTROPY_EMA=126;\
MOE_LAMBDA_EFF=128;\
Q_DRIFT_RATE=129;\
FOLD_WARMUP_FACTOR=130;\
TARGET_Q_BOUND=131;\
ATOM_POS_BOUND_BRANCH_0=132;ATOM_POS_BOUND_BRANCH_1=133;ATOM_POS_BOUND_BRANCH_2=134;ATOM_POS_BOUND_BRANCH_3=135;\
WEIGHT_BOUND_GROUP_0=136;WEIGHT_BOUND_GROUP_1=137;WEIGHT_BOUND_GROUP_2=138;WEIGHT_BOUND_GROUP_3=139;\
WEIGHT_BOUND_GROUP_4=140;WEIGHT_BOUND_GROUP_5=141;WEIGHT_BOUND_GROUP_6=142;WEIGHT_BOUND_GROUP_7=143;\
ADAM_M_BOUND_GROUP_0=144;ADAM_M_BOUND_GROUP_1=145;ADAM_M_BOUND_GROUP_2=146;ADAM_M_BOUND_GROUP_3=147;\
ADAM_M_BOUND_GROUP_4=148;ADAM_M_BOUND_GROUP_5=149;ADAM_M_BOUND_GROUP_6=150;ADAM_M_BOUND_GROUP_7=151;\
ADAM_V_BOUND_GROUP_0=152;ADAM_V_BOUND_GROUP_1=153;ADAM_V_BOUND_GROUP_2=154;ADAM_V_BOUND_GROUP_3=155;\
ADAM_V_BOUND_GROUP_4=156;ADAM_V_BOUND_GROUP_5=157;ADAM_V_BOUND_GROUP_6=158;ADAM_V_BOUND_GROUP_7=159;\
WD_RATE_GROUP_0=160;WD_RATE_GROUP_1=161;WD_RATE_GROUP_2=162;WD_RATE_GROUP_3=163;\
WD_RATE_GROUP_4=164;WD_RATE_GROUP_5=165;WD_RATE_GROUP_6=166;WD_RATE_GROUP_7=167;\
GRAD_CLIP_BOUND=168;H_S2_BOUND=169;L1_LAMBDA_TRUNK=170;\
BW_D_H_S2_BOUND=171;Q_DIR_GRAD_BOUND=172;\
SP5_BASE=174;ATOM_V_CENTER_BASE=174;ATOM_V_HALF_BASE=178;ATOM_HEADROOM_BASE=182;\
ATOM_CLIP_RATE_BASE=186;BUDGET_C51_BASE=190;BUDGET_IQN_BASE=194;BUDGET_CQL_BASE=198;\
BUDGET_ENS_BASE=202;FLATNESS_BASE=206;NOISY_SIGMA_BASE=210;SIGMA_FRACTION_BASE=214;\
BRANCH_ENTROPY_BASE=218;Q_VAR_PER_BRANCH_BASE=222;ADAM_BETA1_BASE=226;ADAM_BETA2_BASE=234;\
ADAM_EPS_BASE=242;IQN_TAU_BASE=250;TRAIL_DIST_PER_DIR_BASE=270;ATOM_NUM_ATOMS_BASE=274;\
KELLY_F_SMOOTH=280;CONVICTION_SMOOTH=281;TRADE_VAR_SMOOTH=282;\
KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\
PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\
HEALTH_SCORE=290;Q_GAP_NORM=291;Q_VAR_NORM=292;GRAD_NORM_NORM=293;\
TRAINING_SHARPE_EMA=294;MAX_DD_EMA=295;LOW_DD_RATIO=296;\
LB_DIFF_VAR_CQL_BASE=297;LB_SAMPLE_VAR_CQL_BASE=301;\
LB_DIFF_VAR_C51_BASE=305;LB_SAMPLE_VAR_C51_BASE=309;\
LB_CQL_ACTIVE_BASE=313;LB_C51_ACTIVE_BASE=317;\
TRAIN_ACTIVE_FRAC_INDEX=321;\
LB_MAX_BUDGET_CQL_BASE=322;LB_MAX_BUDGET_C51_BASE=326;\
KELLY_WARMUP_FLOOR=330;Q_VAR_MAG_EMA=331;\
INTENT_EVAL_DIVERGENCE=332;KELLY_SAMPLE_COUNT_TARGET=333;\
KELLY_DIVERGENCE_TARGET=334;KELLY_TEMPORAL_TARGET=335;\
EVAL_DIST_Q=336;EVAL_DIST_H=337;EVAL_DIST_F=338;\
EVAL_THOMPSON_TEMP=339;\
REWARD_POPART_WEIGHT=340;REWARD_CF_WEIGHT=341;REWARD_TRAIL_WEIGHT=342;\
REWARD_MICRO_WEIGHT=343;REWARD_OPP_COST_WEIGHT=344;REWARD_BONUS_WEIGHT=345;\
CURIOSITY_PRESSURE=346;SABOTEUR_INTENSITY_MULT=347;\
REWARD_WEIGHT_FLOOR=348;CURIOSITY_BOUND=349;\
VAL_SHARPE_DELTA_EMA=350;VAL_SHARPE_VAR_EMA=351;\
REWARD_COMPONENT_MAG_RATIO_BASE=352;\
SABOTEUR_ENGAGEMENT_RATE=358;PNL_REWARD_MAGNITUDE_EMA=359;\
POPART_COMPONENT_MAG_EMA=360;\
REWARD_COMPONENT_VAR_EMA_BASE=361;\
TARGET_DIR_ACC=372;AUX_DIR_ACC_SHORT_EMA=373;AUX_DIR_ACC_LONG_EMA=374;\
AUX_DIR_PREDICTION=375;DIR_SKILL_BONUS_ALPHA=376;DIR_SKILL_BONUS_BETA=377;\
LUCK_WIN_DISCOUNT=378;SKILL_BONUS_CAP_RATIO=379;\
HOLD_COST=380;HOLD_RATE_TARGET=381;HOLD_RATE_OBSERVED_EMA=382;\
ISV_TOTAL_DIM=383;\
PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\
PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\
PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\
PARAM_GAMMA_H_S2=11;PARAM_BETA_H_S2=12;\
PARAM_W_V1=13;PARAM_B_V1=14;PARAM_W_V2=15;PARAM_B_V2=16;\
PARAM_W_B0FC=17;PARAM_B_B0FC=18;PARAM_W_B0OUT=19;PARAM_B_B0OUT=20;\
PARAM_W_B1FC=21;PARAM_B_B1FC=22;PARAM_W_B1OUT=23;PARAM_B_B1OUT=24;\
PARAM_W_B2FC=25;PARAM_B_B2FC=26;PARAM_W_B2OUT=27;PARAM_B_B2OUT=28;\
PARAM_W_B3FC=29;PARAM_B_B3FC=30;PARAM_W_B3OUT=31;PARAM_B_B3OUT=32;\
PARAM_W_BN=33;PARAM_B_BN=34;\
PARAM_W_VSN1_0=35;PARAM_W_VSN2_0=36;\
PARAM_W_VSN1_1=37;PARAM_W_VSN2_1=38;\
PARAM_W_VSN1_2=39;PARAM_W_VSN2_2=40;\
PARAM_W_VSN1_3=41;PARAM_W_VSN2_3=42;\
PARAM_W_GATE_0=43;PARAM_B_GATE_0=44;\
PARAM_W_GATE_1=45;PARAM_B_GATE_1=46;\
PARAM_W_GATE_2=47;PARAM_B_GATE_2=48;\
PARAM_W_GATE_3=49;PARAM_B_GATE_3=50;\
PARAM_KAN_COEFF_0=51;PARAM_KAN_RESID_0=52;\
PARAM_KAN_COEFF_1=53;PARAM_KAN_RESID_1=54;\
PARAM_KAN_COEFF_2=55;PARAM_KAN_RESID_2=56;\
PARAM_KAN_COEFF_3=57;PARAM_KAN_RESID_3=58;\
PARAM_W_REGIME=59;PARAM_B_REGIME=60;\
PARAM_SPACING_RAW_0=61;PARAM_SPACING_RAW_1=62;\
PARAM_SPACING_RAW_2=63;PARAM_SPACING_RAW_3=64;\
PARAM_W_V1_5BAR=65;PARAM_B_V1_5BAR=66;\
PARAM_W_V2_5BAR=67;PARAM_B_V2_5BAR=68;\
PARAM_W_V1_20BAR=69;PARAM_B_V1_20BAR=70;\
PARAM_W_V2_20BAR=71;PARAM_B_V2_20BAR=72;\
PARAM_W_RISK_FC=73;PARAM_B_RISK_FC=74;\
PARAM_W_RISK_OUT=75;PARAM_B_RISK_OUT=76;\
PARAM_W_ISV_FC1=77;PARAM_B_ISV_FC1=78;\
PARAM_W_ISV_FC2=79;PARAM_B_ISV_FC2=80;\
PARAM_W_ISV_GATE=81;PARAM_B_ISV_GATE=82;\
PARAM_W_ISV_GAMMA=83;PARAM_B_ISV_GAMMA=84;\
PARAM_W_CONF_FC=85;PARAM_B_CONF_FC=86;\
PARAM_W_FEATURE_GATE=87;PARAM_B_FEATURE_GATE=88;\
PARAM_W_TEMPORAL_ROUTE=89;PARAM_B_TEMPORAL_ROUTE=90;\
PARAM_W_PLAN_FC=91;PARAM_B_PLAN_FC=92;\
PARAM_W_PLAN_OUT=93;PARAM_B_PLAN_OUT=94;\
PARAM_VSN_W1_G0=95;PARAM_VSN_B1_G0=96;PARAM_VSN_W2_G0=97;PARAM_VSN_B2_G0=98;\
PARAM_VSN_W1_G1=99;PARAM_VSN_B1_G1=100;PARAM_VSN_W2_G1=101;PARAM_VSN_B2_G1=102;\
PARAM_VSN_W1_G2=103;PARAM_VSN_B1_G2=104;PARAM_VSN_W2_G2=105;PARAM_VSN_B2_G2=106;\
PARAM_VSN_W1_G3=107;PARAM_VSN_B1_G3=108;PARAM_VSN_W2_G3=109;PARAM_VSN_B2_G3=110;\
PARAM_VSN_W1_G4=111;PARAM_VSN_B1_G4=112;PARAM_VSN_W2_G4=113;PARAM_VSN_B2_G4=114;\
PARAM_VSN_W1_G5=115;PARAM_VSN_B1_G5=116;PARAM_VSN_W2_G5=117;PARAM_VSN_B2_G5=118;\
PARAM_AUX_NB_W1=119;PARAM_AUX_NB_B1=120;PARAM_AUX_NB_W2=121;PARAM_AUX_NB_B2=122;\
PARAM_AUX_RG_W1=123;PARAM_AUX_RG_B1=124;PARAM_AUX_RG_W2=125;PARAM_AUX_RG_B2=126;\
PARAM_MOE_GATE_W1=127;PARAM_MOE_GATE_B1=128;PARAM_MOE_GATE_W2=129;PARAM_MOE_GATE_B2=130;\
PARAM_MOE_EXPERT_0_W1=131;PARAM_MOE_EXPERT_0_B1=132;PARAM_MOE_EXPERT_0_W2=133;PARAM_MOE_EXPERT_0_B2=134;\
PARAM_MOE_EXPERT_1_W1=135;PARAM_MOE_EXPERT_1_B1=136;PARAM_MOE_EXPERT_1_W2=137;PARAM_MOE_EXPERT_1_B2=138;\
PARAM_MOE_EXPERT_2_W1=139;PARAM_MOE_EXPERT_2_B1=140;PARAM_MOE_EXPERT_2_W2=141;PARAM_MOE_EXPERT_2_B2=142;\
PARAM_MOE_EXPERT_3_W1=143;PARAM_MOE_EXPERT_3_B1=144;PARAM_MOE_EXPERT_3_W2=145;PARAM_MOE_EXPERT_3_B2=146;\
PARAM_MOE_EXPERT_4_W1=147;PARAM_MOE_EXPERT_4_B1=148;PARAM_MOE_EXPERT_4_W2=149;PARAM_MOE_EXPERT_4_B2=150;\
PARAM_MOE_EXPERT_5_W1=151;PARAM_MOE_EXPERT_5_B1=152;PARAM_MOE_EXPERT_5_W2=153;PARAM_MOE_EXPERT_5_B2=154;\
PARAM_MOE_EXPERT_6_W1=155;PARAM_MOE_EXPERT_6_B1=156;PARAM_MOE_EXPERT_6_W2=157;PARAM_MOE_EXPERT_6_B2=158;\
PARAM_MOE_EXPERT_7_W1=159;PARAM_MOE_EXPERT_7_B1=160;PARAM_MOE_EXPERT_7_W2=161;PARAM_MOE_EXPERT_7_B2=162;\
PARAM_TOTAL_TENSORS=163";
#[test]
fn fingerprint_bumped_from_pre_b1_1a() {
use ml::cuda_pipeline::gpu_dqn_trainer::LAYOUT_FINGERPRINT_CURRENT;
let pre_b1_1a_fp = fnv1a_64(PRE_B1_1A_SEED);
assert_ne!(
LAYOUT_FINGERPRINT_CURRENT, pre_b1_1a_fp,
"B1.1a fingerprint must differ from pre-B1.1a value (the W2/B2 \
rename to _K2 must change the seed hash). \
current = {:#018x}, pre-B1.1a = {:#018x}",
LAYOUT_FINGERPRINT_CURRENT, pre_b1_1a_fp
);
}
// ─── Test 18: HEALTH_DIAG snapshot size unchanged at 149*4 ────────────────
/// B1.1a does NOT touch the HEALTH_DIAG snap-words layout — the aux
/// block remained at 3 words after B1.0 retired the `aux_label_scale`
/// entry (B1.0 dropped the size from 150 → 149). B1.1a's struct flips
/// (aux_nb_label_buf f32→i32, K=1→2 buffer growth, dir_acc_buf 3→6)
/// are all GPU-only: HEALTH_DIAG-snap remains 149 floats × 4 bytes =
/// 596 bytes. This test guards against accidental snap-layout
/// regressions during the cascade.
#[test]
fn health_diag_snap_size_stable_at_149_floats() {
use ml::cuda_pipeline::health_diag::HealthDiagSnapshot;
let n = std::mem::size_of::<HealthDiagSnapshot>();
assert_eq!(
n, 149 * 4,
"expected HealthDiagSnapshot to be 149*4 = 596 bytes (B1.0 stable, \
B1.1a unchanged), got {n}"
);
}

331
crates/ml/tests/sp13_phase0_oracle_tests.rs
View File

@@ -5,15 +5,18 @@
//! Validates the single-block tree-reduce kernel introduced by P0a.T2:
//! `aux_dir_acc_reduce_kernel.cu`
//!
//! Six oracle cases drive synthetic `(aux_pred, next_bar_label)` pairs
//! through the production kernel and verify the GPU-computed
//! `(dir_acc, pos_pred_frac, pos_label_frac)` triple against analytically-
//! known expected values. Per `feedback_no_cpu_test_fallbacks.md`: GPU
//! oracle only, no CPU reference impl. Per
//! SP13 B1.1a (2026-05-05): the kernel ABI flipped — reads
//! `softmax [B, K=2]` + i32 labels (was regression-mode `aux_pred f32 +
//! label_f32`) and writes 6 floats (was 3): tail of the output gained
//! `n_down`, `n_up`, `n_skip` per-class counts for HEALTH_DIAG. The
//! tests below drive synthetic `(softmax, i32 labels)` pairs through
//! the production kernel and verify the 6-element output against
//! analytically-known expected values. Per `feedback_no_cpu_test_fallbacks.md`:
//! GPU oracle only, no CPU reference impl. Per
//! `feedback_no_htod_htoh_only_mapped_pinned.md`: every CPU↔GPU buffer
//! is a `MappedF32Buffer` (cuMemHostAlloc with DEVICEMAP|PORTABLE);
//! zero `htod_copy`, zero `dtoh_sync_copy`. Tests are NOT exempt from
//! this rule.
//! is a `MappedF32Buffer` / `MappedI32Buffer` (cuMemHostAlloc with
//! DEVICEMAP|PORTABLE); zero `htod_copy`, zero `dtoh_sync_copy`. Tests
//! are NOT exempt from this rule.
//!
//! All tests are `#[ignore = "requires GPU"]`-gated to match every other
//! GPU oracle test in this crate (sp4/sp5/sp11/sp12). Run on a GPU host:
@@ -22,16 +25,18 @@
//! cargo test -p ml --test sp13_phase0_oracle_tests --features cuda \
//! -- --ignored --nocapture
//!
//! Sign convention under test (kernel side: `(p > 0.0f) ? 1 : 0`):
//! - `p > 0` → predicts class 1 (positive)
//! - `p <= 0` → predicts class 0 (negative) — incl. `0.0` and `-0.0`
//! - `label == 0.0` → no-signal bar, excluded from denominator entirely
//! Argmax convention (kernel side: `(softmax[1] > softmax[0]) ? 1 : 0`):
//! - `softmax[1] > softmax[0]` → predicts class 1 (up)
//! - `softmax[1] <= softmax[0]` → predicts class 0 (down) — incl. ties
//! - `label == -1` → mask/skip bar, excluded from denominator entirely;
//! contributes to `n_skip` only.
//!
//! Sentinel:
//! - When the entire batch is invalid (empty / all-zero labels),
//! - When the entire batch is masked (empty / all-mask labels),
//! `(dir_acc, pos_pred_frac, pos_label_frac) = (0.5, 0.5, 0.5)`.
//! Matches `DIR_ACC_EMA_SENTINEL` in `sp13_isv_slots.rs` per
//! `pearl_first_observation_bootstrap.md`.
//! `pearl_first_observation_bootstrap.md`. The trailing
//! `(n_down, n_up, n_skip)` triple emits actual counts regardless.
#![cfg(feature = "cuda")]
@@ -72,97 +77,132 @@ fn load_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
/// 3-element output through `dev_ptr` with `__threadfence_system()`; the
/// test reads through `read_all()`'s volatile read after stream sync.
///
/// Block dim 256, shared mem `4 × 256 × sizeof(int) = 4 KB` — matches
/// the production launcher's config exactly.
/// Block dim 256, shared mem `6 × 256 × sizeof(int) = 6 KB` — matches
/// the production launcher's config exactly (SP13 B1.1a grew shmem
/// 4 → 6 arrays for n_down/n_up/n_skip).
///
/// Returns the full 6-element output:
/// `[dir_acc, pos_pred_frac, pos_label_frac, n_down, n_up, n_skip]`
fn run_dir_acc(
stream: &Arc<CudaStream>,
f: &CudaFunction,
pred: &[f32],
label: &[f32],
) -> [f32; 3] {
softmax: &[f32],
labels: &[i32],
k: i32,
) -> [f32; 6] {
let n = labels.len();
assert_eq!(
pred.len(),
label.len(),
"pred and label must have matching length; got {} vs {}",
pred.len(),
label.len()
softmax.len(),
n * k as usize,
"softmax must be [B, K]; got {} for B={} K={}",
softmax.len(),
n,
k
);
let n = pred.len();
// Empty-batch case: even if `n == 0` the kernel still needs valid
// device pointers for the input args (CUDA dereferences them only
// inside the strided loop, which doesn't execute, but the launch
// builder still passes the pointers). Allocate 1-element placeholders
// so the dev_ptr is non-zero.
let alloc_len = n.max(1);
let alloc_len_smx = (n * k as usize).max(1);
let alloc_len_lbl = n.max(1);
// Safety: CUDA context active on this thread (resolved via the
// stream's context above). All CPU↔GPU buffers are mapped-pinned.
let pred_buf = unsafe { MappedF32Buffer::new(alloc_len) }
.expect("alloc pred buffer (mapped-pinned)");
let label_buf = unsafe { MappedF32Buffer::new(alloc_len) }
.expect("alloc label buffer (mapped-pinned)");
let softmax_buf = unsafe { MappedF32Buffer::new(alloc_len_smx) }
.expect("alloc softmax buffer (mapped-pinned)");
let labels_buf = unsafe { MappedI32Buffer::new(alloc_len_lbl) }
.expect("alloc labels buffer (mapped-pinned)");
if n > 0 {
pred_buf.write_from_slice(pred);
label_buf.write_from_slice(label);
softmax_buf.write_from_slice(softmax);
labels_buf.write_from_slice(labels);
}
let out_buf = unsafe { MappedF32Buffer::new(3) }
.expect("alloc out buffer (3 f32 mapped-pinned)");
let out_buf = unsafe { MappedF32Buffer::new(6) }
.expect("alloc out buffer (6 f32 mapped-pinned)");
let pred_dev = pred_buf.dev_ptr;
let label_dev = label_buf.dev_ptr;
let out_dev = out_buf.dev_ptr;
let softmax_dev = softmax_buf.dev_ptr;
let labels_dev = labels_buf.dev_ptr;
let out_dev = out_buf.dev_ptr;
let batch_size_i32: i32 = n as i32;
let bdim: u32 = 256;
unsafe {
stream
.launch_builder(f)
.arg(&pred_dev)
.arg(&label_dev)
.arg(&softmax_dev)
.arg(&labels_dev)
.arg(&batch_size_i32)
.arg(&k)
.arg(&out_dev)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (bdim, 1, 1),
shared_mem_bytes: 4 * bdim * std::mem::size_of::<i32>() as u32,
shared_mem_bytes: 6 * bdim * std::mem::size_of::<i32>() as u32,
})
.expect("launch aux_dir_acc_reduce_kernel");
}
stream.synchronize().expect("sync after aux_dir_acc_reduce launch");
let out = out_buf.read_all();
[out[0], out[1], out[2]]
[out[0], out[1], out[2], out[3], out[4], out[5]]
}
// ─── Six P0a.T2 oracle cases ─────────────────────────────────────────────
/// Build a one-hot `[B, K=2]` softmax tile from per-bar argmax bits.
/// `pred_class[b] == 1` produces `[0.0, 1.0]` in row b (predicts up);
/// `pred_class[b] == 0` produces `[1.0, 0.0]` (predicts down).
/// One-hot tiles deliberately exercise the kernel's argmax branch
/// without softmax noise (the kernel only compares the two probs;
/// any strict ordering of `[p0, p1]` that matches the desired argmax
/// is sufficient).
fn one_hot_softmax(pred_class: &[i32]) -> Vec<f32> {
let mut out = Vec::with_capacity(pred_class.len() * 2);
for &c in pred_class {
if c == 1 {
out.push(0.0);
out.push(1.0);
} else {
out.push(1.0);
out.push(0.0);
}
}
out
}
/// All four bars predict the correct sign → dir_acc = 1.0.
// ─── Six P0a.T2 oracle cases (SP13 B1.1a-rewritten) ─────────────────────
/// All four bars predict the correct class → dir_acc = 1.0.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_all_correct_returns_one() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let label = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
// pred [1, 1, 0, 0], label [1, 1, 0, 0] → all correct.
let softmax = one_hot_softmax(&[1, 1, 0, 0]);
let labels = vec![1_i32, 1, 0, 0];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
assert!(
(out[0] - 1.0).abs() < 1e-6,
"expected dir_acc = 1.0, got {}",
out[0]
);
// Tail counts: 2 down + 2 up + 0 skip.
assert!((out[3] - 2.0).abs() < 1e-6, "n_down expected 2, got {}", out[3]);
assert!((out[4] - 2.0).abs() < 1e-6, "n_up expected 2, got {}", out[4]);
assert!((out[5] - 0.0).abs() < 1e-6, "n_skip expected 0, got {}", out[5]);
}
/// All four bars predict the wrong sign → dir_acc = 0.0.
/// All four bars predict the wrong class → dir_acc = 0.0.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_all_wrong_returns_zero() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let label = vec![-1.0_f32, -1.0, 1.0, 1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
// pred [1, 1, 0, 0], label [0, 0, 1, 1] → all wrong.
let softmax = one_hot_softmax(&[1, 1, 0, 0]);
let labels = vec![0_i32, 0, 1, 1];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
assert!(
out[0].abs() < 1e-6,
"expected dir_acc = 0.0, got {}",
@@ -172,17 +212,17 @@ fn dir_acc_all_wrong_returns_zero() {
/// Half right + asymmetric prediction/label distributions.
///
/// pred = [+, +, +, +] → 4 positive predictions / 4 valid bars = 1.0
/// label = [+, +, -, -] → 2 positive labels / 4 valid bars = 0.5
/// pred = [1, 1, 1, 1] → 4 up predictions / 4 valid bars = 1.0
/// label = [1, 1, 0, 0] → 2 up labels / 4 valid bars = 0.5
/// dir_acc = 2/4 = 0.5 (bars 0+1 correct, bars 2+3 wrong).
#[test]
#[ignore = "requires GPU"]
fn dir_acc_half_correct() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, 1.0, 1.0];
let label = vec![ 1.0_f32, 1.0, -1.0, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
let softmax = one_hot_softmax(&[1, 1, 1, 1]);
let labels = vec![1_i32, 1, 0, 0];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
assert!(
(out[0] - 0.5).abs() < 1e-6,
"expected dir_acc = 0.5, got {}",
@@ -190,75 +230,79 @@ fn dir_acc_half_correct() {
);
assert!(
(out[1] - 1.0).abs() < 1e-6,
"expected pos_pred_frac = 1.0 (4/4 positive predictions), got {}",
"expected pos_pred_frac = 1.0 (4/4 up predictions), got {}",
out[1]
);
assert!(
(out[2] - 0.5).abs() < 1e-6,
"expected pos_label_frac = 0.5 (2/4 positive labels), got {}",
"expected pos_label_frac = 0.5 (2/4 up labels), got {}",
out[2]
);
}
/// Bars whose label is exactly 0.0 are excluded from the denominator
/// entirely.
/// Bars whose label is exactly -1 are excluded from the denominator
/// entirely (B1.1a mask sentinel).
///
/// pred = [+, +, -]
/// label = [+, 0, -] ← bar 1 is no-signal and skipped
/// pred = [1, 1, 0]
/// label = [1, -1, 0] ← bar 1 is masked and skipped
/// valid bars = 2 (indexes 0 and 2). Both correct → dir_acc = 2/2 = 1.0.
/// If the kernel mistakenly counted bar 1 as `label_pos = 0` (matching
/// `pred_pos = 1`?) it would be wrong; if it counted bar 1 as valid with
/// `label_pos = 0` and `pred_pos = 1`, dir_acc would drop to 2/3 ≈ 0.667.
/// Either failure mode is caught by the strict 1.0 expectation.
/// n_skip should equal 1 (the mask bar contributes to n_skip only).
#[test]
#[ignore = "requires GPU"]
fn dir_acc_skips_zero_labels() {
fn dir_acc_skips_mask_labels() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 1.0_f32, 1.0, -1.0];
let label = vec![ 1.0_f32, 0.0, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
let softmax = one_hot_softmax(&[1, 1, 0]);
let labels = vec![1_i32, -1, 0];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
assert!(
(out[0] - 1.0).abs() < 1e-6,
"expected dir_acc = 1.0 (bar 1 with label=0 excluded), got {}",
"expected dir_acc = 1.0 (bar 1 with label=-1 excluded), got {}",
out[0]
);
assert!((out[5] - 1.0).abs() < 1e-6, "n_skip expected 1, got {}", out[5]);
assert!((out[3] - 1.0).abs() < 1e-6, "n_down expected 1, got {}", out[3]);
assert!((out[4] - 1.0).abs() < 1e-6, "n_up expected 1, got {}", out[4]);
}
/// Sign convention: `(p > 0.0f) ? 1 : 0` treats both `+0.0` and `-0.0`
/// as the negative class.
/// Argmax convention: `(softmax[1] > softmax[0]) ? 1 : 0` treats ties
/// (both equal) as predicting class 0 (down). The B1.1a softmax is
/// produced by a stable softmax over logits — equal logits produce
/// equal probabilities, which the kernel must round to class 0.
///
/// pred = [0.0, -0.0] ← both predict class 0 (negative)
/// label = [-1.0, -1.0] ← both negative
/// softmax = [[0.5, 0.5], [0.5, 0.5]] ← both predict class 0 (tie → 0)
/// label = [0, 0] ← both down
/// dir_acc = 2/2 = 1.0 (both correct).
#[test]
#[ignore = "requires GPU"]
fn dir_acc_zero_pred_is_negative() {
fn dir_acc_softmax_tie_is_class_zero() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred = vec![ 0.0_f32, -0.0];
let label = vec![-1.0_f32, -1.0];
let out = run_dir_acc(&stream, &f, &pred, &label);
let softmax = vec![0.5_f32, 0.5, 0.5, 0.5];
let labels = vec![0_i32, 0];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
assert!(
(out[0] - 1.0).abs() < 1e-6,
"expected dir_acc = 1.0 (zero pred → class 0 → matches negative label), got {}",
"expected dir_acc = 1.0 (tied softmax → class 0 → matches down label), got {}",
out[0]
);
}
/// Empty batch (or all-zero-label batch) returns the random-baseline
/// sentinel 0.5 across all three outputs. Matches
/// Empty batch (or all-mask-label batch) returns the random-baseline
/// sentinel 0.5 across the first three outputs. Matches
/// `DIR_ACC_EMA_SENTINEL` in `sp13_isv_slots.rs` per
/// `pearl_first_observation_bootstrap.md` — the downstream EMA's first-
/// observation replacement consumes the sentinel naturally.
/// observation replacement consumes the sentinel naturally. The trailing
/// `(n_down, n_up, n_skip)` triple emits 0/0/0 for empty (no rows
/// counted), matching the actual count semantics.
#[test]
#[ignore = "requires GPU"]
fn dir_acc_empty_batch_returns_sentinel() {
let stream = make_test_stream();
let f = load_kernel(&stream);
let pred: Vec<f32> = vec![];
let label: Vec<f32> = vec![];
let out = run_dir_acc(&stream, &f, &pred, &label);
let softmax: Vec<f32> = vec![];
let labels: Vec<i32> = vec![];
let out = run_dir_acc(&stream, &f, &softmax, &labels, 2);
assert!(
(out[0] - 0.5).abs() < 1e-6,
"expected dir_acc sentinel = 0.5 on empty batch, got {}",
@@ -274,6 +318,10 @@ fn dir_acc_empty_batch_returns_sentinel() {
"expected pos_label_frac sentinel = 0.5 on empty batch, got {}",
out[2]
);
// n_down/n_up/n_skip = 0/0/0 on truly empty batch.
assert!((out[3]).abs() < 1e-6, "n_down expected 0 on empty, got {}", out[3]);
assert!((out[4]).abs() < 1e-6, "n_up expected 0 on empty, got {}", out[4]);
assert!((out[5]).abs() < 1e-6, "n_skip expected 0 on empty, got {}", out[5]);
}
// ─── SP13 v3 P0a.T3 Hold-rate observer oracle cases ─────────────────────
@@ -606,15 +654,19 @@ fn fixed_alpha_ema_sentinel_bootstrap() {
);
}
// ─── SP13 P0a.T4 aux-pred → ISV[375] tanh oracle cases ──────────────────
// ─── SP13 P0a.T4 aux-pred → ISV[375] oracle cases (B1.1a-rewritten) ────
//
// Validates the SP13 aux-head per-bar prediction → ISV[AUX_DIR_PREDICTION_INDEX=375]
// tanh-bounded scalar producer kernel introduced by P0a.T4:
// bounded scalar producer kernel introduced by P0a.T4 and rewritten by
// SP13 B1.1a (2026-05-05) to read the K=2 softmax tile directly:
// `aux_pred_to_isv_tanh_kernel.cu`
//
// Drives synthetic per-bar `aux_pred [B]` tiles through the production
// kernel and verifies the GPU-computed `mean(tanh(aux_pred[i]))` ∈ [-1, +1]
// against analytically-known expected values.
// Drives synthetic `softmax [B, 2]` tiles through the production kernel
// and verifies the GPU-computed `mean(softmax[:, 1] - softmax[:, 0])`
// ∈ [-1, +1] against analytically-known expected values. The kernel
// filename + ISV slot name are preserved for layout stability across
// the B1.1a rewrite (the "tanh" in the name is historical — the runtime
// squash is retired since softmax provides the bound structurally).
const SP13_AUX_PRED_TO_ISV_TANH_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/aux_pred_to_isv_tanh_kernel.cubin"));
@@ -629,37 +681,39 @@ fn load_aux_pred_to_isv_tanh(stream: &Arc<CudaStream>) -> CudaFunction {
.expect("load aux_pred_to_isv_tanh_kernel function")
}
/// Drive one launch of the aux-pred → ISV[375] tanh producer with the
/// given per-bar prediction slice. Returns ISV slot 0 post-launch.
/// Drive one launch of the aux-softmax → ISV[375] producer with the
/// given `softmax [B, 2]` tile. Returns ISV slot 0 post-launch.
fn run_aux_pred_to_isv_tanh(
stream: &Arc<CudaStream>,
f: &CudaFunction,
aux_pred: &[f32],
softmax: &[f32],
k: i32,
) -> f32 {
let n = aux_pred.len();
let alloc_len = n.max(1);
let n = if k > 0 { softmax.len() / k as usize } else { 0 };
let alloc_len = (n * k as usize).max(1);
let pred_buf = unsafe { MappedF32Buffer::new(alloc_len) }
.expect("alloc pred buffer (mapped-pinned)");
if n > 0 {
pred_buf.write_from_slice(aux_pred);
let softmax_buf = unsafe { MappedF32Buffer::new(alloc_len) }
.expect("alloc softmax buffer (mapped-pinned)");
if !softmax.is_empty() {
softmax_buf.write_from_slice(softmax);
}
// 1-elem ISV scratch — kernel only writes slot 0.
let isv_buf = unsafe { MappedF32Buffer::new(1) }
.expect("alloc isv buffer (mapped-pinned)");
let pred_dev = pred_buf.dev_ptr;
let softmax_dev = softmax_buf.dev_ptr;
let isv_dev = isv_buf.dev_ptr;
let batch_i32: i32 = n as i32;
let batch_i32: i32 = n as i32;
let isv_off_i32: i32 = 0;
let bdim: u32 = 256;
unsafe {
stream
.launch_builder(f)
.arg(&pred_dev)
.arg(&softmax_dev)
.arg(&batch_i32)
.arg(&k)
.arg(&isv_off_i32)
.arg(&isv_dev)
.launch(LaunchConfig {
@@ -674,58 +728,65 @@ fn run_aux_pred_to_isv_tanh(
isv_buf.read_all()[0]
}
/// Symmetric batch around zero → mean(tanh) ≈ 0.0.
/// `tanh(±x)` is odd, so an exactly-symmetric input sums to zero.
/// Balanced batch (every row is `[0.5, 0.5]`) → mean diff = 0.0.
/// Mirrors the regression-mode "symmetric input → mean=0" test, but
/// for the B1.1a softmax tile the equivalent is "every row is split
/// 50/50 between classes".
#[test]
#[ignore = "requires GPU"]
fn aux_pred_tanh_balanced_softmax_returns_zero() {
let stream = make_test_stream();
let f = load_aux_pred_to_isv_tanh(&stream);
// 4 rows of [0.5, 0.5] → diff = 0 per row → mean = 0.
let softmax = vec![0.5_f32; 8];
let out = run_aux_pred_to_isv_tanh(&stream, &f, &softmax, 2);
assert!(
out.abs() < 1e-6,
"expected mean(softmax[1]-softmax[0]) = 0.0 on balanced batch, got {}",
out
);
}
/// Symmetric (mixed) batch: 2 rows fully up + 2 rows fully down →
/// mean = (1.0 + 1.0 + (-1.0) + (-1.0)) / 4 = 0.0.
#[test]
#[ignore = "requires GPU"]
fn aux_pred_tanh_symmetric_batch_returns_zero() {
let stream = make_test_stream();
let f = load_aux_pred_to_isv_tanh(&stream);
let pred = vec![1.0_f32, -1.0, 0.5, -0.5, 2.0, -2.0];
let out = run_aux_pred_to_isv_tanh(&stream, &f, &pred);
let softmax = one_hot_softmax(&[1, 1, 0, 0]);
let out = run_aux_pred_to_isv_tanh(&stream, &f, &softmax, 2);
assert!(
out.abs() < 1e-6,
"expected mean(tanh) = 0.0 on symmetric batch, got {}",
"expected mean = 0.0 on symmetric batch (2 up + 2 down), got {}",
out
);
}
/// All-zero batch → mean(tanh(0)) = 0.0.
#[test]
#[ignore = "requires GPU"]
fn aux_pred_tanh_all_zero_returns_zero() {
let stream = make_test_stream();
let f = load_aux_pred_to_isv_tanh(&stream);
let pred = vec![0.0_f32; 16];
let out = run_aux_pred_to_isv_tanh(&stream, &f, &pred);
assert!(
out.abs() < 1e-6,
"expected mean(tanh(0)) = 0.0, got {}",
out
);
}
/// Bounded by [-1, +1] regardless of input magnitude. tanh saturates
/// at ±1 for large |x|, so mean(tanh(large)) ≈ ±1; large negative
/// inputs land near -1 (well within the [-1, +1] bound).
/// Saturated batch (every row strongly predicts up) →
/// mean ≈ +1.0; the structural bound `[-1, +1]` holds without any
/// runtime tanh squash. Mirror test for fully-down case.
#[test]
#[ignore = "requires GPU"]
fn aux_pred_tanh_large_magnitude_saturates_in_bounds() {
let stream = make_test_stream();
let f = load_aux_pred_to_isv_tanh(&stream);
let pred = vec![10.0_f32; 32]; // tanh(10) ≈ 1.0 to ~5 nines
let out = run_aux_pred_to_isv_tanh(&stream, &f, &pred);
// All up: every row is [0.0, 1.0] → diff = +1.0 per row → mean = +1.0.
let softmax_up = one_hot_softmax(&[1; 32]);
let out_up = run_aux_pred_to_isv_tanh(&stream, &f, &softmax_up, 2);
assert!(
(out - 1.0).abs() < 1e-4 && out <= 1.0,
"expected mean(tanh(10)) ≈ 1.0 (saturated), got {}",
out
(out_up - 1.0).abs() < 1e-6 && out_up <= 1.0,
"expected mean = +1.0 on all-up batch, got {}",
out_up
);
let pred_neg = vec![-10.0_f32; 32];
let out_neg = run_aux_pred_to_isv_tanh(&stream, &f, &pred_neg);
// All down: every row is [1.0, 0.0] → diff = -1.0 per row → mean = -1.0.
let softmax_dn = one_hot_softmax(&[0; 32]);
let out_dn = run_aux_pred_to_isv_tanh(&stream, &f, &softmax_dn, 2);
assert!(
(out_neg + 1.0).abs() < 1e-4 && out_neg >= -1.0,
"expected mean(tanh(-10)) ≈ -1.0 (saturated), got {}",
out_neg
(out_dn + 1.0).abs() < 1e-6 && out_dn >= -1.0,
"expected mean = -1.0 on all-down batch, got {}",
out_dn
);
}

87
docs/dqn-wire-up-audit.md
View File

@@ -5977,3 +5977,90 @@ Build: `cargo check --workspace --tests` clean. `snapshot_size_is_stable` passes
### 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.
## SP13 Layer B — Commit B1.1a: K=1→2 + softmax CE kernel rewrites + struct flips (2026-05-05)
**Why split B1.1a from B1.1b**: B1.1 itself was decomposed after six full-B1 dispatches confirmed agent-session capacity is the bottleneck, not cascade understanding. B1.1 is now split into:
- **B1.1a** (this commit): kernel ABI flips + struct field flips + tests (kernel-side cascade, contract-consistent atomic).
- **B1.1b** (next commit): producer kernel for `aux_sign_labels` (the i32 -1/0/1 label producer driven by the price trajectory), replay direct-path 8th gather, experience-collector hoist, end-to-end round-trip tests, Smoke A.
B1.1a is contract-consistent atomic kernel-side: every consumer of K-flip / softmax tile / CE / i32 label dtype migrates atomically within this commit per `feedback_no_partial_refactor`. Labels stay zero-init (`alloc_zeros<i32>`) until B1.1b lands the producer — every sample receives label 0 ("down"), so the model converges on "predict class 0 everywhere". This degraded state is intentional and known; the cascade is internally consistent and the tests validate the kernels in isolation, NOT the trained behavior.
**No L40S smoke between B1.1a and B1.1b**. Local unit tests on the RTX 3050 Ti validate B1.1a kernel correctness (CE loss, CE backward, dir_acc argmax, isv_tanh structural bound, fingerprint bump). L40S smoke runs after B1.1b lands the producer kernel + remaining tests.
### Four contracts (atomic in this commit)
| Contract | Producer | Consumers |
|---|---|---|
| **K_NB flip 1 → 2** | `AUX_NEXT_BAR_K = 2` (was 1) | `compute_param_sizes` ([121]/[122]); fingerprint seed (rename `_W2/_B2 → _W2_K2/_B2_K2`); forward kernel; backward kernel; loss reduce; saxpy-spec table; partial-buf allocs (nb_w2 `[B,H]→[B,K,H]`, nb_b2 `[B]→[B,K]`); `max_aux_tensor_len`; pred_buf rename `aux_nb_pred_buf → aux_nb_logits_buf` |
| **Softmax tile output** | `aux_next_bar_forward` writes `[B, K]` softmax via in-kernel stable-softmax | `aux_next_bar_loss_reduce` (CE numerator); `aux_next_bar_backward` (`d_logits = softmax - one_hot`); `aux_dir_acc_reduce_kernel` (argmax); `aux_pred_to_isv_tanh_kernel` (mean diff). NEW field `aux_nb_softmax_buf [B, K]`. |
| **MSE → CE loss flip** | `aux_next_bar_loss_reduce` reads softmax + i32 labels, masks `-1`, divides by `B_valid` (mean-over-valid-rows), writes loss + `B_valid` | `aux_heads_loss_ema_update` (unchanged consumer); `aux_next_bar_backward` reads `B_valid` so loss + grad share `1/B_valid` divisor. NEW field `aux_nb_valid_count_buf [1]`. |
| **i32 label dtype** | `aux_nb_label_buf: CudaSlice<i32>` (was `f32`); `alloc_zeros::<i32>` (no producer wired in B1.1a — B1.1b lands it) | `aux_next_bar_loss_reduce`, `aux_next_bar_backward`, `aux_dir_acc_reduce_kernel` all read i32 labels with `-1` mask sentinel. The strided_gather of `next_states[:, 0]` retired entirely. |
### Audit-verified call site cardinality
| Site | Count | Action |
|---|---|---|
| `aux_next_bar_forward` callers | 1 | gain `K`, `softmax_out` args |
| `aux_next_bar_loss_reduce` callers | 1 | gain `K`, swap `pred → softmax`, `label_f32 → labels_i32`, gain `valid_count_out` |
| `aux_next_bar_backward` callers | 1 | gain `K`, swap `pred → softmax`, `label_f32 → labels_i32`, gain `valid_count` |
| `aux_dir_acc_reduce_kernel` callers | 1 launcher + 1 orchestrator | swap `aux_pred → softmax`, swap `f32 label → i32 labels`, gain `K`, output 3 → 6 floats |
| `aux_pred_to_isv_tanh_kernel` callers | 1 launcher + 1 orchestrator | swap `aux_pred → softmax`, gain `K`, drop runtime tanh (structural bound via softmax) |
| `aux_dir_acc_buf` size | 1 alloc | grew 3 → 6 floats (added `n_down/n_up/n_skip`) |
| `aux_nb_pred_buf` rename | 4 (struct field, 2 fwd reads, 2 bwd reads) | renamed to `aux_nb_logits_buf` per `feedback_no_legacy_aliases` |
| `aux_nb_label_buf` dtype | 1 (struct field + alloc) | `f32 → i32`; `alloc_zeros<i32>`; producer (strided_gather) deleted |
| New buffers | 2 | `aux_nb_softmax_buf [B, K]`, `aux_nb_valid_count_buf [1]` |
| Saxpy-spec partial sizes | 2 entries (121, 122) | `(aux_h)→(aux_knb*aux_h)`, `(1)→(aux_knb)` |
| Layout fingerprint seed | 1 line (line 2144) | rename `_W2/_B2 → _W2_K2/_B2_K2`; fingerprint hash bumps |
| HEALTH_DIAG | 1 new emit line | `aux_b1_diag` reads `aux_dir_acc_buf [3..6]` for `n_down/n_up/n_skip + mask_frac` |
| HEALTH_DIAG snap-words layout | 0 | unchanged — B1.1a is GPU-side; HEALTH_DIAG snap stays at 149 floats × 4 bytes = 596 bytes |
| Existing dir_acc + isv_tanh oracle tests | 9 (6 dir_acc + 3 isv_tanh) | rewritten in-place to new ABI per `feedback_no_partial_refactor` |
| New B1.1a oracle tests | 11 | 5 CE loss/backward + 2 dir_acc + 2 isv_tanh + 2 layout regression |
### Hard rules upheld
- `feedback_no_partial_refactor`: every consumer of K-flip / softmax tile / CE / i32 label migrates within this commit — kernel signatures, orchestrator launchers, struct fields, training-loop diag, existing oracle tests, all atomic
- `feedback_no_atomicadd`: block tree-reduce only; CE loss reduce uses 2 parallel partial-reduction strips (mirrors regime CE shape); CE backward emits per-sample partials → existing `aux_param_grad_reduce` collapses
- `feedback_cpu_is_read_only`: `aux_nb_label_buf` is GPU-resident `CudaSlice<i32>`, populated by the (forthcoming B1.1b) producer kernel; HEALTH_DIAG aux_b1_diag reads via mapped-pinned `aux_dir_acc_buf` (no DtoH)
- `feedback_no_stubs`: every new buffer + kernel arg is wired through to a real consumer in this commit; the CE forward / loss / backward chain executes end-to-end against the placeholder labels (model converges on "predict 0 everywhere" — degraded but real, not a stub)
- `feedback_no_todo_fixme` / `feedback_no_hiding` / `feedback_no_quickfixes` / `feedback_no_feature_flags`: no markers, no `_` underscores, no `enable_*` toggles
- `feedback_no_legacy_aliases`: `aux_nb_pred_buf` renamed to `aux_nb_logits_buf` directly at every call site (no shim); `PARAM_AUX_NB_W2/_B2` renamed to `_W2_K2/_B2_K2` in seed (no `_DEPRECATED` alias); existing dir_acc/isv_tanh oracle tests rewritten in-place rather than carrying old-ABI shadows
- `feedback_no_cpu_test_fallbacks`: every B1.1a oracle test is GPU-side, `#[ignore = "requires GPU"]`-gated (the 2 layout regression tests are CPU-only because they exercise `pub const LAYOUT_FINGERPRINT_CURRENT` and `size_of::<HealthDiagSnapshot>()`, both pure-Rust)
- `feedback_no_htod_htoh_only_mapped_pinned`: every CPU↔GPU buffer in tests + production is `MappedF32Buffer` / `MappedI32Buffer`; zero `htod_copy` / `dtoh_sync_copy`
- `feedback_isv_for_adaptive_bounds`: no hardcoded thresholds added (the CE numerical floor 1e-30 prevents `-log(0) = +inf` and is a numerical-stability floor, not a tuned threshold)
- `feedback_trust_code_not_docs`: 8/8 Phase 0 anchors verified against current code at HEAD `75e94858c` before editing
- `pearl_bounded_modifier_outputs_require_structural_activation`: softmax IS the structural activation; the `aux_pred_to_isv_tanh_kernel` runtime tanh squash is removed entirely (the [-1, +1] bound for ISV[375] now comes from `softmax[1] - softmax[0]` where each component is in [0, 1] and they sum to 1)
- `pearl_build_rs_rerun_if_env_changed`: N/A for B1.1a (no new cubin file added; existing cubins recompile via existing `aux_heads_kernel.cu` / `aux_dir_acc_reduce_kernel.cu` / `aux_pred_to_isv_tanh_kernel.cu` glob patterns in `build.rs`)
### Files (atomic B1.1a commit)
8 source + 1 test + 1 audit doc:
- `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` (~+150 LOC) — `aux_next_bar_forward` gains `K` + `logits_out` + `softmax_out` (in-kernel stable softmax); `aux_next_bar_loss_reduce` ABI flipped (softmax + i32 labels, mean-over-valid CE, `valid_count_out`); `aux_next_bar_backward` ABI flipped (softmax + i32 labels + `valid_count`, K-fanout d_logits, masked rows zero)
- `crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu` (~+50 LOC) — read softmax + i32 labels, argmax over K, output grew 3 → 6 floats (added `n_down/n_up/n_skip`); shmem 4 → 6 int arrays
- `crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu` (~−15 LOC) — read softmax tile, compute `mean(softmax[:, 1] - softmax[:, 0])`; tanh transcend retired (structural bound via softmax components)
- `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs` (~+50 LOC net) — `AUX_NEXT_BAR_K: 1 → 2`; `forward_next_bar` gains `K`, `logits_out`, `softmax_out` args; `next_bar_loss_reduce` gains `K`, `valid_count_out`; `backward_next_bar` gains `softmax_in`, `labels_i32_in`, `valid_count_in`, `K`
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (~+80 LOC net) — `compute_param_sizes` ([121]/[122]); fingerprint seed rename `_W2/_B2 → _W2_K2/_B2_K2`; struct fields (`aux_nb_logits_buf`, `aux_nb_softmax_buf`, `aux_nb_label_buf: CudaSlice<i32>`, `aux_nb_valid_count_buf`); `aux_dir_acc_buf [6]`; `aux_partial_nb_w2 [B,K,H]`, `aux_partial_nb_b2 [B,K]`; `max_aux_tensor_len` extended; saxpy spec table; orchestrator launchers gain `K` arg; strided_gather block deleted entirely
- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (~+30 LOC) — `aux_b1_diag` HEALTH_DIAG line reads `aux_dir_acc_buf [3..6]`; doc comment update for the per-step aux dir-metrics block
- `crates/ml/tests/sp13_phase0_oracle_tests.rs` (~+50 LOC net) — `run_dir_acc` scaffold rewritten for softmax + i32 + 6-float output; 6 dir_acc tests rewritten; 3 isv_tanh tests rewritten; helper `one_hot_softmax`
- `crates/ml/tests/sp13_layer_b_oracle_tests.rs` (NEW, ~700 LOC) — 11 B1.1a tests: 5 CE (single-row + batch-mixed + all-skip-NaN + backward single-row + backward batch-mixed), 2 dir_acc (handcrafted argmax + all-skip NaN-safe), 2 isv_tanh (bounded fuzz + mean handcrafted), 2 layout regression (fingerprint bump + HEALTH_DIAG snap stable)
- `docs/dqn-wire-up-audit.md` — this section
### Build + test verification
- `SQLX_OFFLINE=true cargo check --workspace` — clean (only pre-existing warnings)
- `SQLX_OFFLINE=true cargo check --workspace --tests` — clean
- `SQLX_OFFLINE=true cargo test -p ml --lib --no-run` — compiles
- `SQLX_OFFLINE=true cargo test -p ml-dqn --lib --no-run` — compiles
- `SQLX_OFFLINE=true cargo test -p ml --lib snapshot_size_is_stable` — passes (149*4=596 bytes; B1.0-stable, B1.1a unchanged)
- `SQLX_OFFLINE=true cargo test -p ml --test sp13_layer_b_oracle_tests fingerprint_bumped_from_pre_b1_1a` — passes (proves the W2/B2 rename flipped the seed hash)
- `SQLX_OFFLINE=true cargo test -p ml --test sp13_layer_b_oracle_tests health_diag_snap_size_stable_at_149_floats` — passes
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp13_layer_b_oracle_tests --features cuda -- --ignored --nocapture` — 9 GPU tests pass on local RTX 3050 Ti (B=1/4 — small enough that 4GB VRAM has headroom)
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp13_phase0_oracle_tests --features cuda -- --ignored --nocapture` — 14 tests pass (6 dir_acc + 3 isv_tanh rewrites + 5 unchanged)
Pre-existing test failures (14) on `cargo test -p ml --lib` are environmental (OFI test data missing for `test_feature_vector_to_state`, etc.) and reproduce identically at HEAD `75e94858c` (B1.0) — verified via `git stash` round-trip. Not B1.1a-introduced.
### Next: B1.1b — producer kernel + replay direct path + experience collector hoist + remaining tests + Smoke A
B1.1b covers: the producer kernel that fills `aux_sign_labels` with real {-1, 0, 1} from the 30-bar price trajectory (replaces the placeholder zero-init); replay direct-path 8th gather (the i32 column carried by B0's plumbing finally feeds into `aux_nb_label_buf`); experience collector hoist (so the producer kernel runs at experience-write time, not training-step time); the remaining 6 producer + 1 round-trip oracle tests (6 producer-kernel correctness + 1 end-to-end forward → loss → backward → param-update round trip on a small fold); Smoke A on L40S validates the live training behavior with real labels.