Add a missing memory fence to a SIMD test. (#432)

JuliaGPU · Sep 26, 2024 · 5b983a5 · 5b983a5
1 parent 5628d43
commit 5b983a5
Showing 1 changed file with 4 additions and 1 deletion.
diff --git a/test/device/intrinsics.jl b/test/device/intrinsics.jl
@@ -373,10 +373,13 @@ end
 
             tg_a = MtlThreadGroupArray(T, (8, 8))
             tg_a[pos.x, pos.y] = a[pos.x, pos.y]
-            sg_a = simdgroup_load(tg_a)
+            threadgroup_barrier(Metal.MemoryFlagThreadGroup)
 
+            sg_a = simdgroup_load(tg_a)
             tg_b = MtlThreadGroupArray(T, (8, 8))
             simdgroup_store(sg_a, tg_b)
+
+            threadgroup_barrier(Metal.MemoryFlagThreadGroup)
             b[pos.x, pos.y] = tg_b[pos.x, pos.y]
 
             return