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;
         }