#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

typedef packed_float3 packed_float2x3[2];
typedef packed_float3 packed_rm_float3x2[2];

struct SSBO1
{
    float2x4 a;
    float2x4 a2;
};

struct SSBO2
{
    packed_float2x3 b;
    packed_rm_float3x2 b2;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);

static inline __attribute__((always_inline))
void load_store_column(device SSBO1& _21)
{
    float2 u = _21.a[0].xy;
    float2 v = _21.a[1].xy;
    u += v;
    (device float2&)_21.a[0] = u;
    (device float2&)_21.a[1] = v;
}

static inline __attribute__((always_inline))
void load_store_row(device SSBO1& _21)
{
    float2 u = float2(_21.a2[0][0], _21.a2[1][0]);
    float2 v = float2(_21.a2[0][1], _21.a2[1][1]);
    u += v;
    ((device float*)&_21.a2[0])[0] = u.x;
    ((device float*)&_21.a2[1])[0] = u.y;
    ((device float*)&_21.a2[0])[1] = v.x;
    ((device float*)&_21.a2[1])[1] = v.y;
}

static inline __attribute__((always_inline))
void load_store_packed_column(device SSBO2& _58)
{
    float3 u = float3(_58.b[0]);
    float3 v = float3(_58.b[1]);
    u += v;
    _58.b[0] = u;
    _58.b[1] = v;
}

static inline __attribute__((always_inline))
void load_store_packed_row(device SSBO2& _58)
{
    float2 u = float2(_58.b2[0][0], _58.b2[1][0]);
    float2 v = float2(_58.b2[0][1], _58.b2[1][1]);
    u += v;
    ((device float*)&_58.b2[0])[0] = u.x;
    ((device float*)&_58.b2[1])[0] = u.y;
    ((device float*)&_58.b2[0])[1] = v.x;
    ((device float*)&_58.b2[1])[1] = v.y;
}

kernel void main0(device SSBO1& _21 [[buffer(0)]], device SSBO2& _58 [[buffer(1)]])
{
    load_store_column(_21);
    load_store_row(_21);
    load_store_packed_column(_58);
    load_store_packed_row(_58);
}