| #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; |
| } |
| } |