blob: a6e9dba7a433509eda98fab9c1a4192fa34888c8 [file] [log] [blame]
#include <metal_stdlib>
using namespace metal;
struct tint_symbol_2 {
uint tint_symbol_1;
};
struct tint_module_vars_struct {
threadgroup uint* sh_atomic_failed;
device uint* output;
};
kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], uint3 local_id [[thread_position_in_threadgroup]], uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]], device uint* output [[buffer(4)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sh_atomic_failed=(&(*v).tint_symbol_1), .output=output};
if ((tint_local_index == 0u)) {
(*tint_module_vars.sh_atomic_failed) = 0u;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_threadgroup);
uint const v_1 = (*tint_module_vars.sh_atomic_failed);
threadgroup_barrier(mem_flags::mem_threadgroup);
uint const failed = v_1;
if ((local_id[0u] == 0u)) {
(*tint_module_vars.output) = failed;
}
}