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