Add exhaustive barrier tests for MSL.

Seems incorrect, need to be fixed later.

reference/opt/shaders-msl/comp/barriers.comp

@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+
+kernel void main0()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    threadgroup_barrier(mem_flags::mem_texture);
+    threadgroup_barrier(mem_flags::mem_none);
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    threadgroup_barrier(mem_flags::mem_texture);
+    threadgroup_barrier(mem_flags::mem_none);
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+

reference/shaders-msl/comp/barriers.comp

@@ -0,0 +1,79 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+
+void barrier_shared()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+
+void full_barrier()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+
+void image_barrier()
+{
+    threadgroup_barrier(mem_flags::mem_texture);
+}
+
+void buffer_barrier()
+{
+    threadgroup_barrier(mem_flags::mem_none);
+}
+
+void group_barrier()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+
+void barrier_shared_exec()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+
+void full_barrier_exec()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+
+void image_barrier_exec()
+{
+    threadgroup_barrier(mem_flags::mem_texture);
+}
+
+void buffer_barrier_exec()
+{
+    threadgroup_barrier(mem_flags::mem_none);
+}
+
+void group_barrier_exec()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+
+void exec_barrier()
+{
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+}
+
+kernel void main0()
+{
+    barrier_shared();
+    full_barrier();
+    image_barrier();
+    buffer_barrier();
+    group_barrier();
+    barrier_shared_exec();
+    full_barrier_exec();
+    image_barrier_exec();
+    buffer_barrier_exec();
+    group_barrier_exec();
+    exec_barrier();
+}
+

shaders-msl/comp/barriers.comp

@@ -0,0 +1,79 @@
+#version 310 es
+layout(local_size_x = 4) in;
+
+void barrier_shared()
+{
+	memoryBarrierShared();
+}
+
+void full_barrier()
+{
+	memoryBarrier();
+}
+
+void image_barrier()
+{
+	memoryBarrierImage();
+}
+
+void buffer_barrier()
+{
+	memoryBarrierBuffer();
+}
+
+void group_barrier()
+{
+	groupMemoryBarrier();
+}
+
+void barrier_shared_exec()
+{
+	memoryBarrierShared();
+	barrier();
+}
+
+void full_barrier_exec()
+{
+	memoryBarrier();
+	barrier();
+}
+
+void image_barrier_exec()
+{
+	memoryBarrierImage();
+	barrier();
+}
+
+void buffer_barrier_exec()
+{
+	memoryBarrierBuffer();
+	barrier();
+}
+
+void group_barrier_exec()
+{
+	groupMemoryBarrier();
+	barrier();
+}
+
+void exec_barrier()
+{
+	barrier();
+}
+
+void main()
+{
+	barrier_shared();
+	full_barrier();
+	image_barrier();
+	buffer_barrier();
+	group_barrier();
+
+	barrier_shared_exec();
+	full_barrier_exec();
+	image_barrier_exec();
+	buffer_barrier_exec();
+	group_barrier_exec();
+
+	exec_barrier();
+}