1f6653ec07
If we have emitted block IO lowering at the end of vertex shader, we will end up using the wrong name. Forcing a v_ prefix does not solve any actual problems since the intentifier already has to be valid.
92 lines
2.8 KiB
Plaintext
92 lines
2.8 KiB
Plaintext
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
typedef packed_float3 packed_rm_float3x2[2];
|
|
|
|
struct SSBOCol
|
|
{
|
|
float3x2 col_major0;
|
|
float3x2 col_major1;
|
|
};
|
|
|
|
struct SSBORow
|
|
{
|
|
packed_rm_float3x2 row_major0;
|
|
packed_rm_float3x2 row_major1;
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
|
|
static inline __attribute__((always_inline))
|
|
void load_store_to_variable_col_major(device SSBOCol& _29)
|
|
{
|
|
float3x2 loaded = _29.col_major0;
|
|
_29.col_major1 = loaded;
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void load_store_to_variable_row_major(device SSBORow& _41)
|
|
{
|
|
float3x2 loaded = transpose(float2x3(float3(_41.row_major0[0]), float3(_41.row_major0[1])));
|
|
_41.row_major0[0] = float3(loaded[0][0], loaded[1][0], loaded[2][0]);
|
|
_41.row_major0[1] = float3(loaded[0][1], loaded[1][1], loaded[2][1]);
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void copy_col_major_to_col_major(device SSBOCol& _29)
|
|
{
|
|
_29.col_major0 = _29.col_major1;
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void copy_col_major_to_row_major(device SSBOCol& _29, device SSBORow& _41)
|
|
{
|
|
_41.row_major0[0] = float3(_29.col_major0[0][0], _29.col_major0[1][0], _29.col_major0[2][0]);
|
|
_41.row_major0[1] = float3(_29.col_major0[0][1], _29.col_major0[1][1], _29.col_major0[2][1]);
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void copy_row_major_to_col_major(device SSBOCol& _29, device SSBORow& _41)
|
|
{
|
|
_29.col_major0 = transpose(float2x3(float3(_41.row_major0[0]), float3(_41.row_major0[1])));
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void copy_row_major_to_row_major(device SSBORow& _41)
|
|
{
|
|
_41.row_major0[0] = float2x3(float3(_41.row_major1[0]), float3(_41.row_major1[1]))[0];
|
|
_41.row_major0[1] = float2x3(float3(_41.row_major1[0]), float3(_41.row_major1[1]))[1];
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void copy_columns(device SSBOCol& _29, device SSBORow& _41)
|
|
{
|
|
_29.col_major0[1] = float2(_41.row_major0[0][1], _41.row_major0[1][1]);
|
|
((device float*)&_41.row_major0[0])[1] = _29.col_major0[1].x;
|
|
((device float*)&_41.row_major0[1])[1] = _29.col_major0[1].y;
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void copy_elements(device SSBOCol& _29, device SSBORow& _41)
|
|
{
|
|
((device float*)&_29.col_major0[0])[1u] = _41.row_major0[1u][0];
|
|
_41.row_major0[1u][0] = ((device float*)&_29.col_major0[0])[1u];
|
|
}
|
|
|
|
kernel void main0(device SSBOCol& _29 [[buffer(0)]], device SSBORow& _41 [[buffer(1)]])
|
|
{
|
|
load_store_to_variable_col_major(_29);
|
|
load_store_to_variable_row_major(_41);
|
|
copy_col_major_to_col_major(_29);
|
|
copy_col_major_to_row_major(_29, _41);
|
|
copy_row_major_to_col_major(_29, _41);
|
|
copy_row_major_to_row_major(_41);
|
|
copy_columns(_29, _41);
|
|
copy_elements(_29, _41);
|
|
}
|
|
|