From ca7dfa9f08c9093f7e968db086c09b3fb12412af Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 17 Apr 2026 13:19:39 +0200 Subject: [PATCH] =?UTF-8?q?fix(critical):=20c51=5Floss=5Freduce=20missing?= =?UTF-8?q?=20threadIdx=20guard=20=E2=80=94=2032-way=20write=20race?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit SECOND race condition from block_dim (1,1,1)→(32,1,1) change: c51_loss_reduce had 32 threads all computing sum and writing to total_loss[0] simultaneously. On sm_90 (H100) → corrupted loss → NaN in C51 gradient → graph_mega replay hangs. This kernel runs INSIDE graph_mega (submit_forward_ops_main), called twice (MSE reduce + C51 reduce). Both corrupted. Also fixed stochastic_depth_rng in previous commit (same issue, runs in pre-replay). Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml/src/cuda_pipeline/c51_loss_kernel.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index bd2f33ab1..21a55e4a2 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -708,6 +708,7 @@ extern "C" __global__ void c51_loss_reduce( float* __restrict__ total_loss, /* [1] output */ int batch_size ) { + if (threadIdx.x != 0) return; /* single writer — total_loss[0] */ float sum = 0.0f; for (int i = 0; i < batch_size; i++) { sum += per_sample_loss[i];