From 38d19821d4cf772bf1b886faede9cadf2dd2b991 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Tue, 11 Sep 2018 12:58:03 +0200 Subject: [PATCH] MSL: Support copying array of arrays. --- .../frag/lut-promotion-initializer.asm.frag | 7 +- .../comp/composite-array-initialization.comp | 7 +- .../shaders-msl/comp/composite-construct.comp | 7 +- .../comp/copy-array-of-arrays.comp | 15 +++ .../opt/shaders-msl/frag/constant-array.frag | 5 +- .../shaders-msl/frag/constant-composites.frag | 5 +- .../opt/shaders-msl/frag/lut-promotion.frag | 7 +- .../frag/lut-promotion-initializer.asm.frag | 7 +- .../comp/composite-array-initialization.comp | 7 +- .../shaders-msl/comp/composite-construct.comp | 7 +- .../comp/copy-array-of-arrays.comp | 73 +++++++++++++ .../shaders-msl/frag/constant-array.frag | 5 +- .../shaders-msl/frag/constant-composites.frag | 5 +- reference/shaders-msl/frag/lut-promotion.frag | 7 +- reference/shaders-msl/vert/return-array.vert | 9 +- shaders-msl/comp/copy-array-of-arrays.comp | 15 +++ spirv_msl.cpp | 100 ++++++++++++++++-- spirv_msl.hpp | 11 +- 18 files changed, 241 insertions(+), 58 deletions(-) create mode 100644 reference/opt/shaders-msl/comp/copy-array-of-arrays.comp create mode 100644 reference/shaders-msl/comp/copy-array-of-arrays.comp create mode 100644 shaders-msl/comp/copy-array-of-arrays.comp diff --git a/reference/opt/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag b/reference/opt/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag index 6563cf25..610d447a 100644 --- a/reference/opt/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag +++ b/reference/opt/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag @@ -21,14 +21,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -62,7 +61,7 @@ fragment main0_out main0(main0_in in [[stage_in]]) } int _37 = in.index & 3; out.FragColor += foobar[_37].z; - spvArrayCopyConstant(baz, _90); + spvArrayCopyFromConstant1(baz, _90); out.FragColor += baz[_37].z; return out; } diff --git a/reference/opt/shaders-msl/comp/composite-array-initialization.comp b/reference/opt/shaders-msl/comp/composite-array-initialization.comp index 2003efa3..9e47e897 100644 --- a/reference/opt/shaders-msl/comp/composite-array-initialization.comp +++ b/reference/opt/shaders-msl/comp/composite-array-initialization.comp @@ -30,14 +30,13 @@ constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } }; // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -47,7 +46,7 @@ kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadg Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } }; Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } }; Data data2[2]; - spvArrayCopy(data2, _31); + spvArrayCopyFromStack1(data2, _31); _53.outdata[gl_WorkGroupID.x].a = data[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a; _53.outdata[gl_WorkGroupID.x].b = data[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b; } diff --git a/reference/opt/shaders-msl/comp/composite-construct.comp b/reference/opt/shaders-msl/comp/composite-construct.comp index 2d9fdcca..f5da3df3 100644 --- a/reference/opt/shaders-msl/comp/composite-construct.comp +++ b/reference/opt/shaders-msl/comp/composite-construct.comp @@ -12,14 +12,13 @@ struct SSBO0 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -28,7 +27,7 @@ kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO0& _32 [[buffer(1) { float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.as[gl_GlobalInvocationID.x] }; float4 values[2]; - spvArrayCopy(values, _37); + spvArrayCopyFromStack1(values, _37); _16.as[0] = values[gl_LocalInvocationIndex]; _32.as[1] = float4(40.0); } diff --git a/reference/opt/shaders-msl/comp/copy-array-of-arrays.comp b/reference/opt/shaders-msl/comp/copy-array-of-arrays.comp new file mode 100644 index 00000000..e23e2468 --- /dev/null +++ b/reference/opt/shaders-msl/comp/copy-array-of-arrays.comp @@ -0,0 +1,15 @@ +#include +#include + +using namespace metal; + +struct BUF +{ + int a; +}; + +kernel void main0(device BUF& o [[buffer(0)]]) +{ + o.a = 4; +} + diff --git a/reference/opt/shaders-msl/frag/constant-array.frag b/reference/opt/shaders-msl/frag/constant-array.frag index 1e0c27ec..0ba2781b 100644 --- a/reference/opt/shaders-msl/frag/constant-array.frag +++ b/reference/opt/shaders-msl/frag/constant-array.frag @@ -29,14 +29,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } diff --git a/reference/opt/shaders-msl/frag/constant-composites.frag b/reference/opt/shaders-msl/frag/constant-composites.frag index b3076cfe..2c91b111 100644 --- a/reference/opt/shaders-msl/frag/constant-composites.frag +++ b/reference/opt/shaders-msl/frag/constant-composites.frag @@ -26,14 +26,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } diff --git a/reference/opt/shaders-msl/frag/lut-promotion.frag b/reference/opt/shaders-msl/frag/lut-promotion.frag index 7af74aba..ddba6083 100644 --- a/reference/opt/shaders-msl/frag/lut-promotion.frag +++ b/reference/opt/shaders-msl/frag/lut-promotion.frag @@ -21,14 +21,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -62,7 +61,7 @@ fragment main0_out main0(main0_in in [[stage_in]]) int _91 = in.index & 3; out.FragColor += foobar[_91].z; float4 baz[4] = { float4(0.0), float4(1.0), float4(8.0), float4(5.0) }; - spvArrayCopyConstant(baz, _104); + spvArrayCopyFromConstant1(baz, _104); out.FragColor += baz[_91].z; return out; } diff --git a/reference/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag b/reference/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag index bc9a3bb5..48f3317d 100644 --- a/reference/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag +++ b/reference/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag @@ -21,14 +21,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -60,7 +59,7 @@ fragment main0_out main0(main0_in in [[stage_in]]) foobar[1].z = 20.0; } out.FragColor += foobar[in.index & 3].z; - spvArrayCopyConstant(baz, _90); + spvArrayCopyFromConstant1(baz, _90); out.FragColor += baz[in.index & 3].z; return out; } diff --git a/reference/shaders-msl/comp/composite-array-initialization.comp b/reference/shaders-msl/comp/composite-array-initialization.comp index 3e67fcbf..1a901dc8 100644 --- a/reference/shaders-msl/comp/composite-array-initialization.comp +++ b/reference/shaders-msl/comp/composite-array-initialization.comp @@ -30,14 +30,13 @@ constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } }; // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -52,7 +51,7 @@ kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadg Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } }; Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } }; Data data2[2]; - spvArrayCopy(data2, _31); + spvArrayCopyFromStack1(data2, _31); Data param = data[gl_LocalInvocationID.x]; Data param_1 = data2[gl_LocalInvocationID.x]; Data _73 = combine(param, param_1); diff --git a/reference/shaders-msl/comp/composite-construct.comp b/reference/shaders-msl/comp/composite-construct.comp index 1864f04c..4b5ea37e 100644 --- a/reference/shaders-msl/comp/composite-construct.comp +++ b/reference/shaders-msl/comp/composite-construct.comp @@ -25,14 +25,13 @@ constant float4 _43[2] = { float4(20.0), float4(40.0) }; // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -41,7 +40,7 @@ kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO1& _32 [[buffer(1) { float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.bs[gl_GlobalInvocationID.x] }; float4 values[2]; - spvArrayCopy(values, _37); + spvArrayCopyFromStack1(values, _37); Composite c = Composite{ values[0], _43[1] }; _16.as[0] = values[gl_LocalInvocationIndex]; _32.bs[1] = c.b; diff --git a/reference/shaders-msl/comp/copy-array-of-arrays.comp b/reference/shaders-msl/comp/copy-array-of-arrays.comp new file mode 100644 index 00000000..fd48fbe3 --- /dev/null +++ b/reference/shaders-msl/comp/copy-array-of-arrays.comp @@ -0,0 +1,73 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct BUF +{ + int a; +}; + +constant float _16[2] = { 1.0, 2.0 }; +constant float _19[2] = { 3.0, 4.0 }; +constant float _20[2][2] = { { 1.0, 2.0 }, { 3.0, 4.0 } }; +constant float _21[2][2][2] = { { { 1.0, 2.0 }, { 3.0, 4.0 } }, { { 1.0, 2.0 }, { 3.0, 4.0 } } }; + +// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. +template +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) +{ + for (uint i = 0; i < N; dst[i] = src[i], i++); +} + +template +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) +{ + for (uint i = 0; i < N; dst[i] = src[i], i++); +} + +template +void spvArrayCopyFromStack2(thread T (&dst)[A][B], thread const T (&src)[A][B]) +{ + for (uint i = 0; i < A; i++) + { + spvArrayCopyFromStack1(dst[i], src[i]); + } +} + +template +void spvArrayCopyFromConstant2(thread T (&dst)[A][B], constant T (&src)[A][B]) +{ + for (uint i = 0; i < A; i++) + { + spvArrayCopyFromConstant1(dst[i], src[i]); + } +} + +template +void spvArrayCopyFromStack3(thread T (&dst)[A][B][C], thread const T (&src)[A][B][C]) +{ + for (uint i = 0; i < A; i++) + { + spvArrayCopyFromStack2(dst[i], src[i]); + } +} + +template +void spvArrayCopyFromConstant3(thread T (&dst)[A][B][C], constant T (&src)[A][B][C]) +{ + for (uint i = 0; i < A; i++) + { + spvArrayCopyFromConstant2(dst[i], src[i]); + } +} + +kernel void main0(device BUF& o [[buffer(0)]]) +{ + float c[2][2][2]; + spvArrayCopyFromConstant3(c, _21); + o.a = int(c[1][1][1]); +} + diff --git a/reference/shaders-msl/frag/constant-array.frag b/reference/shaders-msl/frag/constant-array.frag index 78674b7c..f6989a2a 100644 --- a/reference/shaders-msl/frag/constant-array.frag +++ b/reference/shaders-msl/frag/constant-array.frag @@ -29,14 +29,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } diff --git a/reference/shaders-msl/frag/constant-composites.frag b/reference/shaders-msl/frag/constant-composites.frag index b3076cfe..2c91b111 100644 --- a/reference/shaders-msl/frag/constant-composites.frag +++ b/reference/shaders-msl/frag/constant-composites.frag @@ -26,14 +26,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } diff --git a/reference/shaders-msl/frag/lut-promotion.frag b/reference/shaders-msl/frag/lut-promotion.frag index 49ea0773..f7e51edb 100644 --- a/reference/shaders-msl/frag/lut-promotion.frag +++ b/reference/shaders-msl/frag/lut-promotion.frag @@ -21,14 +21,13 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } @@ -60,7 +59,7 @@ fragment main0_out main0(main0_in in [[stage_in]]) } out.FragColor += foobar[in.index & 3].z; float4 baz[4] = { float4(0.0), float4(1.0), float4(8.0), float4(5.0) }; - spvArrayCopyConstant(baz, _104); + spvArrayCopyFromConstant1(baz, _104); out.FragColor += baz[in.index & 3].z; return out; } diff --git a/reference/shaders-msl/vert/return-array.vert b/reference/shaders-msl/vert/return-array.vert index c51da87e..cd06fdda 100644 --- a/reference/shaders-msl/vert/return-array.vert +++ b/reference/shaders-msl/vert/return-array.vert @@ -20,21 +20,20 @@ struct main0_in // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment. template -void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N]) +void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } -// An overload for constant arrays. template -void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N]) +void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N]) { for (uint i = 0; i < N; dst[i] = src[i], i++); } void test(thread float4 (&SPIRV_Cross_return_value)[2]) { - spvArrayCopyConstant(SPIRV_Cross_return_value, _20); + spvArrayCopyFromConstant1(SPIRV_Cross_return_value, _20); } void test2(thread float4 (&SPIRV_Cross_return_value)[2], thread float4& vInput0, thread float4& vInput1) @@ -42,7 +41,7 @@ void test2(thread float4 (&SPIRV_Cross_return_value)[2], thread float4& vInput0, float4 foobar[2]; foobar[0] = vInput0; foobar[1] = vInput1; - spvArrayCopy(SPIRV_Cross_return_value, foobar); + spvArrayCopyFromStack1(SPIRV_Cross_return_value, foobar); } vertex main0_out main0(main0_in in [[stage_in]]) diff --git a/shaders-msl/comp/copy-array-of-arrays.comp b/shaders-msl/comp/copy-array-of-arrays.comp new file mode 100644 index 00000000..a23bfb1b --- /dev/null +++ b/shaders-msl/comp/copy-array-of-arrays.comp @@ -0,0 +1,15 @@ +#version 450 +layout(local_size_x = 1) in; + +layout(set = 0, binding = 0, std430) buffer BUF +{ + int a; +} o; + +void main() +{ + const float a[2][2][2] = float[][][](float[][](float[](1.0, 2.0), float[](3.0, 4.0)), float[][](float[](1.0, 2.0), float[](3.0, 4.0))); + float b[2][2][2] = a; + float c[2][2][2] = b; + o.a = int(c[1][1][1]); +} diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 207ebc3c..d1551e16 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -1168,6 +1168,10 @@ void CompilerMSL::add_typedef_line(const string &line) // Emits any needed custom function bodies. void CompilerMSL::emit_custom_functions() { + for (uint32_t i = SPVFuncImplArrayCopyMultidimMax; i >= 2; i--) + if (spv_function_implementations.count(static_cast(SPVFuncImplArrayCopyMultidimBase + i))) + spv_function_implementations.insert(static_cast(SPVFuncImplArrayCopyMultidimBase + i - 1)); + for (auto &spv_func : spv_function_implementations) { switch (spv_func) @@ -1237,21 +1241,72 @@ void CompilerMSL::emit_custom_functions() statement("// Implementation of an array copy function to cover GLSL's ability to copy an array via " "assignment."); statement("template"); - statement("void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])"); + statement("void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])"); begin_scope(); statement("for (uint i = 0; i < N; dst[i] = src[i], i++);"); end_scope(); statement(""); - statement("// An overload for constant arrays."); statement("template"); - statement("void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])"); + statement("void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])"); begin_scope(); statement("for (uint i = 0; i < N; dst[i] = src[i], i++);"); end_scope(); statement(""); break; + case SPVFuncImplArrayOfArrayCopy2Dim: + case SPVFuncImplArrayOfArrayCopy3Dim: + case SPVFuncImplArrayOfArrayCopy4Dim: + case SPVFuncImplArrayOfArrayCopy5Dim: + case SPVFuncImplArrayOfArrayCopy6Dim: + { + static const char *function_name_tags[] = { + "FromStack", + "FromConstant", + }; + + static const char *src_address_space[] = { + "thread const", + "constant", + }; + + for (uint32_t variant = 0; variant < 2; variant++) + { + uint32_t dimensions = spv_func - SPVFuncImplArrayCopyMultidimBase; + string tmp = "templateremapped_variable && var->statically_assigned && + ids[var->static_expression].get_type() == TypeConstant) + { + is_constant = true; + } + + const char *tag = is_constant ? "FromConstant" : "FromStack"; + statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ", ", to_expression(rhs_id), ");"); } // Since MSL does not allow arrays to be copied via simple variable assignment, @@ -4380,7 +4449,13 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o case OpFunctionCall: { auto &return_type = compiler.get(args[0]); - if (!return_type.array.empty()) + if (return_type.array.size() > 1) + { + if (return_type.array.size() > SPVFuncImplArrayCopyMultidimMax) + SPIRV_CROSS_THROW("Cannot support this many dimensions for arrays of arrays."); + return static_cast(SPVFuncImplArrayCopyMultidimBase + return_type.array.size()); + } + else if (return_type.array.size() > 0) return SPVFuncImplArrayCopy; break; @@ -4414,7 +4489,16 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o bool static_expression_lhs = var && var->storage == StorageClassFunction && var->statically_assigned && var->remapped_variable; if (type && compiler.is_array(*type) && !static_expression_lhs) - return SPVFuncImplArrayCopy; + { + if (type->array.size() > 1) + { + if (type->array.size() > SPVFuncImplArrayCopyMultidimMax) + SPIRV_CROSS_THROW("Cannot support this many dimensions for arrays of arrays."); + return static_cast(SPVFuncImplArrayCopyMultidimBase + type->array.size()); + } + else + return SPVFuncImplArrayCopy; + } break; } diff --git a/spirv_msl.hpp b/spirv_msl.hpp index e7a6ad85..00350ed4 100644 --- a/spirv_msl.hpp +++ b/spirv_msl.hpp @@ -222,7 +222,15 @@ public: SPVFuncImplFindILsb, SPVFuncImplFindSMsb, SPVFuncImplFindUMsb, - SPVFuncImplArrayCopy, + SPVFuncImplArrayCopyMultidimBase, + // Unfortunately, we cannot use recursive templates in the MSL compiler properly, + // so stamp out variants up to some arbitrary maximum. + SPVFuncImplArrayCopy = SPVFuncImplArrayCopyMultidimBase + 1, + SPVFuncImplArrayOfArrayCopy2Dim = SPVFuncImplArrayCopyMultidimBase + 2, + SPVFuncImplArrayOfArrayCopy3Dim = SPVFuncImplArrayCopyMultidimBase + 3, + SPVFuncImplArrayOfArrayCopy4Dim = SPVFuncImplArrayCopyMultidimBase + 4, + SPVFuncImplArrayOfArrayCopy5Dim = SPVFuncImplArrayCopyMultidimBase + 5, + SPVFuncImplArrayOfArrayCopy6Dim = SPVFuncImplArrayCopyMultidimBase + 6, SPVFuncImplTexelBufferCoords, SPVFuncImplInverse4x4, SPVFuncImplInverse3x3, @@ -233,6 +241,7 @@ public: SPVFuncImplRowMajor3x4, SPVFuncImplRowMajor4x2, SPVFuncImplRowMajor4x3, + SPVFuncImplArrayCopyMultidimMax = 6 }; // Constructs an instance to compile the SPIR-V code into Metal Shading Language,