A potential nondeterministic issue

[yfyang86]

I am using this project to enhance other VLM (A Qwen2 variant) and found a potential nondeterministic issue, i.e. the causal-GQA kernel proved nondeterministic under repeated identical dispatches (29/30 runs differed at sequence 662). Fix this will greatly reduce the performance. In my case 47 ms vs 5.5 ms at sequence length of 2,013.

If deterministic is required in some case, some one may find this short patch useful, but it needs a re-design to get better performance.

--- a/h3_shaders.metal	2026-08-20 17:01:10
+++ b/h3_shaders.metal	2026-08-21 10:54:11
@@ -4003,6 +4003,11 @@
         threadgroup_barrier(mem_flags::mem_threadgroup);
     }
     float maximum = reductions[0];
+    /* fix: WAR hazard —  reuses reductions[]; without this
+     * barrier one thread can overwrite reductions[0] before another thread
+     * has read the max, making the softmax nondeterministic at long
+     * sequences (observed: 29/30 identical-input runs differed at seq 662). */
+    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);