test/tint/bug/tint/1926.wgsl.expected.msl

#include <metal_stdlib>
using namespace metal;

struct tint_module_vars_struct {
  threadgroup uint* sh_atomic_failed;
  device uint* output;
};

struct tint_symbol_1 {
  uint tint_symbol;
};

void main_inner(uint3 global_id, uint3 local_id, uint tint_local_index, tint_module_vars_struct tint_module_vars) {
  if ((tint_local_index < 1u)) {
    (*tint_module_vars.sh_atomic_failed) = 0u;
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  uint const v = (*tint_module_vars.sh_atomic_failed);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  uint const failed = v;
  if ((local_id.x == 0u)) {
    (*tint_module_vars.output) = failed;
  }
}

[[max_total_threads_per_threadgroup(256)]]
kernel void v_1(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_1* v_2 [[threadgroup(0)]], device uint* output [[buffer(0)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sh_atomic_failed=(&(*v_2).tint_symbol), .output=output};
  main_inner(global_id, local_id, tint_local_index, tint_module_vars);
}