From 104fe81ca86502d9699b2a4f95ab11f2b3524921 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 29 May 2026 10:49:58 +0200 Subject: [PATCH] =?UTF-8?q?fix(cuda):=20VSN=20stride=20mismatch=20?= =?UTF-8?q?=E2=80=94=20read=20window=5Ftensor=20at=20stride=2056?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Root cause of the step-4 NaN that has blocked the Phase 2.1 smoke gate since session start. `window_tensor_d` is allocated [B, K, ENCODER_INPUT_DIM=56]: - Features [0..40) are per-snapshot market features (snap_feature_assemble_batched) - Features [40..56) are per-batch broadcast context (rl_encoder_context_broadcast writes via offset `idx * 56 + 40`) But `variable_selection_fwd` and `_bwd` read x with stride `VSN_FEATURE_DIM=40`. For row 0 the stride happens to align with the [B, K, 56] layout. For row 1 onwards the kernel reads MIXED data across row boundaries — broadcast context bleeds into VSN's "snap features" view, and snap features bleed across K boundaries. The bleed produces nonsense for most steps but is bounded enough to train through. Around step 4, accumulated trade_context magnitudes (time_in_trade, unrealized_R, position_lots) get large enough to overflow VSN's softmax in the bleed-affected rows → NaN cascades through the encoder forward → backward computes NaN gradients → AdamW applies them → step 5's forward graph sees NaN weights → G8 abort. Fix: introduce `VSN_X_ROW_STRIDE = 56` (= ENCODER_INPUT_DIM) and use it for all input reads of `x` in both forward and backward: - variable_selection_fwd:41 (x_row = x + row * 56) - variable_selection_bwd:173 (x_i read at stride 56) - variable_selection_bwd:204 (xj read at stride 56) Output strides (gates_out, y, grad_W, grad_b, grad_x) remain at VSN_FEATURE_DIM=40 because: - gates_out / y feed Mamba2 L1 which reads at in_dim=40 stride - grad_W / grad_b accumulate into [FEATURE_DIM, FEATURE_DIM] / [FEATURE_DIM] - grad_x is write-only (audit: `vsn_grad_x_d` has zero downstream consumers per grep of crates/ml-alpha/) Verification: - 10/10 runs clean at seed=16962, n-steps=10 (was 4-7/10 NaN before) - 1000-step smoke clean at seed=16962: l_q≤1.76, l_v≤1.18, wr climbing 0.36→0.55 (the surfer pattern emerging as expected per pearl_dd049d9a4_surfer_baseline_verified) - Phase 2.1 smoke gate (l_q<3.0, l_v<2.0, no NaN over 1000 steps) PASSED Per pearl_atomicadd_masks_v_instability commits 60e96bf55..4e629830b for the full diagnostic chain (49 nan_scan labels, 4 sub-agent dispatches, 13 commits localizing the bug from "somewhere in the trainer" down to this single line). Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/cuda/variable_selection.cu | 24 +++++++++++++++++++--- 1 file changed, 21 insertions(+), 3 deletions(-) diff --git a/crates/ml-alpha/cuda/variable_selection.cu b/crates/ml-alpha/cuda/variable_selection.cu index b86b69bfa..6c25bec86 100644 --- a/crates/ml-alpha/cuda/variable_selection.cu +++ b/crates/ml-alpha/cuda/variable_selection.cu @@ -26,6 +26,24 @@ #define VSN_FEATURE_DIM 40 #define VSN_BLOCK 64 // round up to warp-multiple; threads i >= FEATURE_DIM idle. +// 2026-05-29 stride-mismatch fix. +// VSN's input buffer (window_tensor_d) is allocated [B, K, ENCODER_INPUT_DIM=56] +// by perception.rs (snap features [0..40) + per-batch broadcast context +// [40..56) written by rl_encoder_context_broadcast). VSN only processes the +// first VSN_FEATURE_DIM=40 features per row (snap features), but the input +// rows are spaced 56 floats apart, not 40. The kernel originally indexed x +// with stride VSN_FEATURE_DIM=40 — correct ONLY for row 0; every subsequent +// row read mixed broadcast-context + snap features across the [B, K, 56] +// row boundaries. Symptoms: intermittent step-4 NaN as accumulating trade +// context magnitudes overflowed VSN's softmax via the bleed. +// +// Fix: use VSN_X_ROW_STRIDE=56 for reading x in both forward and backward. +// Output (gates, y) and gradient outputs (grad_W, grad_b, grad_x) remain at +// VSN_FEATURE_DIM=40 because the downstream consumers (Mamba2 L1 with +// in_dim=40) read at compact 40-stride. grad_x is unused downstream (see +// `vsn_grad_x_d` audit — write-only), so its stride doesn't matter. +#define VSN_X_ROW_STRIDE 56 // = ENCODER_INPUT_DIM in heads.rs / perception.rs + extern "C" __global__ void variable_selection_fwd( const float* __restrict__ W_vsn, // [FEATURE_DIM, FEATURE_DIM] const float* __restrict__ b_vsn, // [FEATURE_DIM] @@ -38,7 +56,7 @@ extern "C" __global__ void variable_selection_fwd( int tid = threadIdx.x; if (row >= n_rows) return; - const float* x_row = x + (long long)row * VSN_FEATURE_DIM; + const float* x_row = x + (long long)row * VSN_X_ROW_STRIDE; // Shared mem: gate_logit + max-reduce scratch + sum-reduce scratch. __shared__ float s_logit[VSN_FEATURE_DIM]; @@ -170,7 +188,7 @@ extern "C" __global__ void variable_selection_bwd( float dy_i = 0.0f; if (tid < VSN_FEATURE_DIM) { gates_i = gates[(long long)row * VSN_FEATURE_DIM + tid]; - x_i = x[(long long)row * VSN_FEATURE_DIM + tid]; + x_i = x[(long long)row * VSN_X_ROW_STRIDE + tid]; dy_i = grad_y[(long long)row * VSN_FEATURE_DIM + tid]; s_gates[tid] = gates_i; s_dgates[tid] = dy_i * x_i; @@ -201,7 +219,7 @@ extern "C" __global__ void variable_selection_bwd( const float dl_t = s_dlogit[tid]; #pragma unroll for (int j = 0; j < VSN_FEATURE_DIM; ++j) { - const float xj = x[(long long)row * VSN_FEATURE_DIM + j]; + const float xj = x[(long long)row * VSN_X_ROW_STRIDE + j]; grad_W_vsn_scratch[row_FF + (long long)tid * VSN_FEATURE_DIM + j] += dl_t * xj; }