Add MemoryFlagDevice to KA.jl's synchronization primitive. (#609)

KernelAbstractions (influenced by CUDA) states

> After a  statement all read and writes to global and local memory
> from each thread in the workgroup are visible in from all other threads in the
> workgroup.

Author: maleadt

src/MetalKernels.jl

@@ -189,7 +189,7 @@ end
 ## other
 
 @device_override @inline function KA.__synchronize()
-    threadgroup_barrier(Metal.MemoryFlagThreadGroup)
+    threadgroup_barrier(Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup)
 end
 
 @device_override @inline function KA.__print(args...)

src/device/intrinsics/synchronization.jl

@@ -34,21 +34,21 @@ end
 
 
 """
-    threadgroup_barrier(flag::MemoryFlags=MemoryFlagNone)
+    threadgroup_barrier(flag=MemoryFlagNone)
 
 Synchronize all threads in a threadgroup.
 
 Possible flags that affect the memory synchronization behavior are found in [`MemoryFlags`](@ref)
 """
-@inline threadgroup_barrier(flag::MemoryFlags=MemoryFlagNone) =
+@inline threadgroup_barrier(flag=MemoryFlagNone) =
     ccall("extern air.wg.barrier", llvmcall, Cvoid, (Cuint, Cuint, ), flag, UInt32(1))
 
 """
-    simdgroup_barrier(flag::MemoryFlags=MemoryFlagNone)
+    simdgroup_barrier(flag=MemoryFlagNone)
 
 Synchronize all threads in a SIMD-group.
 
 Possible flags that affect the memory synchronization behavior are found in [`MemoryFlags`](@ref)
 """
-@inline simdgroup_barrier(flag::MemoryFlags=MemoryFlagNone) =
+@inline simdgroup_barrier(flag=MemoryFlagNone) =
     ccall("extern air.simdgroup.barrier", llvmcall, Cvoid, (Cuint, Cuint, ), flag, UInt32(1))