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

#include <metal_stdlib>
using namespace metal;

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

struct tint_symbol_1 {
  half4x2 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) = half4x2(half2(0.0h), half2(0.0h), half2(0.0h), half2(0.0h));
  }
  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 half4x2* 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);
}