89 lines
3.8 KiB
Plaintext
89 lines
3.8 KiB
Plaintext
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
typedef packed_float3 packed_float4x3[4];
|
|
|
|
struct SSBOCol
|
|
{
|
|
packed_float4x3 col_major0;
|
|
packed_float4x3 col_major1;
|
|
};
|
|
|
|
struct SSBORow
|
|
{
|
|
float3x4 row_major0;
|
|
float3x4 row_major1;
|
|
};
|
|
|
|
void load_store_to_variable_col_major(device SSBOCol& v_29)
|
|
{
|
|
float4x3 loaded = float4x3(float3(v_29.col_major0[0]), float3(v_29.col_major0[1]), float3(v_29.col_major0[2]), float3(v_29.col_major0[3]));
|
|
v_29.col_major1[0] = loaded[0];
|
|
v_29.col_major1[1] = loaded[1];
|
|
v_29.col_major1[2] = loaded[2];
|
|
v_29.col_major1[3] = loaded[3];
|
|
}
|
|
|
|
void load_store_to_variable_row_major(device SSBORow& v_41)
|
|
{
|
|
float4x3 loaded = transpose(v_41.row_major0);
|
|
v_41.row_major0 = transpose(loaded);
|
|
}
|
|
|
|
void copy_col_major_to_col_major(device SSBOCol& v_29)
|
|
{
|
|
v_29.col_major0[0] = float4x3(float3(v_29.col_major1[0]), float3(v_29.col_major1[1]), float3(v_29.col_major1[2]), float3(v_29.col_major1[3]))[0];
|
|
v_29.col_major0[1] = float4x3(float3(v_29.col_major1[0]), float3(v_29.col_major1[1]), float3(v_29.col_major1[2]), float3(v_29.col_major1[3]))[1];
|
|
v_29.col_major0[2] = float4x3(float3(v_29.col_major1[0]), float3(v_29.col_major1[1]), float3(v_29.col_major1[2]), float3(v_29.col_major1[3]))[2];
|
|
v_29.col_major0[3] = float4x3(float3(v_29.col_major1[0]), float3(v_29.col_major1[1]), float3(v_29.col_major1[2]), float3(v_29.col_major1[3]))[3];
|
|
}
|
|
|
|
void copy_col_major_to_row_major(device SSBOCol& v_29, device SSBORow& v_41)
|
|
{
|
|
v_41.row_major0 = transpose(float4x3(float4(v_29.col_major0[0]), float4(v_29.col_major0[1]), float4(v_29.col_major0[2])));
|
|
}
|
|
|
|
void copy_row_major_to_col_major(device SSBOCol& v_29, device SSBORow& v_41)
|
|
{
|
|
v_29.col_major0[0] = float3(float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[0][0], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[1][0], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[2][0]));
|
|
v_29.col_major0[1] = float3(float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[0][1], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[1][1], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[2][1]));
|
|
v_29.col_major0[2] = float3(float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[0][2], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[1][2], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[2][2]));
|
|
v_29.col_major0[3] = float3(float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[0][3], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[1][3], float4x3(v_41.row_major0[0], v_41.row_major0[1], v_41.row_major0[2])[2][3]));
|
|
}
|
|
|
|
void copy_row_major_to_row_major(device SSBORow& v_41)
|
|
{
|
|
v_41.row_major0 = v_41.row_major1;
|
|
}
|
|
|
|
void copy_columns(device SSBOCol& v_29, device SSBORow& v_41)
|
|
{
|
|
v_29.col_major0[1] = float3(v_41.row_major0[0][1], v_41.row_major0[1][1], v_41.row_major0[2][1]);
|
|
v_41.row_major0[0][1] = v_29.col_major0[1][0];
|
|
v_41.row_major0[1][1] = v_29.col_major0[1][1];
|
|
v_41.row_major0[2][1] = v_29.col_major0[1][2];
|
|
}
|
|
|
|
void copy_elements(device SSBOCol& v_29, device SSBORow& v_41)
|
|
{
|
|
v_29.col_major0[0][1u] = v_41.row_major0[1u][0];
|
|
v_41.row_major0[1u][0] = v_29.col_major0[0][1u];
|
|
}
|
|
|
|
kernel void main0(device SSBOCol& v_29 [[buffer(0)]], device SSBORow& v_41 [[buffer(1)]])
|
|
{
|
|
load_store_to_variable_col_major(v_29);
|
|
load_store_to_variable_row_major(v_41);
|
|
copy_col_major_to_col_major(v_29);
|
|
copy_col_major_to_row_major(v_29, v_41);
|
|
copy_row_major_to_col_major(v_29, v_41);
|
|
copy_row_major_to_row_major(v_41);
|
|
copy_columns(v_29, v_41);
|
|
copy_elements(v_29, v_41);
|
|
}
|
|
|