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

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;
  float3 const v_4 = float3(v_3[0u].packed);
  float3 const v_5 = float3(v_3[1u].packed);
  return Inner{.b=v_1, .v=v_2, .m=float3x3(v_4, v_5, 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_6 = tint_load_struct_packed_vec3_1((&(*from)[0u]));
  Inner const v_7 = tint_load_struct_packed_vec3_1((&(*from)[1u]));
  Inner const v_8 = tint_load_struct_packed_vec3_1((&(*from)[2u]));
  return tint_array<Inner, 4>{v_6, v_7, v_8, 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_9 = tint_load_struct_packed_vec3(tint_module_vars.v);
  threadgroup_barrier(mem_flags::mem_threadgroup);
  return v_9;
}