test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl

#include <metal_stdlib>
using namespace metal;

struct UBO {
  /* 0x0000 */ int dynamic_idx;
};

struct Result {
  /* 0x0000 */ int out;
};

struct S {
  array<int, 64> data;
};

struct tint_module_vars_struct {
  const constant UBO* ubo;
  device Result* result;
  threadgroup S* s;
};

struct tint_symbol_1 {
  S tint_symbol;
};

void f_inner(uint tint_local_index, tint_module_vars_struct tint_module_vars) {
  {
    uint idx = tint_local_index;
    while(true) {
      uint const v = idx;
      if ((v >= 64u)) {
        break;
      }
      (*tint_module_vars.s).data[v] = 0;
      {
        idx = (idx + 1u);
      }
    }
  }
  (threadgroup_barrier(mem_flags::mem_threadgroup));
  (*tint_module_vars.s).data[min(uint((*tint_module_vars.ubo).dynamic_idx), 63u)] = 1;
  (*tint_module_vars.result).out = (*tint_module_vars.s).data[3];
}

[[max_total_threads_per_threadgroup(1)]]
kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant UBO* ubo [[buffer(0)]], device Result* result [[buffer(1)]], threadgroup tint_symbol_1* v_1 [[threadgroup(0)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.ubo=ubo, .result=result, .s=(&(*v_1).tint_symbol)};
  (f_inner(tint_local_index, tint_module_vars));
}