SPIRV-Cross/reference/shaders-msl-no-opt/packing/load-store-col-rows.comp

#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& v_21)
{
    float2 u = v_21.a[0].xy;
    float2 v = v_21.a[1].xy;
    u += v;
    (device float2&)v_21.a[0] = u;
    (device float2&)v_21.a[1] = v;
}

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

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

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

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