diff --git a/crates/ml-alpha/cuda/compute_advantage_return.cu b/crates/ml-alpha/cuda/compute_advantage_return.cu index 45d016064..e850923ed 100644 --- a/crates/ml-alpha/cuda/compute_advantage_return.cu +++ b/crates/ml-alpha/cuda/compute_advantage_return.cu @@ -19,13 +19,26 @@ extern "C" __global__ void compute_advantage_return( const int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= b_size) return; - const float gamma = isv[RL_GAMMA_INDEX]; - const float r = rewards[b]; - const float done = dones[b]; - const float vt = v_t[b]; - const float vtp1 = v_tp1[b]; + const float gamma = isv[RL_GAMMA_INDEX]; + const float r = rewards[b]; + const float is_done = (dones[b] > 0.5f); + const float vt = v_t[b]; + const float vtp1 = v_tp1[b]; - const float ret = r + gamma * (1.0f - done) * vtp1; - returns[b] = ret; - advantages[b] = done * (ret - vt); + // 2026-05-29: branch-gate (not multiplication-gate) the V_tp1 term + // and the (ret - vt) advantage. Multiplication-gating fails on IEEE + // `0 * Inf = NaN`: if any batch's encoder produces a non-finite V + // prediction early in training (random init outliers), the prior + // `done * (ret - vt)` multiplication propagated NaN to ALL non-done + // batches' advantages, which then broke compute_advantage_rms (sum + // of A² → NaN) and PPO (A/RMS → NaN, ratio×A → NaN, l_pi → NaN). + // Branch-gating ensures non-finite vtp1/vt are NEVER mixed into a + // non-done batch's advantage. Validated by the smoke bisect at + // 10d4614fb (deterministic NaN at step 4 with l_v=6.329 stable — + // proving V REGRESSION is fine while V FORWARD has at least one + // non-finite output that the 0*Inf trick was promoting to NaN + // everywhere). Per `pearl_atomicadd_masks_v_instability`. + const float ret = is_done ? r : (r + gamma * vtp1); + returns[b] = ret; + advantages[b] = is_done ? (ret - vt) : 0.0f; }