SPIRV-Cross/reference/shaders-msl/comp/barriers.comp

#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);

static inline __attribute__((always_inline))
void barrier_shared()
{
    threadgroup_barrier(mem_flags::mem_threadgroup);
}

static inline __attribute__((always_inline))
void full_barrier()
{
    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
}

static inline __attribute__((always_inline))
void image_barrier()
{
    threadgroup_barrier(mem_flags::mem_texture);
}

static inline __attribute__((always_inline))
void buffer_barrier()
{
    threadgroup_barrier(mem_flags::mem_device);
}

static inline __attribute__((always_inline))
void group_barrier()
{
    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
}

static inline __attribute__((always_inline))
void barrier_shared_exec()
{
    threadgroup_barrier(mem_flags::mem_threadgroup);
}

static inline __attribute__((always_inline))
void full_barrier_exec()
{
    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
}

static inline __attribute__((always_inline))
void image_barrier_exec()
{
    threadgroup_barrier(mem_flags::mem_texture);
}

static inline __attribute__((always_inline))
void buffer_barrier_exec()
{
    threadgroup_barrier(mem_flags::mem_device);
}

static inline __attribute__((always_inline))
void group_barrier_exec()
{
    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
}

static inline __attribute__((always_inline))
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();
}