test/tint/builtins/workgroupUniformLoad/structures.wgsl.expected.msl

#include <metal_stdlib>
using namespace metal;

struct Inner {
  bool b;
  int4 v;
  float3x3 m;
};

template<typename T, size_t N>
struct tint_array {
  const constant T& operator[](size_t i) const constant { return elements[i]; }
  device T& operator[](size_t i) device { return elements[i]; }
  const device T& operator[](size_t i) const device { return elements[i]; }
  thread T& operator[](size_t i) thread { return elements[i]; }
  const thread T& operator[](size_t i) const thread { return elements[i]; }
  threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
  const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
  T elements[N];
};

struct tint_packed_vec3_f32_array_element {
  packed_float3 packed;
};

struct Inner_packed_vec3 {
  bool b;
  int4 v;
  tint_array<tint_packed_vec3_f32_array_element, 3> m;
};

struct Outer {
  tint_array<Inner, 4> a;
};

struct Outer_packed_vec3 {
  tint_array<Inner_packed_vec3, 4> a;
};

struct tint_module_vars_struct {
  threadgroup Outer_packed_vec3* v;
};

struct tint_symbol_1 {
  Outer_packed_vec3 tint_symbol;
};

Inner tint_load_struct_packed_vec3_1(threadgroup Inner_packed_vec3* const from) {
  bool const v_1 = (*from).b;
  int4 const v_2 = (*from).v;
  tint_array<tint_packed_vec3_f32_array_element, 3> const v_3 = (*from).m;
  return Inner{.b=v_1, .v=v_2, .m=float3x3(float3(v_3[0u].packed), float3(v_3[1u].packed), float3(v_3[2u].packed))};
}

tint_array<Inner, 4> tint_load_array_packed_vec3(threadgroup tint_array<Inner_packed_vec3, 4>* const from) {
  Inner const v_4 = tint_load_struct_packed_vec3_1((&(*from)[0u]));
  Inner const v_5 = tint_load_struct_packed_vec3_1((&(*from)[1u]));
  Inner const v_6 = tint_load_struct_packed_vec3_1((&(*from)[2u]));
  return tint_array<Inner, 4>{v_4, v_5, v_6, tint_load_struct_packed_vec3_1((&(*from)[3u]))};
}

Outer tint_load_struct_packed_vec3(threadgroup Outer_packed_vec3* const from) {
  return Outer{.a=tint_load_array_packed_vec3((&(*from).a))};
}

Outer foo(tint_module_vars_struct tint_module_vars) {
  threadgroup_barrier(mem_flags::mem_threadgroup);
  Outer const v_7 = tint_load_struct_packed_vec3(tint_module_vars.v);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  return v_7;
}

void tint_store_array_packed_vec3(threadgroup Inner_packed_vec3* const to, Inner value) {
  (*to).b = value.b;
  (*to).v = value.v;
  (*to).m[0u].packed = packed_float3(value.m[0u]);
  (*to).m[1u].packed = packed_float3(value.m[1u]);
  (*to).m[2u].packed = packed_float3(value.m[2u]);
}

void main_inner(uint tint_local_index, tint_module_vars_struct tint_module_vars) {
  {
    uint v_8 = 0u;
    v_8 = tint_local_index;
    while(true) {
      uint const v_9 = v_8;
      if ((v_9 >= 4u)) {
        break;
      }
      tint_store_array_packed_vec3((&(*tint_module_vars.v).a[v_9]), Inner{});
      {
        v_8 = (v_9 + 1u);
      }
      continue;
    }
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  foo(tint_module_vars);
}

[[max_total_threads_per_threadgroup(1)]]
kernel void v_10(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v_11 [[threadgroup(0)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_11).tint_symbol)};
  main_inner(tint_local_index, tint_module_vars);
}