From 1b5afccec7524c8806597d0837a9aa125995e90b Mon Sep 17 00:00:00 2001 From: ZYP Date: Sat, 12 Sep 2026 23:49:30 +0800 Subject: [PATCH] Barrier before reusing GQA reduction scratch After the max reduction, reductions[0] must be read by every thread before the array is overwritten for the sum pass. Missing the barrier made causal GQA softmax nondeterministic at long sequences (#52). --- h3_shaders.metal | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/h3_shaders.metal b/h3_shaders.metal index 36316f52..cab6c750 100644 --- a/h3_shaders.metal +++ b/h3_shaders.metal @@ -4003,6 +4003,10 @@ kernel void h3_gqa_causal_bf16( threadgroup_barrier(mem_flags::mem_threadgroup); } float maximum = reductions[0]; + /* WAR hazard: reductions[] is reused for the sum reduction. Without + * this barrier one thread can overwrite reductions[0] before another + * has read the max (antirez/h3.c#52). */ + threadgroup_barrier(mem_flags::mem_threadgroup); float local_sum = 0.0f; for (uint key_row = tid; key_row < key_count; key_row += threads) { float probability = exp(scores[key_row] - maximum);