test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.msl

#include <metal_stdlib>
using namespace metal;

struct tint_module_vars_struct {
  const constant float3x2* u;
  threadgroup float3x2* w;
};

struct tint_symbol_1 {
  float3x2 tint_symbol;
};

void f_inner(uint tint_local_index, tint_module_vars_struct tint_module_vars) {
  if ((tint_local_index < 1u)) {
    (*tint_module_vars.w) = float3x2(float2(0.0f), float2(0.0f), float2(0.0f));
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  (*tint_module_vars.w) = (*tint_module_vars.u);
  (*tint_module_vars.w)[1u] = (*tint_module_vars.u)[0u];
  (*tint_module_vars.w)[1u] = (*tint_module_vars.u)[0u].yx;
  (*tint_module_vars.w)[0u].y = (*tint_module_vars.u)[1u].x;
}

kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v [[threadgroup(0)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
  f_inner(tint_local_index, tint_module_vars);
}