A potential nondeterministic issue

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.

```diff
--- 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);
```

Start in h3_shaders.metal around the causal-GQA kernel near the shown reduction code. Reproduce repeated identical dispatches at the reported sequence lengths, then evaluate the supplied synchronization change and its performance cost. Done means identical inputs produce consistent results without an unacceptable regression from the reported 47 ms versus 5.5 ms behavior.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.