diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index 311fc5204..ac7713079 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -411,6 +411,9 @@ extern "C" __global__ void c51_loss_batched( if (tid == 0) { float clamped_ce = fminf(avg_ce, MAX_PER_SAMPLE_CE); float weighted_loss = clamped_ce * is_weight; + /* NaN guard: zero out poisoned samples (bf16 logit overflow from GemmEx) */ + if (!fast_isfinite(weighted_loss)) weighted_loss = 0.0f; + if (!fast_isfinite(clamped_ce)) clamped_ce = 0.0f; per_sample_loss[sample_id] = bf16(weighted_loss); td_errors[sample_id] = bf16(clamped_ce); atomicAdd(total_loss, weighted_loss / (float)batch_size); diff --git a/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu b/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu index d6cc3191c..460439c9f 100644 --- a/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu @@ -355,6 +355,13 @@ extern "C" __global__ void mse_loss_batched( if (tid == 0) { float weighted_loss = avg_mse * is_weight; + /* NaN guard: if any per-sample computation produced NaN, zero it out. + * This prevents one poisoned sample from killing the entire batch. + * The NaN source is bf16 logit inputs from cuBLAS GemmEx — when the + * f32 accumulated dot product exceeds bf16 max, the bf16 C-matrix write + * produces Inf/NaN which propagates through softmax expected-Q. */ + if (!fast_isfinite(weighted_loss)) weighted_loss = 0.0f; + if (!fast_isfinite(avg_td)) avg_td = 0.0f; per_sample_loss[sample_id] = bf16(weighted_loss); td_errors[sample_id] = bf16(avg_td); atomicAdd(total_loss, weighted_loss / (float)batch_size);