From e406fb215cdddba985e2df2e0f0a2f22f0920fa8 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Tue, 21 May 2024 13:52:15 +0200 Subject: [PATCH] Fixup invalid shader that should not pass. --- .../comp/basic.dynamic-buffer.msl2.comp | 66 ++++++++++++++ .../basic.dynamic-buffer.msl2.invalid.comp | 90 ------------------- ...id.comp => basic.dynamic-buffer.msl2.comp} | 6 +- 3 files changed, 69 insertions(+), 93 deletions(-) create mode 100644 reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp delete mode 100644 reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp rename shaders-msl-no-opt/comp/{basic.dynamic-buffer.msl2.invalid.comp => basic.dynamic-buffer.msl2.comp} (73%) diff --git a/reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp b/reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp new file mode 100644 index 00000000..3bc1ae24 --- /dev/null +++ b/reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp @@ -0,0 +1,66 @@ +#include +#include + +using namespace metal; + +struct Baz +{ + int e; + int f; +}; + +struct Foo +{ + int a; + int b; +}; + +struct Bar +{ + int c; + int d; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, 3u, 2u); + +struct spvDescriptorSetBuffer0 +{ + constant Foo* m_34 [[id(0)]]; + constant Bar* m_40 [[id(1)]]; +}; + +struct spvDescriptorSetBuffer1 +{ + device Baz* baz [[id(0)]][18]; +}; + +kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant uint* spvDynamicOffsets [[buffer(23)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + constant auto& _34 = *(constant Foo* )((constant char* )spvDescriptorSet0.m_34 + spvDynamicOffsets[0]); + device Baz* baz[18] = + { + (device Baz* )((device char* )spvDescriptorSet1.baz[0] + spvDynamicOffsets[1]), + (device Baz* )((device char* )spvDescriptorSet1.baz[1] + spvDynamicOffsets[2]), + (device Baz* )((device char* )spvDescriptorSet1.baz[2] + spvDynamicOffsets[3]), + (device Baz* )((device char* )spvDescriptorSet1.baz[3] + spvDynamicOffsets[4]), + (device Baz* )((device char* )spvDescriptorSet1.baz[4] + spvDynamicOffsets[5]), + (device Baz* )((device char* )spvDescriptorSet1.baz[5] + spvDynamicOffsets[6]), + (device Baz* )((device char* )spvDescriptorSet1.baz[6] + spvDynamicOffsets[7]), + (device Baz* )((device char* )spvDescriptorSet1.baz[7] + spvDynamicOffsets[8]), + (device Baz* )((device char* )spvDescriptorSet1.baz[8] + spvDynamicOffsets[9]), + (device Baz* )((device char* )spvDescriptorSet1.baz[9] + spvDynamicOffsets[10]), + (device Baz* )((device char* )spvDescriptorSet1.baz[10] + spvDynamicOffsets[11]), + (device Baz* )((device char* )spvDescriptorSet1.baz[11] + spvDynamicOffsets[12]), + (device Baz* )((device char* )spvDescriptorSet1.baz[12] + spvDynamicOffsets[13]), + (device Baz* )((device char* )spvDescriptorSet1.baz[13] + spvDynamicOffsets[14]), + (device Baz* )((device char* )spvDescriptorSet1.baz[14] + spvDynamicOffsets[15]), + (device Baz* )((device char* )spvDescriptorSet1.baz[15] + spvDynamicOffsets[16]), + (device Baz* )((device char* )spvDescriptorSet1.baz[16] + spvDynamicOffsets[17]), + (device Baz* )((device char* )spvDescriptorSet1.baz[17] + spvDynamicOffsets[18]), + }; + + uint3 coords = gl_GlobalInvocationID; + baz[(coords.x + coords.y) + coords.z]->e = _34.a + (*spvDescriptorSet0.m_40).c; + baz[(coords.x + coords.y) + coords.z]->f = _34.b * (*spvDescriptorSet0.m_40).d; +} + diff --git a/reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp b/reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp deleted file mode 100644 index ae8c5b02..00000000 --- a/reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp +++ /dev/null @@ -1,90 +0,0 @@ -#include -#include - -using namespace metal; - -struct Baz -{ - int e; - int f; -}; - -struct Foo -{ - int a; - int b; -}; - -struct Bar -{ - int c; - int d; -}; - -constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, 3u, 2u); - -struct spvDescriptorSetBuffer0 -{ - constant Foo* m_34 [[id(0)]]; - constant Bar* m_40 [[id(1)]]; -}; - -struct spvDescriptorSetBuffer1 -{ - device Baz* baz [[id(0)]][3][3][2]; -}; - -kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant uint* spvDynamicOffsets [[buffer(23)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) -{ - constant auto& _34 = *(constant Foo* )((constant char* )spvDescriptorSet0.m_34 + spvDynamicOffsets[0]); - device Baz* baz[3][3][2] = - { - { - { - (device Baz* )((device char* )spvDescriptorSet1.baz[0][0][0] + spvDynamicOffsets[1]), - (device Baz* )((device char* )spvDescriptorSet1.baz[0][0][1] + spvDynamicOffsets[2]), - }, - { - (device Baz* )((device char* )spvDescriptorSet1.baz[0][1][0] + spvDynamicOffsets[3]), - (device Baz* )((device char* )spvDescriptorSet1.baz[0][1][1] + spvDynamicOffsets[4]), - }, - { - (device Baz* )((device char* )spvDescriptorSet1.baz[0][2][0] + spvDynamicOffsets[5]), - (device Baz* )((device char* )spvDescriptorSet1.baz[0][2][1] + spvDynamicOffsets[6]), - }, - }, - { - { - (device Baz* )((device char* )spvDescriptorSet1.baz[1][0][0] + spvDynamicOffsets[7]), - (device Baz* )((device char* )spvDescriptorSet1.baz[1][0][1] + spvDynamicOffsets[8]), - }, - { - (device Baz* )((device char* )spvDescriptorSet1.baz[1][1][0] + spvDynamicOffsets[9]), - (device Baz* )((device char* )spvDescriptorSet1.baz[1][1][1] + spvDynamicOffsets[10]), - }, - { - (device Baz* )((device char* )spvDescriptorSet1.baz[1][2][0] + spvDynamicOffsets[11]), - (device Baz* )((device char* )spvDescriptorSet1.baz[1][2][1] + spvDynamicOffsets[12]), - }, - }, - { - { - (device Baz* )((device char* )spvDescriptorSet1.baz[2][0][0] + spvDynamicOffsets[13]), - (device Baz* )((device char* )spvDescriptorSet1.baz[2][0][1] + spvDynamicOffsets[14]), - }, - { - (device Baz* )((device char* )spvDescriptorSet1.baz[2][1][0] + spvDynamicOffsets[15]), - (device Baz* )((device char* )spvDescriptorSet1.baz[2][1][1] + spvDynamicOffsets[16]), - }, - { - (device Baz* )((device char* )spvDescriptorSet1.baz[2][2][0] + spvDynamicOffsets[17]), - (device Baz* )((device char* )spvDescriptorSet1.baz[2][2][1] + spvDynamicOffsets[18]), - }, - }, - }; - - uint3 coords = gl_GlobalInvocationID; - baz[coords.x][coords.y][coords.z]->e = _34.a + (*spvDescriptorSet0.m_40).c; - baz[coords.x][coords.y][coords.z]->f = _34.b * (*spvDescriptorSet0.m_40).d; -} - diff --git a/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp b/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp similarity index 73% rename from shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp rename to shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp index c2965731..abf100f4 100644 --- a/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp +++ b/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp @@ -17,11 +17,11 @@ layout(set = 1, binding = 2) buffer Baz { int e; int f; -} baz[3][3][2]; +} baz[3 * 3 * 2]; void main() { uvec3 coords = gl_GlobalInvocationID; - baz[coords.x][coords.y][coords.z].e = a + c; - baz[coords.x][coords.y][coords.z].f = b * d; + baz[coords.x + coords.y + coords.z].e = a + c; + baz[coords.x + coords.y + coords.z].f = b * d; }