91 lines
2.0 KiB
Plaintext
91 lines
2.0 KiB
Plaintext
#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();
|
|
}
|
|
|