607b0d6d42
Support MSL typedefs to declare 3-row row-major matrices as 3-column matrices. Allow those matrices to be decorated as packed. Support transposing those matrices when used. Modify how member alignments are calculated.
106 lines
2.3 KiB
Plaintext
106 lines
2.3 KiB
Plaintext
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
typedef float3x2 packed_float2x3;
|
|
|
|
struct S0
|
|
{
|
|
float2 a[1];
|
|
float b;
|
|
};
|
|
|
|
struct S1
|
|
{
|
|
packed_float3 a;
|
|
float b;
|
|
};
|
|
|
|
struct S2
|
|
{
|
|
float3 a[1];
|
|
float b;
|
|
};
|
|
|
|
struct S3
|
|
{
|
|
float2 a;
|
|
float b;
|
|
};
|
|
|
|
struct S4
|
|
{
|
|
float2 c;
|
|
};
|
|
|
|
struct Content
|
|
{
|
|
S0 m0s[1];
|
|
S1 m1s[1];
|
|
S2 m2s[1];
|
|
S0 m0;
|
|
S1 m1;
|
|
S2 m2;
|
|
S3 m3;
|
|
char pad7[4];
|
|
float m4;
|
|
S4 m3s[8];
|
|
};
|
|
|
|
struct SSBO1
|
|
{
|
|
Content content;
|
|
Content content1[2];
|
|
Content content2;
|
|
char pad3[8];
|
|
float2x2 m0;
|
|
float2x2 m1;
|
|
float2x3 m2[4];
|
|
float3x2 m3;
|
|
float2x2 m4;
|
|
float2x2 m5[9];
|
|
packed_float2x3 m6[4][2];
|
|
char pad10[8];
|
|
float3x2 m7;
|
|
char pad11[8];
|
|
float array[1];
|
|
};
|
|
|
|
struct SSBO0
|
|
{
|
|
Content content;
|
|
Content content1[2];
|
|
Content content2;
|
|
float array[1];
|
|
};
|
|
|
|
kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]])
|
|
{
|
|
ssbo_430.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
|
|
ssbo_430.content.m0s[0].b = ssbo_140.content.m0s[0].b;
|
|
ssbo_430.content.m1s[0].a = ssbo_140.content.m1s[0].a;
|
|
ssbo_430.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
|
ssbo_430.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
|
|
ssbo_430.content.m2s[0].b = ssbo_140.content.m2s[0].b;
|
|
ssbo_430.content.m0.a[0] = ssbo_140.content.m0.a[0];
|
|
ssbo_430.content.m0.b = ssbo_140.content.m0.b;
|
|
ssbo_430.content.m1.a = ssbo_140.content.m1.a;
|
|
ssbo_430.content.m1.b = ssbo_140.content.m1.b;
|
|
ssbo_430.content.m2.a[0] = ssbo_140.content.m2.a[0];
|
|
ssbo_430.content.m2.b = ssbo_140.content.m2.b;
|
|
ssbo_430.content.m3.a = ssbo_140.content.m3.a;
|
|
ssbo_430.content.m3.b = ssbo_140.content.m3.b;
|
|
ssbo_430.content.m4 = ssbo_140.content.m4;
|
|
ssbo_430.content.m3s[0].c = ssbo_140.content.m3s[0].c;
|
|
ssbo_430.content.m3s[1].c = ssbo_140.content.m3s[1].c;
|
|
ssbo_430.content.m3s[2].c = ssbo_140.content.m3s[2].c;
|
|
ssbo_430.content.m3s[3].c = ssbo_140.content.m3s[3].c;
|
|
ssbo_430.content.m3s[4].c = ssbo_140.content.m3s[4].c;
|
|
ssbo_430.content.m3s[5].c = ssbo_140.content.m3s[5].c;
|
|
ssbo_430.content.m3s[6].c = ssbo_140.content.m3s[6].c;
|
|
ssbo_430.content.m3s[7].c = ssbo_140.content.m3s[7].c;
|
|
ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
|
|
}
|
|
|