From 7d10ea8b3e84ca1127dfaf67eafcb5adb1f45153 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 5 May 2026 14:09:48 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp13):=20B1.1a=20=E2=80=94=20K=3D1?= =?UTF-8?q?=E2=86=922=20+=20softmax=20CE=20kernel=20rewrites=20+=20struct?= =?UTF-8?q?=20flips?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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` is `alloc_zeros` 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; 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) --- .../aux_dir_acc_reduce_kernel.cu | 179 ++-- .../ml/src/cuda_pipeline/aux_heads_kernel.cu | 302 +++++-- .../aux_pred_to_isv_tanh_kernel.cu | 87 +- crates/ml/src/cuda_pipeline/gpu_aux_heads.rs | 107 ++- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 437 ++++++---- .../src/trainers/dqn/trainer/training_loop.rs | 58 +- crates/ml/tests/sp13_layer_b_oracle_tests.rs | 821 ++++++++++++++++++ crates/ml/tests/sp13_phase0_oracle_tests.rs | 331 ++++--- docs/dqn-wire-up-audit.md | 87 ++ 9 files changed, 1865 insertions(+), 544 deletions(-) create mode 100644 crates/ml/tests/sp13_layer_b_oracle_tests.rs diff --git a/crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu b/crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu index f21bf1bda..2d7eb8662 100644 --- a/crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu @@ -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 #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(); } } diff --git a/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu b/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu index 7569aabbd..bd3c73cdf 100644 --- a/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu @@ -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) { diff --git a/crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu b/crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu index 5bb2762dd..b395a7b31 100644 --- a/crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_pred_to_isv_tanh_kernel.cu @@ -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 #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(); diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs b/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs index 6e11ba68a..050b7b07b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_heads.rs @@ -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::() as u32; + let k_i32 = k as i32; + let smem_bytes = (AUX_HIDDEN_DIM as u32 + k as u32) * std::mem::size_of::() 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, - 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::() 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::() 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::() 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::() 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) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index dd90faa4e..408d937c5 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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, - /// `[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, - /// `[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, - /// `[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, + /// `[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, + /// `[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, + /// `[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, + /// `[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, /// `[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, @@ -4353,11 +4379,14 @@ pub struct GpuDqnTrainer { aux_partial_nb_w1: CudaSlice, /// `[B, AUX_HIDDEN_DIM]` per-sample db1 partial — next-bar head. aux_partial_nb_b1: CudaSlice, - /// `[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, - /// `[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, /// `[B, AUX_HIDDEN_DIM, SH2]` per-sample dW1 partial — regime head. aux_partial_rg_w1: CudaSlice, @@ -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::() as u32, + shared_mem_bytes: 6 * bdim * std::mem::size_of::() 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, ¶m_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, ¶m_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::(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::(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, diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 185a94cd9..9f508ae3b 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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. diff --git a/crates/ml/tests/sp13_layer_b_oracle_tests.rs b/crates/ml/tests/sp13_layer_b_oracle_tests.rs new file mode 100644 index 000000000..ff0f2f6e0 --- /dev/null +++ b/crates/ml/tests/sp13_layer_b_oracle_tests.rs @@ -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 { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() +} + +fn load_loss_reduce(stream: &Arc) -> 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) -> 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) -> 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) -> 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, + 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::() 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, + f: &CudaFunction, + softmax: &[f32], + labels: &[i32], + valid_count: f32, + b: usize, + sh2: usize, + k: i32, +) -> Vec { + 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::() 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, + 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::() 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, + 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::() 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; 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::(); + assert_eq!( + n, 149 * 4, + "expected HealthDiagSnapshot to be 149*4 = 596 bytes (B1.0 stable, \ + B1.1a unchanged), got {n}" + ); +} diff --git a/crates/ml/tests/sp13_phase0_oracle_tests.rs b/crates/ml/tests/sp13_phase0_oracle_tests.rs index 7073dfa59..4362d69e8 100644 --- a/crates/ml/tests/sp13_phase0_oracle_tests.rs +++ b/crates/ml/tests/sp13_phase0_oracle_tests.rs @@ -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) -> 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, 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::() as u32, + shared_mem_bytes: 6 * bdim * std::mem::size_of::() 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 { + 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 = vec![]; - let label: Vec = vec![]; - let out = run_dir_acc(&stream, &f, &pred, &label); + let softmax: Vec = vec![]; + let labels: Vec = 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) -> 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, 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 ); } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 5e79ae8b5..751b8499e 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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`) 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` (was `f32`); `alloc_zeros::` (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`; 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`, 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::()`, 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`, `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.