1#pragma clang diagnostic ignored "-Wmissing-prototypes"
2
3#include <metal_stdlib>
4#include <simd/simd.h>
5
6using namespace metal;
7
8constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
9
10static inline __attribute__((always_inline))
11void barrier_shared()
12{
13    threadgroup_barrier(mem_flags::mem_threadgroup);
14}
15
16static inline __attribute__((always_inline))
17void full_barrier()
18{
19    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
20}
21
22static inline __attribute__((always_inline))
23void image_barrier()
24{
25    threadgroup_barrier(mem_flags::mem_texture);
26}
27
28static inline __attribute__((always_inline))
29void buffer_barrier()
30{
31    threadgroup_barrier(mem_flags::mem_device);
32}
33
34static inline __attribute__((always_inline))
35void group_barrier()
36{
37    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
38}
39
40static inline __attribute__((always_inline))
41void barrier_shared_exec()
42{
43    threadgroup_barrier(mem_flags::mem_threadgroup);
44}
45
46static inline __attribute__((always_inline))
47void full_barrier_exec()
48{
49    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
50}
51
52static inline __attribute__((always_inline))
53void image_barrier_exec()
54{
55    threadgroup_barrier(mem_flags::mem_texture);
56}
57
58static inline __attribute__((always_inline))
59void buffer_barrier_exec()
60{
61    threadgroup_barrier(mem_flags::mem_device);
62}
63
64static inline __attribute__((always_inline))
65void group_barrier_exec()
66{
67    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
68}
69
70static inline __attribute__((always_inline))
71void exec_barrier()
72{
73    threadgroup_barrier(mem_flags::mem_threadgroup);
74}
75
76kernel void main0()
77{
78    barrier_shared();
79    full_barrier();
80    image_barrier();
81    buffer_barrier();
82    group_barrier();
83    barrier_shared_exec();
84    full_barrier_exec();
85    image_barrier_exec();
86    buffer_barrier_exec();
87    group_barrier_exec();
88    exec_barrier();
89}
90
91