1c18078811
Support Workgroup (threadgroup) variables. Mark if SPIRConstant is used as an array length, since it cannot be specialized. Resolve specialized array length constants. Support passing an array to MSL function. Support emitting GLSL array assignments in MSL via an array copy function. Support for memory and control barriers. Struct packing enhancements, including packing nested structs. Enhancements to replacing illegal MSL variable and function names. Add Compiler::get_entry_point_name_map() function to retrieve entry point renamings. Remove CompilerGLSL::clean_func_name() as obsolete. Fixes to types in bitcast MSL functions. Add Variant::get_id() member function. Add CompilerMSL::Options::msl_version option. Add numerous MSL compute tests.
101 lines
2.2 KiB
Plaintext
101 lines
2.2 KiB
Plaintext
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
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];
|
|
float2x3 m6[4][2];
|
|
float3x2 m7;
|
|
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;
|
|
}
|
|
|