| #include <metal_stdlib> | 
 |  | 
 | using namespace metal; | 
 | struct S { | 
 |   int x; | 
 |   atomic_uint a; | 
 |   uint y; | 
 | }; | 
 |  | 
 | void tint_zero_workgroup_memory(uint local_idx, threadgroup S* const tint_symbol) { | 
 |   if ((local_idx < 1u)) { | 
 |     (*(tint_symbol)).x = 0; | 
 |     atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed); | 
 |     (*(tint_symbol)).y = 0u; | 
 |   } | 
 |   threadgroup_barrier(mem_flags::mem_threadgroup); | 
 | } | 
 |  | 
 | void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol_1) { | 
 |   tint_zero_workgroup_memory(local_invocation_index, tint_symbol_1); | 
 |   atomic_store_explicit(&((*(tint_symbol_1)).a), 1u, memory_order_relaxed); | 
 | } | 
 |  | 
 | kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) { | 
 |   threadgroup S tint_symbol_2; | 
 |   compute_main_inner(local_invocation_index, &(tint_symbol_2)); | 
 |   return; | 
 | } | 
 |  |