From cf1bf1c6aeada991e28fbbd55908f30875f16194 Mon Sep 17 00:00:00 2001 From: Ryan Harrison Date: Wed, 18 Sep 2019 15:56:51 -0400 Subject: [PATCH] Update external/ to SPIR-V 1.5 Rolled the hashes used for glslang, SPIRV-Tools, and SPIRV-Headers to HEAD, which includes the update to 1.5. Added passing '--amb' to glslang, so I didn't have to explicitly set bindings in a large number of test shaders that currently don't, and now glslang considers them invalid. Marked all shaders that no longer pass spirv-val as .invalid. --- checkout_glslang_spirv_tools.sh | 6 +- .../opt/shaders-hlsl/comp/access-chains.comp | 2 + .../shaders-hlsl/comp/address-buffers.comp | 2 + reference/opt/shaders-hlsl/comp/atomic.comp | 2 + .../shaders-hlsl/comp/globallycoherent.comp | 2 + reference/opt/shaders-hlsl/comp/image.comp | 2 + reference/opt/shaders-hlsl/comp/inverse.comp | 2 + .../comp/num-workgroups-alone.comp | 2 + .../comp/num-workgroups-with-builtins.comp | 2 + .../opt/shaders-hlsl/comp/outer-product.comp | 2 + .../opt/shaders-hlsl/comp/rmw-matrix.comp | 2 + .../shaders-hlsl/comp/rwbuffer-matrix.comp | 104 +++++++++--------- ...alar-std450-distance-length-normalize.comp | 2 + .../comp/spec-constant-op-member-array.comp | 1 + .../shaders-hlsl/comp/ssbo-array-length.comp | 2 + .../opt/shaders-hlsl/comp/ssbo-array.comp | 2 + .../comp/subgroups.invalid.nofxc.sm60.comp | 2 + .../access-private-workgroup-in-function.comp | 2 + ...ffers-discrete.msl2.argument.discrete.comp | 2 + .../opt/shaders-msl/comp/array-length.comp | 2 + .../array-length.msl2.argument.discrete.comp | 2 + reference/opt/shaders-msl/comp/atomic.comp | 2 + reference/opt/shaders-msl/comp/basic.comp | 2 + .../comp/basic.dispatchbase.msl11.comp | 4 +- ...=> basic.dynamic-buffer.msl2.invalid.comp} | 0 .../comp/bitcast-16bit-1.invalid.comp | 2 + .../comp/bitcast-16bit-2.invalid.comp | 2 + .../opt/shaders-msl/comp/coherent-block.comp | 2 + .../opt/shaders-msl/comp/coherent-image.comp | 2 + .../shaders-msl/comp/composite-construct.comp | 2 + .../comp/copy-array-of-arrays.comp | 2 + .../opt/shaders-msl/comp/defer-parens.comp | 2 + reference/opt/shaders-msl/comp/dowhile.comp | 18 +-- .../comp/image-cube-array-load-store.comp | 2 + reference/opt/shaders-msl/comp/image.comp | 2 + reference/opt/shaders-msl/comp/insert.comp | 22 ++-- .../shaders-msl/comp/int64.invalid.msl22.comp | 18 +-- reference/opt/shaders-msl/comp/inverse.comp | 2 + reference/opt/shaders-msl/comp/mat3.comp | 2 + reference/opt/shaders-msl/comp/mod.comp | 2 + reference/opt/shaders-msl/comp/modf.comp | 2 + .../opt/shaders-msl/comp/outer-product.comp | 2 + .../opt/shaders-msl/comp/read-write-only.comp | 2 + .../opt/shaders-msl/comp/rmw-matrix.comp | 2 + reference/opt/shaders-msl/comp/rmw-opt.comp | 2 + ...alar-std450-distance-length-normalize.comp | 2 + .../comp/spec-constant-op-member-array.comp | 1 + .../storage-buffer-std140-vector-array.comp | 50 +++++---- .../opt/shaders-msl/comp/struct-layout.comp | 2 + .../opt/shaders-msl/comp/struct-nested.comp | 2 + .../opt/shaders-msl/comp/struct-packing.comp | 2 + .../opt/shaders-msl/comp/torture-loop.comp | 30 ++--- .../opt/shaders-msl/comp/type-alias.comp | 2 + reference/opt/shaders-msl/comp/udiv.comp | 2 + .../comp/extended-arithmetic.desktop.comp | 2 + .../frag/shader-arithmetic-8bit.frag | 57 ++++++---- ...ct-packing-scalar.nocompat.invalid.vk.comp | 2 + .../subgroups.nocompat.invalid.vk.msl21.comp | 2 + ...bgroups.nocompat.invalid.vk.msl21.ios.comp | 2 + ...e-to-helper.vk.nocompat.msl21.invalid.frag | 1 + reference/opt/shaders/comp/cfg.comp | 14 +-- reference/opt/shaders/comp/dowhile.comp | 16 +-- reference/opt/shaders/comp/insert.comp | 20 ++-- reference/opt/shaders/comp/torture-loop.comp | 28 ++--- .../comp/buffer-reference.nocompat.vk.comp.vk | 22 +--- ...shader-arithmetic-8bit.nocompat.vk.frag.vk | 57 ++++++---- ...h-block-case-fallthrough.asm.invalid.frag} | 0 .../comp/glsl.std450.fxconly.comp | 2 + .../shaders-hlsl/comp/access-chains.comp | 2 + .../shaders-hlsl/comp/address-buffers.comp | 2 + reference/shaders-hlsl/comp/atomic.comp | 2 + .../shaders-hlsl/comp/globallycoherent.comp | 2 + reference/shaders-hlsl/comp/image.comp | 2 + reference/shaders-hlsl/comp/inverse.comp | 2 + .../comp/num-workgroups-alone.comp | 2 + .../comp/num-workgroups-with-builtins.comp | 2 + .../shaders-hlsl/comp/outer-product.comp | 2 + reference/shaders-hlsl/comp/rmw-matrix.comp | 2 + .../shaders-hlsl/comp/rwbuffer-matrix.comp | 2 + ...alar-std450-distance-length-normalize.comp | 2 + .../comp/spec-constant-op-member-array.comp | 1 + .../shaders-hlsl/comp/ssbo-array-length.comp | 2 + reference/shaders-hlsl/comp/ssbo-array.comp | 2 + .../comp/subgroups.invalid.nofxc.sm60.comp | 2 + ...h-block-case-fallthrough.asm.invalid.frag} | 0 .../shaders-msl-no-opt/comp/glsl.std450.comp | 2 + reference/shaders-msl-no-opt/comp/loop.comp | 2 + reference/shaders-msl-no-opt/comp/return.comp | 2 + .../packing/matrix-2x2-scalar.comp | 2 + .../packing/matrix-2x2-std140.comp | 2 + .../packing/matrix-2x2-std430.comp | 2 + .../packing/matrix-2x3-scalar.comp | 2 + .../packing/matrix-2x3-std140.comp | 2 + .../packing/matrix-2x3-std430.comp | 2 + .../packing/matrix-2x4-scalar.comp | 2 + .../packing/matrix-2x4-std140.comp | 2 + .../packing/matrix-2x4-std430.comp | 2 + .../packing/matrix-3x2-scalar.comp | 2 + .../packing/matrix-3x2-std140.comp | 2 + .../packing/matrix-3x2-std430.comp | 2 + .../packing/matrix-3x3-scalar.comp | 2 + .../packing/matrix-3x3-std140.comp | 2 + .../packing/matrix-3x3-std430.comp | 2 + .../packing/matrix-3x4-scalar.comp | 2 + .../packing/matrix-3x4-std140.comp | 2 + .../packing/matrix-3x4-std430.comp | 2 + .../packing/matrix-4x2-scalar.comp | 2 + .../packing/matrix-4x2-std140.comp | 2 + .../packing/matrix-4x2-std430.comp | 2 + .../packing/matrix-4x3-scalar.comp | 2 + .../packing/matrix-4x3-std140.comp | 2 + .../packing/matrix-4x3-std430.comp | 2 + .../packing/matrix-4x4-scalar.comp | 2 + .../packing/matrix-4x4-std140.comp | 2 + .../packing/matrix-4x4-std430.comp | 2 + .../packing/matrix-multiply-row-major.comp | 2 + .../matrix-multiply-unpacked-col-major-2.comp | 2 + .../matrix-multiply-unpacked-col-major.comp | 2 + .../matrix-multiply-unpacked-row-major-2.comp | 2 + .../matrix-multiply-unpacked-row-major.comp | 2 + .../packing/member-padding.comp | 2 + .../packing/std140-array-of-vectors.comp | 2 + .../packing/struct-alignment.comp | 2 + .../struct-packing-array-of-scalar.comp | 2 + .../packing/struct-packing-recursive.comp | 2 + .../packing/struct-packing.comp | 2 + .../struct-size-padding-array-of-array.comp | 2 + .../packing/struct-size-padding.comp | 2 + .../access-private-workgroup-in-function.comp | 2 + ...ffers-discrete.msl2.argument.discrete.comp | 2 + reference/shaders-msl/comp/array-length.comp | 2 + .../array-length.msl2.argument.discrete.comp | 2 + reference/shaders-msl/comp/atomic.comp | 2 + reference/shaders-msl/comp/basic.comp | 2 + .../comp/basic.dispatchbase.msl11.comp | 4 +- ...=> basic.dynamic-buffer.msl2.invalid.comp} | 0 .../comp/bitcast-16bit-1.invalid.comp | 2 + .../comp/bitcast-16bit-2.invalid.comp | 2 + .../shaders-msl/comp/coherent-block.comp | 2 + .../shaders-msl/comp/coherent-image.comp | 2 + .../shaders-msl/comp/composite-construct.comp | 2 + .../comp/copy-array-of-arrays.comp | 2 + reference/shaders-msl/comp/defer-parens.comp | 2 + reference/shaders-msl/comp/dowhile.comp | 2 + .../comp/image-cube-array-load-store.comp | 2 + reference/shaders-msl/comp/image.comp | 2 + reference/shaders-msl/comp/insert.comp | 2 + .../shaders-msl/comp/int64.invalid.msl22.comp | 2 + reference/shaders-msl/comp/inverse.comp | 2 + reference/shaders-msl/comp/mat3.comp | 2 + reference/shaders-msl/comp/mod.comp | 2 + reference/shaders-msl/comp/modf.comp | 2 + reference/shaders-msl/comp/outer-product.comp | 2 + .../shaders-msl/comp/read-write-only.comp | 2 + reference/shaders-msl/comp/rmw-matrix.comp | 2 + reference/shaders-msl/comp/rmw-opt.comp | 2 + ...alar-std450-distance-length-normalize.comp | 2 + .../comp/spec-constant-op-member-array.comp | 1 + .../storage-buffer-std140-vector-array.comp | 2 + reference/shaders-msl/comp/struct-layout.comp | 2 + reference/shaders-msl/comp/struct-nested.comp | 2 + .../shaders-msl/comp/struct-packing.comp | 2 + reference/shaders-msl/comp/torture-loop.comp | 2 + reference/shaders-msl/comp/type-alias.comp | 2 + reference/shaders-msl/comp/udiv.comp | 2 + .../comp/extended-arithmetic.desktop.comp | 2 + ...ct-packing-scalar.nocompat.invalid.vk.comp | 2 + .../subgroups.nocompat.invalid.vk.msl21.comp | 2 + ...bgroups.nocompat.invalid.vk.msl21.ios.comp | 2 + ...e-sampler-dxc-min16float.asm.invalid.frag} | 0 ...h-block-case-fallthrough.asm.invalid.frag} | 0 ...h-block-case-fallthrough.asm.invalid.frag} | 0 ...h-block-case-fallthrough.asm.invalid.frag} | 0 ...=> basic.dynamic-buffer.msl2.invalid.comp} | 0 ...e-sampler-dxc-min16float.asm.invalid.frag} | 0 ...h-block-case-fallthrough.asm.invalid.frag} | 0 test_shaders.py | 12 +- 177 files changed, 552 insertions(+), 223 deletions(-) rename reference/opt/shaders-msl/comp/{basic.dynamic-buffer.msl2.comp => basic.dynamic-buffer.msl2.invalid.comp} (100%) rename reference/shaders-hlsl-no-opt/asm/frag/{switch-block-case-fallthrough.asm.frag => switch-block-case-fallthrough.asm.invalid.frag} (100%) rename reference/shaders-msl-no-opt/asm/frag/{switch-block-case-fallthrough.asm.frag => switch-block-case-fallthrough.asm.invalid.frag} (100%) rename reference/shaders-msl/comp/{basic.dynamic-buffer.msl2.comp => basic.dynamic-buffer.msl2.invalid.comp} (100%) rename reference/shaders-no-opt/asm/frag/{combined-image-sampler-dxc-min16float.asm.frag => combined-image-sampler-dxc-min16float.asm.invalid.frag} (100%) rename reference/shaders-no-opt/asm/frag/{switch-block-case-fallthrough.asm.frag => switch-block-case-fallthrough.asm.invalid.frag} (100%) rename shaders-hlsl-no-opt/asm/frag/{switch-block-case-fallthrough.asm.frag => switch-block-case-fallthrough.asm.invalid.frag} (100%) rename shaders-msl-no-opt/asm/frag/{switch-block-case-fallthrough.asm.frag => switch-block-case-fallthrough.asm.invalid.frag} (100%) rename shaders-msl/comp/{basic.dynamic-buffer.msl2.comp => basic.dynamic-buffer.msl2.invalid.comp} (100%) rename shaders-no-opt/asm/frag/{combined-image-sampler-dxc-min16float.asm.frag => combined-image-sampler-dxc-min16float.asm.invalid.frag} (100%) rename shaders-no-opt/asm/frag/{switch-block-case-fallthrough.asm.frag => switch-block-case-fallthrough.asm.invalid.frag} (100%) diff --git a/checkout_glslang_spirv_tools.sh b/checkout_glslang_spirv_tools.sh index bd2b5fab..54171f2b 100755 --- a/checkout_glslang_spirv_tools.sh +++ b/checkout_glslang_spirv_tools.sh @@ -1,8 +1,8 @@ #!/bin/bash -GLSLANG_REV=25a508cc735109cc4e382c3a1cc293a9452a41f3 -SPIRV_TOOLS_REV=55adf4cf707bb12c29fc12f784ebeaa29a819e9b -SPIRV_HEADERS_REV=29c11140baaf9f7fdaa39a583672c556bf1795a1 +GLSLANG_REV=4b5159ea8170fa34e29f13448fddebf88e0a722a +SPIRV_TOOLS_REV=bbb29870b510f83f99994358179c9ea6838c3100 +SPIRV_HEADERS_REV=601d738723ac381741311c6c98c36d6170be14a2 if [ -z $PROTOCOL ]; then PROTOCOL=git diff --git a/reference/opt/shaders-hlsl/comp/access-chains.comp b/reference/opt/shaders-hlsl/comp/access-chains.comp index 924e9191..c748200b 100644 --- a/reference/opt/shaders-hlsl/comp/access-chains.comp +++ b/reference/opt/shaders-hlsl/comp/access-chains.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer wo : register(u1); ByteAddressBuffer ro : register(t0); diff --git a/reference/opt/shaders-hlsl/comp/address-buffers.comp b/reference/opt/shaders-hlsl/comp/address-buffers.comp index a252fc8a..7f1c7975 100644 --- a/reference/opt/shaders-hlsl/comp/address-buffers.comp +++ b/reference/opt/shaders-hlsl/comp/address-buffers.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer WriteOnly : register(u2); ByteAddressBuffer ReadOnly : register(t0); RWByteAddressBuffer ReadWrite : register(u1); diff --git a/reference/opt/shaders-hlsl/comp/atomic.comp b/reference/opt/shaders-hlsl/comp/atomic.comp index 72e15bf7..e6ff891e 100644 --- a/reference/opt/shaders-hlsl/comp/atomic.comp +++ b/reference/opt/shaders-hlsl/comp/atomic.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer ssbo : register(u2); RWTexture2D uImage : register(u0); RWTexture2D iImage : register(u1); diff --git a/reference/opt/shaders-hlsl/comp/globallycoherent.comp b/reference/opt/shaders-hlsl/comp/globallycoherent.comp index 1637727d..b5f1e377 100644 --- a/reference/opt/shaders-hlsl/comp/globallycoherent.comp +++ b/reference/opt/shaders-hlsl/comp/globallycoherent.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + globallycoherent RWByteAddressBuffer _29 : register(u3); ByteAddressBuffer _33 : register(t2); RWTexture2D uImageIn : register(u0); diff --git a/reference/opt/shaders-hlsl/comp/image.comp b/reference/opt/shaders-hlsl/comp/image.comp index 6c2b58cd..e2f6b0a3 100644 --- a/reference/opt/shaders-hlsl/comp/image.comp +++ b/reference/opt/shaders-hlsl/comp/image.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWTexture2D uImageInF : register(u0); RWTexture2D uImageOutF : register(u1); RWTexture2D uImageInI : register(u2); diff --git a/reference/opt/shaders-hlsl/comp/inverse.comp b/reference/opt/shaders-hlsl/comp/inverse.comp index 3be954a6..f9ec89aa 100644 --- a/reference/opt/shaders-hlsl/comp/inverse.comp +++ b/reference/opt/shaders-hlsl/comp/inverse.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _15 : register(u0); ByteAddressBuffer _20 : register(t1); diff --git a/reference/opt/shaders-hlsl/comp/num-workgroups-alone.comp b/reference/opt/shaders-hlsl/comp/num-workgroups-alone.comp index dee39e3d..dc87dc84 100644 --- a/reference/opt/shaders-hlsl/comp/num-workgroups-alone.comp +++ b/reference/opt/shaders-hlsl/comp/num-workgroups-alone.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _10 : register(u0); cbuffer SPIRV_Cross_NumWorkgroups : register(b0) { diff --git a/reference/opt/shaders-hlsl/comp/num-workgroups-with-builtins.comp b/reference/opt/shaders-hlsl/comp/num-workgroups-with-builtins.comp index 1c98e5e5..2e2ad55f 100644 --- a/reference/opt/shaders-hlsl/comp/num-workgroups-with-builtins.comp +++ b/reference/opt/shaders-hlsl/comp/num-workgroups-with-builtins.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _10 : register(u0); cbuffer SPIRV_Cross_NumWorkgroups : register(b0) { diff --git a/reference/opt/shaders-hlsl/comp/outer-product.comp b/reference/opt/shaders-hlsl/comp/outer-product.comp index 71613d4f..e58c02fe 100644 --- a/reference/opt/shaders-hlsl/comp/outer-product.comp +++ b/reference/opt/shaders-hlsl/comp/outer-product.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _21 : register(u0); ByteAddressBuffer _26 : register(t1); diff --git a/reference/opt/shaders-hlsl/comp/rmw-matrix.comp b/reference/opt/shaders-hlsl/comp/rmw-matrix.comp index ed666693..30ac03f8 100644 --- a/reference/opt/shaders-hlsl/comp/rmw-matrix.comp +++ b/reference/opt/shaders-hlsl/comp/rmw-matrix.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _11 : register(u0); void comp_main() diff --git a/reference/opt/shaders-hlsl/comp/rwbuffer-matrix.comp b/reference/opt/shaders-hlsl/comp/rwbuffer-matrix.comp index 42103c2b..a6621c97 100644 --- a/reference/opt/shaders-hlsl/comp/rwbuffer-matrix.comp +++ b/reference/opt/shaders-hlsl/comp/rwbuffer-matrix.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _28 : register(u0); cbuffer UBO : register(b1) { @@ -8,57 +10,57 @@ cbuffer UBO : register(b1) void comp_main() { - float4x4 _253 = asfloat(uint4x4(_28.Load(64), _28.Load(80), _28.Load(96), _28.Load(112), _28.Load(68), _28.Load(84), _28.Load(100), _28.Load(116), _28.Load(72), _28.Load(88), _28.Load(104), _28.Load(120), _28.Load(76), _28.Load(92), _28.Load(108), _28.Load(124))); - _28.Store4(0, asuint(_253[0])); - _28.Store4(16, asuint(_253[1])); - _28.Store4(32, asuint(_253[2])); - _28.Store4(48, asuint(_253[3])); - float2x2 _256 = asfloat(uint2x2(_28.Load(144), _28.Load(152), _28.Load(148), _28.Load(156))); - _28.Store2(128, asuint(_256[0])); - _28.Store2(136, asuint(_256[1])); - float2x3 _259 = asfloat(uint2x3(_28.Load(192), _28.Load(200), _28.Load(208), _28.Load(196), _28.Load(204), _28.Load(212))); - _28.Store3(160, asuint(_259[0])); - _28.Store3(176, asuint(_259[1])); - float3x2 _262 = asfloat(uint3x2(_28.Load(240), _28.Load(256), _28.Load(244), _28.Load(260), _28.Load(248), _28.Load(264))); - _28.Store2(216, asuint(_262[0])); - _28.Store2(224, asuint(_262[1])); - _28.Store2(232, asuint(_262[2])); - float4x4 _265 = asfloat(uint4x4(_28.Load4(0), _28.Load4(16), _28.Load4(32), _28.Load4(48))); - _28.Store(64, asuint(_265[0].x)); - _28.Store(68, asuint(_265[1].x)); - _28.Store(72, asuint(_265[2].x)); - _28.Store(76, asuint(_265[3].x)); - _28.Store(80, asuint(_265[0].y)); - _28.Store(84, asuint(_265[1].y)); - _28.Store(88, asuint(_265[2].y)); - _28.Store(92, asuint(_265[3].y)); - _28.Store(96, asuint(_265[0].z)); - _28.Store(100, asuint(_265[1].z)); - _28.Store(104, asuint(_265[2].z)); - _28.Store(108, asuint(_265[3].z)); - _28.Store(112, asuint(_265[0].w)); - _28.Store(116, asuint(_265[1].w)); - _28.Store(120, asuint(_265[2].w)); - _28.Store(124, asuint(_265[3].w)); - float2x2 _268 = asfloat(uint2x2(_28.Load2(128), _28.Load2(136))); - _28.Store(144, asuint(_268[0].x)); - _28.Store(148, asuint(_268[1].x)); - _28.Store(152, asuint(_268[0].y)); - _28.Store(156, asuint(_268[1].y)); - float2x3 _271 = asfloat(uint2x3(_28.Load3(160), _28.Load3(176))); - _28.Store(192, asuint(_271[0].x)); - _28.Store(196, asuint(_271[1].x)); - _28.Store(200, asuint(_271[0].y)); - _28.Store(204, asuint(_271[1].y)); - _28.Store(208, asuint(_271[0].z)); - _28.Store(212, asuint(_271[1].z)); - float3x2 _274 = asfloat(uint3x2(_28.Load2(216), _28.Load2(224), _28.Load2(232))); - _28.Store(240, asuint(_274[0].x)); - _28.Store(244, asuint(_274[1].x)); - _28.Store(248, asuint(_274[2].x)); - _28.Store(256, asuint(_274[0].y)); - _28.Store(260, asuint(_274[1].y)); - _28.Store(264, asuint(_274[2].y)); + float4x4 _257 = asfloat(uint4x4(_28.Load(64), _28.Load(80), _28.Load(96), _28.Load(112), _28.Load(68), _28.Load(84), _28.Load(100), _28.Load(116), _28.Load(72), _28.Load(88), _28.Load(104), _28.Load(120), _28.Load(76), _28.Load(92), _28.Load(108), _28.Load(124))); + _28.Store4(0, asuint(_257[0])); + _28.Store4(16, asuint(_257[1])); + _28.Store4(32, asuint(_257[2])); + _28.Store4(48, asuint(_257[3])); + float2x2 _260 = asfloat(uint2x2(_28.Load(144), _28.Load(152), _28.Load(148), _28.Load(156))); + _28.Store2(128, asuint(_260[0])); + _28.Store2(136, asuint(_260[1])); + float2x3 _263 = asfloat(uint2x3(_28.Load(192), _28.Load(200), _28.Load(208), _28.Load(196), _28.Load(204), _28.Load(212))); + _28.Store3(160, asuint(_263[0])); + _28.Store3(176, asuint(_263[1])); + float3x2 _266 = asfloat(uint3x2(_28.Load(240), _28.Load(256), _28.Load(244), _28.Load(260), _28.Load(248), _28.Load(264))); + _28.Store2(216, asuint(_266[0])); + _28.Store2(224, asuint(_266[1])); + _28.Store2(232, asuint(_266[2])); + float4x4 _269 = asfloat(uint4x4(_28.Load4(0), _28.Load4(16), _28.Load4(32), _28.Load4(48))); + _28.Store(64, asuint(_269[0].x)); + _28.Store(68, asuint(_269[1].x)); + _28.Store(72, asuint(_269[2].x)); + _28.Store(76, asuint(_269[3].x)); + _28.Store(80, asuint(_269[0].y)); + _28.Store(84, asuint(_269[1].y)); + _28.Store(88, asuint(_269[2].y)); + _28.Store(92, asuint(_269[3].y)); + _28.Store(96, asuint(_269[0].z)); + _28.Store(100, asuint(_269[1].z)); + _28.Store(104, asuint(_269[2].z)); + _28.Store(108, asuint(_269[3].z)); + _28.Store(112, asuint(_269[0].w)); + _28.Store(116, asuint(_269[1].w)); + _28.Store(120, asuint(_269[2].w)); + _28.Store(124, asuint(_269[3].w)); + float2x2 _272 = asfloat(uint2x2(_28.Load2(128), _28.Load2(136))); + _28.Store(144, asuint(_272[0].x)); + _28.Store(148, asuint(_272[1].x)); + _28.Store(152, asuint(_272[0].y)); + _28.Store(156, asuint(_272[1].y)); + float2x3 _275 = asfloat(uint2x3(_28.Load3(160), _28.Load3(176))); + _28.Store(192, asuint(_275[0].x)); + _28.Store(196, asuint(_275[1].x)); + _28.Store(200, asuint(_275[0].y)); + _28.Store(204, asuint(_275[1].y)); + _28.Store(208, asuint(_275[0].z)); + _28.Store(212, asuint(_275[1].z)); + float3x2 _278 = asfloat(uint3x2(_28.Load2(216), _28.Load2(224), _28.Load2(232))); + _28.Store(240, asuint(_278[0].x)); + _28.Store(244, asuint(_278[1].x)); + _28.Store(248, asuint(_278[2].x)); + _28.Store(256, asuint(_278[0].y)); + _28.Store(260, asuint(_278[1].y)); + _28.Store(264, asuint(_278[2].y)); _28.Store(_68_index0 * 4 + _68_index1 * 16 + 64, asuint(1.0f)); _28.Store(_68_index0 * 4 + _68_index1 * 8 + 144, asuint(2.0f)); _28.Store(_68_index0 * 4 + _68_index1 * 8 + 192, asuint(3.0f)); diff --git a/reference/opt/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp b/reference/opt/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp index 47f2fe41..db2bbe96 100644 --- a/reference/opt/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp +++ b/reference/opt/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _9 : register(u0); void comp_main() diff --git a/reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp b/reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp index c4537db0..4e7c5e61 100644 --- a/reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp +++ b/reference/opt/shaders-hlsl/comp/spec-constant-op-member-array.comp @@ -28,6 +28,7 @@ static const int d = (c + 50); #define SPIRV_CROSS_CONSTANT_ID_3 400 #endif static const int e = SPIRV_CROSS_CONSTANT_ID_3; +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); RWByteAddressBuffer _22 : register(u0); diff --git a/reference/opt/shaders-hlsl/comp/ssbo-array-length.comp b/reference/opt/shaders-hlsl/comp/ssbo-array-length.comp index 2e3df626..82657cac 100644 --- a/reference/opt/shaders-hlsl/comp/ssbo-array-length.comp +++ b/reference/opt/shaders-hlsl/comp/ssbo-array-length.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _11 : register(u1); void comp_main() diff --git a/reference/opt/shaders-hlsl/comp/ssbo-array.comp b/reference/opt/shaders-hlsl/comp/ssbo-array.comp index d8bce8d5..ee202a22 100644 --- a/reference/opt/shaders-hlsl/comp/ssbo-array.comp +++ b/reference/opt/shaders-hlsl/comp/ssbo-array.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + void comp_main() { } diff --git a/reference/opt/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp b/reference/opt/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp index dabc7df9..d89e3abc 100644 --- a/reference/opt/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp +++ b/reference/opt/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _9 : register(u0, space0); static uint4 gl_SubgroupEqMask; diff --git a/reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp b/reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp index 59fc03a7..e57b2ea1 100644 --- a/reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp +++ b/reference/opt/shaders-msl/comp/access-private-workgroup-in-function.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0() { } diff --git a/reference/opt/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp b/reference/opt/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp index f7757cd1..18cfd68c 100644 --- a/reference/opt/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp +++ b/reference/opt/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp @@ -23,6 +23,8 @@ struct SSBO2 float4 v; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + struct spvDescriptorSetBuffer0 { const device SSBO0* ssbo0 [[id(0)]]; diff --git a/reference/opt/shaders-msl/comp/array-length.comp b/reference/opt/shaders-msl/comp/array-length.comp index 79358eb9..5a284b96 100644 --- a/reference/opt/shaders-msl/comp/array-length.comp +++ b/reference/opt/shaders-msl/comp/array-length.comp @@ -14,6 +14,8 @@ struct SSBO1 float bz[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device SSBO& _14 [[buffer(0)]], device SSBO1* ssbos_0 [[buffer(1)]], device SSBO1* ssbos_1 [[buffer(2)]]) { device SSBO1* ssbos[] = diff --git a/reference/opt/shaders-msl/comp/array-length.msl2.argument.discrete.comp b/reference/opt/shaders-msl/comp/array-length.msl2.argument.discrete.comp index 6ec9b11b..d804e187 100644 --- a/reference/opt/shaders-msl/comp/array-length.msl2.argument.discrete.comp +++ b/reference/opt/shaders-msl/comp/array-length.msl2.argument.discrete.comp @@ -25,6 +25,8 @@ struct SSBO3 float bz[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + struct spvDescriptorSetBuffer0 { device SSBO* m_16 [[id(0)]]; diff --git a/reference/opt/shaders-msl/comp/atomic.comp b/reference/opt/shaders-msl/comp/atomic.comp index 04721502..fca72bfc 100644 --- a/reference/opt/shaders-msl/comp/atomic.comp +++ b/reference/opt/shaders-msl/comp/atomic.comp @@ -12,6 +12,8 @@ struct SSBO int i32; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& ssbo [[buffer(0)]]) { threadgroup uint shared_u32; diff --git a/reference/opt/shaders-msl/comp/basic.comp b/reference/opt/shaders-msl/comp/basic.comp index e37d4fc5..dbb839f5 100644 --- a/reference/opt/shaders-msl/comp/basic.comp +++ b/reference/opt/shaders-msl/comp/basic.comp @@ -21,6 +21,8 @@ struct SSBO3 uint counter; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { float4 _29 = _23.in_data[gl_GlobalInvocationID.x]; diff --git a/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp b/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp index 8c3d2576..2d991f5d 100644 --- a/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp +++ b/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp @@ -21,9 +21,11 @@ struct SSBO3 uint counter; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(constant uint3& spvDispatchBase [[buffer(29)]], const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { - gl_GlobalInvocationID += spvDispatchBase * uint3(1, 1, 1); + gl_GlobalInvocationID += spvDispatchBase * gl_WorkGroupSize; float4 _33 = _27.in_data[gl_GlobalInvocationID.x]; if (dot(_33, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875) { diff --git a/reference/opt/shaders-msl/comp/basic.dynamic-buffer.msl2.comp b/reference/opt/shaders-msl/comp/basic.dynamic-buffer.msl2.invalid.comp similarity index 100% rename from reference/opt/shaders-msl/comp/basic.dynamic-buffer.msl2.comp rename to reference/opt/shaders-msl/comp/basic.dynamic-buffer.msl2.invalid.comp diff --git a/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp b/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp index ad9733a8..60d8d12c 100644 --- a/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp +++ b/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp @@ -13,6 +13,8 @@ struct SSBO1 int4 outputs[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO0& _25 [[buffer(0)]], device SSBO1& _39 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _39.outputs[gl_GlobalInvocationID.x].x = int(as_type(as_type(_25.inputs[gl_GlobalInvocationID.x].xy) + half2(half(1.0)))); diff --git a/reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp b/reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp index a4230b1e..9edb3423 100644 --- a/reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp +++ b/reference/opt/shaders-msl/comp/bitcast-16bit-2.invalid.comp @@ -18,6 +18,8 @@ struct UBO half4 const0; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO1& _21 [[buffer(0)]], device SSBO0& _29 [[buffer(1)]], constant UBO& _40 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { short2 _47 = as_type(_29.inputs[gl_GlobalInvocationID.x].x) + as_type(_40.const0.xy); diff --git a/reference/opt/shaders-msl/comp/coherent-block.comp b/reference/opt/shaders-msl/comp/coherent-block.comp index 580b9e3e..58bbacb7 100644 --- a/reference/opt/shaders-msl/comp/coherent-block.comp +++ b/reference/opt/shaders-msl/comp/coherent-block.comp @@ -8,6 +8,8 @@ struct SSBO float4 value; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(volatile device SSBO& _10 [[buffer(0)]]) { _10.value = float4(20.0); diff --git a/reference/opt/shaders-msl/comp/coherent-image.comp b/reference/opt/shaders-msl/comp/coherent-image.comp index c6af46b5..50904844 100644 --- a/reference/opt/shaders-msl/comp/coherent-image.comp +++ b/reference/opt/shaders-msl/comp/coherent-image.comp @@ -8,6 +8,8 @@ struct SSBO int4 value; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(volatile device SSBO& _10 [[buffer(0)]], texture2d uImage [[texture(0)]]) { _10.value = uImage.read(uint2(int2(10))); diff --git a/reference/opt/shaders-msl/comp/composite-construct.comp b/reference/opt/shaders-msl/comp/composite-construct.comp index c0189371..17dbb8fd 100644 --- a/reference/opt/shaders-msl/comp/composite-construct.comp +++ b/reference/opt/shaders-msl/comp/composite-construct.comp @@ -15,6 +15,8 @@ struct SSBO1 float4 bs[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + template inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A]) { diff --git a/reference/opt/shaders-msl/comp/copy-array-of-arrays.comp b/reference/opt/shaders-msl/comp/copy-array-of-arrays.comp index ea9693ce..cb396cff 100644 --- a/reference/opt/shaders-msl/comp/copy-array-of-arrays.comp +++ b/reference/opt/shaders-msl/comp/copy-array-of-arrays.comp @@ -10,6 +10,8 @@ struct BUF float c; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device BUF& o [[buffer(0)]]) { o.a = 4; diff --git a/reference/opt/shaders-msl/comp/defer-parens.comp b/reference/opt/shaders-msl/comp/defer-parens.comp index 69a8aab9..8c130e3a 100644 --- a/reference/opt/shaders-msl/comp/defer-parens.comp +++ b/reference/opt/shaders-msl/comp/defer-parens.comp @@ -9,6 +9,8 @@ struct SSBO int index; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _13 [[buffer(0)]]) { float4 _17 = _13.data; diff --git a/reference/opt/shaders-msl/comp/dowhile.comp b/reference/opt/shaders-msl/comp/dowhile.comp index 3ebafe0f..b503c948 100644 --- a/reference/opt/shaders-msl/comp/dowhile.comp +++ b/reference/opt/shaders-msl/comp/dowhile.comp @@ -14,21 +14,23 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { - float4 _57; - int _58; - _58 = 0; - _57 = _28.in_data[gl_GlobalInvocationID.x]; + float4 _59; + int _60; + _60 = 0; + _59 = _28.in_data[gl_GlobalInvocationID.x]; float4 _42; for (;;) { - _42 = _28.mvp * _57; - int _44 = _58 + 1; + _42 = _28.mvp * _59; + int _44 = _60 + 1; if (_44 < 16) { - _58 = _44; - _57 = _42; + _60 = _44; + _59 = _42; } else { diff --git a/reference/opt/shaders-msl/comp/image-cube-array-load-store.comp b/reference/opt/shaders-msl/comp/image-cube-array-load-store.comp index 1eeaf87c..41c4dfc1 100644 --- a/reference/opt/shaders-msl/comp/image-cube-array-load-store.comp +++ b/reference/opt/shaders-msl/comp/image-cube-array-load-store.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(texturecube_array uImageIn [[texture(0)]], texturecube_array uImageOut [[texture(1)]]) { uImageOut.write(uImageIn.read(uint2(int3(9, 7, 11).xy), uint(int3(9, 7, 11).z) % 6u, uint(int3(9, 7, 11).z) / 6u), uint2(int3(9, 7, 11).xy), uint(int3(9, 7, 11).z) % 6u, uint(int3(9, 7, 11).z) / 6u); diff --git a/reference/opt/shaders-msl/comp/image.comp b/reference/opt/shaders-msl/comp/image.comp index 447732dd..c875e78d 100644 --- a/reference/opt/shaders-msl/comp/image.comp +++ b/reference/opt/shaders-msl/comp/image.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(texture2d uImageIn [[texture(0)]], texture2d uImageOut [[texture(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { int2 _23 = int2(gl_GlobalInvocationID.xy); diff --git a/reference/opt/shaders-msl/comp/insert.comp b/reference/opt/shaders-msl/comp/insert.comp index 1418ce35..d7392c98 100644 --- a/reference/opt/shaders-msl/comp/insert.comp +++ b/reference/opt/shaders-msl/comp/insert.comp @@ -8,19 +8,21 @@ struct SSBO float4 out_data[1]; }; -constant float4 _52 = {}; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +constant float4 _53 = {}; kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { - float4 _45 = _52; - _45.x = 10.0; - float4 _47 = _45; - _47.y = 30.0; - float4 _49 = _47; - _49.z = 70.0; - float4 _51 = _49; - _51.w = 90.0; - _27.out_data[gl_GlobalInvocationID.x] = _51; + float4 _46 = _53; + _46.x = 10.0; + float4 _48 = _46; + _48.y = 30.0; + float4 _50 = _48; + _50.z = 70.0; + float4 _52 = _50; + _52.w = 90.0; + _27.out_data[gl_GlobalInvocationID.x] = _52; _27.out_data[gl_GlobalInvocationID.x].y = 20.0; } diff --git a/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp b/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp index 13304bd0..8a9beaf9 100644 --- a/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp +++ b/reference/opt/shaders-msl/comp/int64.invalid.msl22.comp @@ -9,16 +9,18 @@ struct SSBO uint u32; }; -constant long _162 = {}; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +constant long _164 = {}; kernel void main0(device SSBO& _96 [[buffer(0)]]) { - long4 _137; - ulong4 _141; - _137 = abs((_137 + long4(30l, 40l, 50l, 60l)) + long4(_141 + ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul))); - _141 += ulong4(long4(999999999999999999l, 8888888888888888l, 77777777777777777l, 6666666666666666l)); - ulong _109 = ulong(_162); - _96.s32 = int(uint(((ulong(_137.x) + _141.y) + _109) + _109)); - _96.u32 = uint(((ulong(_137.y) + _141.z) + ulong(_162 + 1l)) + _109); + long4 _139; + ulong4 _143; + _139 = abs((_139 + long4(30l, 40l, 50l, 60l)) + long4(_143 + ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul))); + _143 += ulong4(long4(999999999999999999l, 8888888888888888l, 77777777777777777l, 6666666666666666l)); + ulong _109 = ulong(_164); + _96.s32 = int(uint(((ulong(_139.x) + _143.y) + _109) + _109)); + _96.u32 = uint(((ulong(_139.y) + _143.z) + ulong(_164 + 1l)) + _109); } diff --git a/reference/opt/shaders-msl/comp/inverse.comp b/reference/opt/shaders-msl/comp/inverse.comp index 3ea692fa..a4ef60c4 100644 --- a/reference/opt/shaders-msl/comp/inverse.comp +++ b/reference/opt/shaders-msl/comp/inverse.comp @@ -19,6 +19,8 @@ struct MatrixIn float4x4 m4in; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + // Returns the determinant of a 2x2 matrix. inline float spvDet2x2(float a1, float a2, float b1, float b2) { diff --git a/reference/opt/shaders-msl/comp/mat3.comp b/reference/opt/shaders-msl/comp/mat3.comp index 72f08dd8..31351ba5 100644 --- a/reference/opt/shaders-msl/comp/mat3.comp +++ b/reference/opt/shaders-msl/comp/mat3.comp @@ -8,6 +8,8 @@ struct SSBO2 float3x3 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _22 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _22.out_data[gl_GlobalInvocationID.x] = float3x3(float3(10.0), float3(20.0), float3(40.0)); diff --git a/reference/opt/shaders-msl/comp/mod.comp b/reference/opt/shaders-msl/comp/mod.comp index e3799571..e8c01f96 100644 --- a/reference/opt/shaders-msl/comp/mod.comp +++ b/reference/opt/shaders-msl/comp/mod.comp @@ -15,6 +15,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + // Implementation of the GLSL mod() function, which is slightly different than Metal fmod() template inline Tx mod(Tx x, Ty y) diff --git a/reference/opt/shaders-msl/comp/modf.comp b/reference/opt/shaders-msl/comp/modf.comp index 39e40233..df19cae5 100644 --- a/reference/opt/shaders-msl/comp/modf.comp +++ b/reference/opt/shaders-msl/comp/modf.comp @@ -13,6 +13,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { float4 i; diff --git a/reference/opt/shaders-msl/comp/outer-product.comp b/reference/opt/shaders-msl/comp/outer-product.comp index 8e32db39..e589642d 100644 --- a/reference/opt/shaders-msl/comp/outer-product.comp +++ b/reference/opt/shaders-msl/comp/outer-product.comp @@ -23,6 +23,8 @@ struct ReadSSBO float4 v4; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _21 [[buffer(0)]], const device ReadSSBO& _26 [[buffer(1)]]) { _21.m22 = float2x2(_26.v2 * _26.v2.x, _26.v2 * _26.v2.y); diff --git a/reference/opt/shaders-msl/comp/read-write-only.comp b/reference/opt/shaders-msl/comp/read-write-only.comp index 7547b417..0cf8d8e3 100644 --- a/reference/opt/shaders-msl/comp/read-write-only.comp +++ b/reference/opt/shaders-msl/comp/read-write-only.comp @@ -21,6 +21,8 @@ struct SSBO1 float4 data3; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _10 [[buffer(0)]], const device SSBO0& _15 [[buffer(1)]], device SSBO1& _21 [[buffer(2)]]) { _10.data4 = _15.data0 + _21.data2; diff --git a/reference/opt/shaders-msl/comp/rmw-matrix.comp b/reference/opt/shaders-msl/comp/rmw-matrix.comp index 150db7ed..b53a3a75 100644 --- a/reference/opt/shaders-msl/comp/rmw-matrix.comp +++ b/reference/opt/shaders-msl/comp/rmw-matrix.comp @@ -13,6 +13,8 @@ struct SSBO float4x4 c1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _11 [[buffer(0)]]) { _11.a *= _11.a1; diff --git a/reference/opt/shaders-msl/comp/rmw-opt.comp b/reference/opt/shaders-msl/comp/rmw-opt.comp index 05e1f6f2..f93967da 100644 --- a/reference/opt/shaders-msl/comp/rmw-opt.comp +++ b/reference/opt/shaders-msl/comp/rmw-opt.comp @@ -8,6 +8,8 @@ struct SSBO int a; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _9 [[buffer(0)]]) { _9.a += 10; diff --git a/reference/opt/shaders-msl/comp/scalar-std450-distance-length-normalize.comp b/reference/opt/shaders-msl/comp/scalar-std450-distance-length-normalize.comp index 312a6f94..0ae6e556 100644 --- a/reference/opt/shaders-msl/comp/scalar-std450-distance-length-normalize.comp +++ b/reference/opt/shaders-msl/comp/scalar-std450-distance-length-normalize.comp @@ -12,6 +12,8 @@ struct SSBO float e; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _9 [[buffer(0)]]) { _9.c = abs(_9.a - _9.b); diff --git a/reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp b/reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp index d3c8b7dc..8f54f052 100644 --- a/reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp +++ b/reference/opt/shaders-msl/comp/spec-constant-op-member-array.comp @@ -40,6 +40,7 @@ struct SSBO constant int e_tmp [[function_constant(3)]]; constant int e = is_function_constant_defined(e_tmp) ? e_tmp : 400; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); kernel void main0(device SSBO& _22 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { diff --git a/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp b/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp index 905222d3..24ce2807 100644 --- a/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp +++ b/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp @@ -16,32 +16,34 @@ struct SSBO Sub sub[2]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { - float _153[2]; - _153[0] = _27.sub[gl_WorkGroupID.x].f[0].x; - _153[1] = _27.sub[gl_WorkGroupID.x].f[1].x; - float2 _154[2]; - _154[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy; - _154[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy; - float3 _155[2]; - _155[0] = _27.sub[gl_WorkGroupID.x].f3[0]; - _155[1] = _27.sub[gl_WorkGroupID.x].f3[1]; - float4 _156[2]; - _156[0] = _27.sub[gl_WorkGroupID.x].f4[0]; - _156[1] = _27.sub[gl_WorkGroupID.x].f4[1]; - _153[gl_GlobalInvocationID.x] += 1.0; - _154[gl_GlobalInvocationID.x] += float2(2.0); - _155[gl_GlobalInvocationID.x] += float3(3.0); - _156[gl_GlobalInvocationID.x] += float4(4.0); - _27.sub[gl_WorkGroupID.x].f[0].x = _153[0]; - _27.sub[gl_WorkGroupID.x].f[1].x = _153[1]; - _27.sub[gl_WorkGroupID.x].f2[0].xy = _154[0]; - _27.sub[gl_WorkGroupID.x].f2[1].xy = _154[1]; - _27.sub[gl_WorkGroupID.x].f3[0] = _155[0]; - _27.sub[gl_WorkGroupID.x].f3[1] = _155[1]; - _27.sub[gl_WorkGroupID.x].f4[0] = _156[0]; - _27.sub[gl_WorkGroupID.x].f4[1] = _156[1]; + float _155[2]; + _155[0] = _27.sub[gl_WorkGroupID.x].f[0].x; + _155[1] = _27.sub[gl_WorkGroupID.x].f[1].x; + float2 _156[2]; + _156[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy; + _156[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy; + float3 _157[2]; + _157[0] = _27.sub[gl_WorkGroupID.x].f3[0]; + _157[1] = _27.sub[gl_WorkGroupID.x].f3[1]; + float4 _158[2]; + _158[0] = _27.sub[gl_WorkGroupID.x].f4[0]; + _158[1] = _27.sub[gl_WorkGroupID.x].f4[1]; + _155[gl_GlobalInvocationID.x] += 1.0; + _156[gl_GlobalInvocationID.x] += float2(2.0); + _157[gl_GlobalInvocationID.x] += float3(3.0); + _158[gl_GlobalInvocationID.x] += float4(4.0); + _27.sub[gl_WorkGroupID.x].f[0].x = _155[0]; + _27.sub[gl_WorkGroupID.x].f[1].x = _155[1]; + _27.sub[gl_WorkGroupID.x].f2[0].xy = _156[0]; + _27.sub[gl_WorkGroupID.x].f2[1].xy = _156[1]; + _27.sub[gl_WorkGroupID.x].f3[0] = _157[0]; + _27.sub[gl_WorkGroupID.x].f3[1] = _157[1]; + _27.sub[gl_WorkGroupID.x].f4[0] = _158[0]; + _27.sub[gl_WorkGroupID.x].f4[1] = _158[1]; _27.sub[0].f[0].x += 5.0; _27.sub[0].f2[1].xy += float2(5.0); } diff --git a/reference/opt/shaders-msl/comp/struct-layout.comp b/reference/opt/shaders-msl/comp/struct-layout.comp index d4413625..0445f5ae 100644 --- a/reference/opt/shaders-msl/comp/struct-layout.comp +++ b/reference/opt/shaders-msl/comp/struct-layout.comp @@ -18,6 +18,8 @@ struct SSBO Foo in_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _23 [[buffer(0)]], const device SSBO& _30 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _23.out_data[gl_GlobalInvocationID.x].m = _30.in_data[gl_GlobalInvocationID.x].m * _30.in_data[gl_GlobalInvocationID.x].m; diff --git a/reference/opt/shaders-msl/comp/struct-nested.comp b/reference/opt/shaders-msl/comp/struct-nested.comp index 6a1419ce..ad706c59 100644 --- a/reference/opt/shaders-msl/comp/struct-nested.comp +++ b/reference/opt/shaders-msl/comp/struct-nested.comp @@ -18,6 +18,8 @@ struct dstbuffer s2 test[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device dstbuffer& _19 [[buffer(0)]]) { _19.test[0].b.a = 0; diff --git a/reference/opt/shaders-msl/comp/struct-packing.comp b/reference/opt/shaders-msl/comp/struct-packing.comp index ad22c257..775bb348 100644 --- a/reference/opt/shaders-msl/comp/struct-packing.comp +++ b/reference/opt/shaders-msl/comp/struct-packing.comp @@ -118,6 +118,8 @@ struct SSBO0 float4 array[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]]) { Content_1 _60 = ssbo_140.content; diff --git a/reference/opt/shaders-msl/comp/torture-loop.comp b/reference/opt/shaders-msl/comp/torture-loop.comp index 4c367d3e..ff7e02e2 100644 --- a/reference/opt/shaders-msl/comp/torture-loop.comp +++ b/reference/opt/shaders-msl/comp/torture-loop.comp @@ -14,29 +14,31 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { - float4 _99; - _99 = _24.in_data[gl_GlobalInvocationID.x]; - for (int _93 = 0; (_93 + 1) < 10; ) + float4 _101; + _101 = _24.in_data[gl_GlobalInvocationID.x]; + for (int _95 = 0; (_95 + 1) < 10; ) { - _99 *= 2.0; - _93 += 2; + _101 *= 2.0; + _95 += 2; continue; } - float4 _98; - _98 = _99; - float4 _103; - for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++) + float4 _100; + _100 = _101; + float4 _105; + for (uint _96 = 0u; _96 < 16u; _100 = _105, _96++) { - _103 = _98; - for (uint _100 = 0u; _100 < 30u; ) + _105 = _100; + for (uint _102 = 0u; _102 < 30u; ) { - _103 = _24.mvp * _103; - _100++; + _105 = _24.mvp * _105; + _102++; continue; } } - _89.out_data[gl_GlobalInvocationID.x] = _98; + _89.out_data[gl_GlobalInvocationID.x] = _100; } diff --git a/reference/opt/shaders-msl/comp/type-alias.comp b/reference/opt/shaders-msl/comp/type-alias.comp index 8a68933d..2f6a0b7b 100644 --- a/reference/opt/shaders-msl/comp/type-alias.comp +++ b/reference/opt/shaders-msl/comp/type-alias.comp @@ -28,6 +28,8 @@ struct SSBO2 float4 outputs[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO0& _36 [[buffer(0)]], device SSBO1& _55 [[buffer(1)]], device SSBO2& _66 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _66.outputs[gl_GlobalInvocationID.x] = _36.s0s[gl_GlobalInvocationID.x].a + _55.s1s[gl_GlobalInvocationID.x].a; diff --git a/reference/opt/shaders-msl/comp/udiv.comp b/reference/opt/shaders-msl/comp/udiv.comp index 32874ad7..7f7315b8 100644 --- a/reference/opt/shaders-msl/comp/udiv.comp +++ b/reference/opt/shaders-msl/comp/udiv.comp @@ -13,6 +13,8 @@ struct SSBO uint inputs[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _10 [[buffer(0)]], device SSBO& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _10.outputs[gl_GlobalInvocationID.x] = _23.inputs[gl_GlobalInvocationID.x] / 29u; diff --git a/reference/opt/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp b/reference/opt/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp index a37fe519..cea12980 100644 --- a/reference/opt/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp +++ b/reference/opt/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp @@ -91,6 +91,8 @@ struct ResType_7 int4 _m1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBOUint& u [[buffer(0)]], device SSBOInt& i [[buffer(1)]]) { ResType _25; diff --git a/reference/opt/shaders-msl/frag/shader-arithmetic-8bit.frag b/reference/opt/shaders-msl/frag/shader-arithmetic-8bit.frag index 30b28d21..060e05e9 100644 --- a/reference/opt/shaders-msl/frag/shader-arithmetic-8bit.frag +++ b/reference/opt/shaders-msl/frag/shader-arithmetic-8bit.frag @@ -37,9 +37,10 @@ fragment main0_out main0(main0_in in [[stage_in]], device SSBO& ssbo [[buffer(0) main0_out out = {}; short _196 = 10; int _197 = 20; - char2 _198 = as_type(_196); - char4 _199 = as_type(_197); - _196 = as_type(_198); + char2 _201 = as_type(10); + char2 _198 = _201; + char4 _199 = as_type(20); + _196 = as_type(_201); _197 = as_type(_199); ssbo.i8[0] = _199.x; ssbo.i8[1] = _199.y; @@ -47,9 +48,10 @@ fragment main0_out main0(main0_in in [[stage_in]], device SSBO& ssbo [[buffer(0) ssbo.i8[3] = _199.w; ushort _220 = 10u; uint _221 = 20u; - uchar2 _222 = as_type(_220); - uchar4 _223 = as_type(_221); - _220 = as_type(_222); + uchar2 _225 = as_type(10u); + uchar2 _222 = _225; + uchar4 _223 = as_type(20u); + _220 = as_type(_225); _221 = as_type(_223); ssbo.u8[0] = _223.x; ssbo.u8[1] = _223.y; @@ -57,21 +59,34 @@ fragment main0_out main0(main0_in in [[stage_in]], device SSBO& ssbo [[buffer(0) ssbo.u8[3] = _223.w; char4 _246 = char4(in.vColor); char4 _244 = _246; - _244 += char4(registers.i8); - _244 += char4(-40); - _244 += char4(-50); - _244 += char4(char(10), char(20), char(30), char(40)); - _244 += char4(ssbo.i8[4]); - _244 += char4(ubo.i8); - out.FragColorInt = int4(_244); - uchar4 _271 = uchar4(_246); - _271 += uchar4(registers.u8); - _271 += uchar4(216); - _271 += uchar4(206); - _271 += uchar4(uchar(10), uchar(20), uchar(30), uchar(40)); - _271 += uchar4(ssbo.u8[4]); - _271 += uchar4(ubo.u8); - out.FragColorUint = uint4(_271); + char4 _251 = _246 + char4(registers.i8); + _244 = _251; + char4 _254 = _251 + char4(-40); + _244 = _254; + char4 _256 = _254 + char4(-50); + _244 = _256; + char4 _258 = _256 + char4(char(10), char(20), char(30), char(40)); + _244 = _258; + char4 _263 = _258 + char4(ssbo.i8[4]); + _244 = _263; + char4 _268 = _263 + char4(ubo.i8); + _244 = _268; + out.FragColorInt = int4(_268); + uchar4 _274 = uchar4(_246); + uchar4 _271 = _274; + uchar4 _279 = _274 + uchar4(registers.u8); + _271 = _279; + uchar4 _282 = _279 + uchar4(216); + _271 = _282; + uchar4 _284 = _282 + uchar4(206); + _271 = _284; + uchar4 _286 = _284 + uchar4(uchar(10), uchar(20), uchar(30), uchar(40)); + _271 = _286; + uchar4 _291 = _286 + uchar4(ssbo.u8[4]); + _271 = _291; + uchar4 _296 = _291 + uchar4(ubo.u8); + _271 = _296; + out.FragColorUint = uint4(_296); return out; } diff --git a/reference/opt/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp b/reference/opt/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp index 2f731356..4f341227 100644 --- a/reference/opt/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp +++ b/reference/opt/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp @@ -121,6 +121,8 @@ struct SSBO2 packed_rm_float3x2 m2; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]]) { ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0].xy; diff --git a/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp b/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp index 948806db..71bce630 100644 --- a/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp +++ b/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp @@ -10,6 +10,8 @@ struct SSBO float FragColor; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline uint4 spvSubgroupBallot(bool value) { simd_vote vote = simd_ballot(value); diff --git a/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp b/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp index 6d32de69..e5c83eb6 100644 --- a/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp +++ b/reference/opt/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp @@ -8,6 +8,8 @@ struct SSBO float FragColor; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[quadgroups_per_threadgroup]], uint gl_SubgroupID [[quadgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_quadgroup]]) { _9.FragColor = float(gl_NumSubgroups); diff --git a/reference/opt/shaders-msl/vulkan/frag/demote-to-helper.vk.nocompat.msl21.invalid.frag b/reference/opt/shaders-msl/vulkan/frag/demote-to-helper.vk.nocompat.msl21.invalid.frag index 429da476..0e0348bf 100644 --- a/reference/opt/shaders-msl/vulkan/frag/demote-to-helper.vk.nocompat.msl21.invalid.frag +++ b/reference/opt/shaders-msl/vulkan/frag/demote-to-helper.vk.nocompat.msl21.invalid.frag @@ -6,5 +6,6 @@ using namespace metal; fragment void main0() { bool _9 = simd_is_helper_thread(); + bool helper = _9; } diff --git a/reference/opt/shaders/comp/cfg.comp b/reference/opt/shaders/comp/cfg.comp index 0b7e0c16..bc1a4178 100644 --- a/reference/opt/shaders/comp/cfg.comp +++ b/reference/opt/shaders/comp/cfg.comp @@ -6,7 +6,7 @@ layout(binding = 0, std430) buffer SSBO float data; } _11; -float _183; +float _187; void main() { @@ -31,14 +31,14 @@ void main() break; } } - float _180; - _180 = _183; - for (int _179 = 0; _179 < 20; ) + float _184; + _184 = _187; + for (int _183 = 0; _183 < 20; ) { - _180 += 10.0; - _179++; + _184 += 10.0; + _183++; continue; } - _11.data = _180; + _11.data = _184; } diff --git a/reference/opt/shaders/comp/dowhile.comp b/reference/opt/shaders/comp/dowhile.comp index 4370ea30..d9a9f77c 100644 --- a/reference/opt/shaders/comp/dowhile.comp +++ b/reference/opt/shaders/comp/dowhile.comp @@ -14,19 +14,19 @@ layout(binding = 1, std430) writeonly buffer SSBO2 void main() { - vec4 _57; - int _58; - _58 = 0; - _57 = _28.in_data[gl_GlobalInvocationID.x]; + vec4 _59; + int _60; + _60 = 0; + _59 = _28.in_data[gl_GlobalInvocationID.x]; vec4 _42; for (;;) { - _42 = _28.mvp * _57; - int _44 = _58 + 1; + _42 = _28.mvp * _59; + int _44 = _60 + 1; if (_44 < 16) { - _58 = _44; - _57 = _42; + _60 = _44; + _59 = _42; } else { diff --git a/reference/opt/shaders/comp/insert.comp b/reference/opt/shaders/comp/insert.comp index 5ff71944..ad394aab 100644 --- a/reference/opt/shaders/comp/insert.comp +++ b/reference/opt/shaders/comp/insert.comp @@ -6,19 +6,19 @@ layout(binding = 0, std430) writeonly buffer SSBO vec4 out_data[]; } _27; -vec4 _52; +vec4 _53; void main() { - vec4 _45 = _52; - _45.x = 10.0; - vec4 _47 = _45; - _47.y = 30.0; - vec4 _49 = _47; - _49.z = 70.0; - vec4 _51 = _49; - _51.w = 90.0; - _27.out_data[gl_GlobalInvocationID.x] = _51; + vec4 _46 = _53; + _46.x = 10.0; + vec4 _48 = _46; + _48.y = 30.0; + vec4 _50 = _48; + _50.z = 70.0; + vec4 _52 = _50; + _52.w = 90.0; + _27.out_data[gl_GlobalInvocationID.x] = _52; _27.out_data[gl_GlobalInvocationID.x].y = 20.0; } diff --git a/reference/opt/shaders/comp/torture-loop.comp b/reference/opt/shaders/comp/torture-loop.comp index 5943966c..9ca2b959 100644 --- a/reference/opt/shaders/comp/torture-loop.comp +++ b/reference/opt/shaders/comp/torture-loop.comp @@ -14,27 +14,27 @@ layout(binding = 1, std430) writeonly buffer SSBO2 void main() { - vec4 _99; - _99 = _24.in_data[gl_GlobalInvocationID.x]; - for (int _93 = 0; (_93 + 1) < 10; ) + vec4 _101; + _101 = _24.in_data[gl_GlobalInvocationID.x]; + for (int _95 = 0; (_95 + 1) < 10; ) { - _99 *= 2.0; - _93 += 2; + _101 *= 2.0; + _95 += 2; continue; } - vec4 _98; - _98 = _99; - vec4 _103; - for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++) + vec4 _100; + _100 = _101; + vec4 _105; + for (uint _96 = 0u; _96 < 16u; _100 = _105, _96++) { - _103 = _98; - for (uint _100 = 0u; _100 < 30u; ) + _105 = _100; + for (uint _102 = 0u; _102 < 30u; ) { - _103 = _24.mvp * _103; - _100++; + _105 = _24.mvp * _105; + _102++; continue; } } - _89.out_data[gl_GlobalInvocationID.x] = _98; + _89.out_data[gl_GlobalInvocationID.x] = _100; } diff --git a/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk b/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk index dfcaac83..31283196 100644 --- a/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk +++ b/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk @@ -19,27 +19,17 @@ layout(set = 0, binding = 0, std430) restrict buffer LinkedList void main() { - Node _45; + Node _112; if (gl_WorkGroupID.x < 4u) { - _45 = _50.head1; + _112 = _50.head1; } else { - _45 = _50.head2; + _112 = _50.head2; } - restrict Node n = _45; - Node param = n.next; - Node param_1 = _50.head1; - Node param_2 = _50.head2; - param.value = param_1.value + param_2.value; - Node param_4 = _50.head1; - Node param_3 = param_4; - n = param_3; - int v = _50.head2.value; - n.value = 20; - n.value = v * 10; - uint64_t uptr = uint64_t(_50.head2.next); - Node unode = Node(uptr); + _112.next.value = _50.head1.value + _50.head2.value; + _50.head1.value = 20; + _50.head1.value = _50.head2.value * 10; } diff --git a/reference/opt/shaders/vulkan/frag/shader-arithmetic-8bit.nocompat.vk.frag.vk b/reference/opt/shaders/vulkan/frag/shader-arithmetic-8bit.nocompat.vk.frag.vk index d09930f3..69558dee 100644 --- a/reference/opt/shaders/vulkan/frag/shader-arithmetic-8bit.nocompat.vk.frag.vk +++ b/reference/opt/shaders/vulkan/frag/shader-arithmetic-8bit.nocompat.vk.frag.vk @@ -30,9 +30,10 @@ void main() { int16_t _196 = 10s; int _197 = 20; - i8vec2 _198 = unpack8(_196); - i8vec4 _199 = unpack8(_197); - _196 = pack16(_198); + i8vec2 _201 = unpack8(10s); + i8vec2 _198 = _201; + i8vec4 _199 = unpack8(20); + _196 = pack16(_201); _197 = pack32(_199); ssbo.i8[0] = _199.x; ssbo.i8[1] = _199.y; @@ -40,9 +41,10 @@ void main() ssbo.i8[3] = _199.w; uint16_t _220 = 10us; uint _221 = 20u; - u8vec2 _222 = unpack8(_220); - u8vec4 _223 = unpack8(_221); - _220 = pack16(_222); + u8vec2 _225 = unpack8(10us); + u8vec2 _222 = _225; + u8vec4 _223 = unpack8(20u); + _220 = pack16(_225); _221 = pack32(_223); ssbo.u8[0] = _223.x; ssbo.u8[1] = _223.y; @@ -50,20 +52,33 @@ void main() ssbo.u8[3] = _223.w; i8vec4 _246 = i8vec4(vColor); i8vec4 _244 = _246; - _244 += i8vec4(registers.i8); - _244 += i8vec4(-40); - _244 += i8vec4(-50); - _244 += i8vec4(int8_t(10), int8_t(20), int8_t(30), int8_t(40)); - _244 += i8vec4(ssbo.i8[4]); - _244 += i8vec4(ubo.i8); - FragColorInt = ivec4(_244); - u8vec4 _271 = u8vec4(_246); - _271 += u8vec4(registers.u8); - _271 += u8vec4(216); - _271 += u8vec4(206); - _271 += u8vec4(uint8_t(10), uint8_t(20), uint8_t(30), uint8_t(40)); - _271 += u8vec4(ssbo.u8[4]); - _271 += u8vec4(ubo.u8); - FragColorUint = uvec4(_271); + i8vec4 _251 = _246 + i8vec4(registers.i8); + _244 = _251; + i8vec4 _254 = _251 + i8vec4(-40); + _244 = _254; + i8vec4 _256 = _254 + i8vec4(-50); + _244 = _256; + i8vec4 _258 = _256 + i8vec4(int8_t(10), int8_t(20), int8_t(30), int8_t(40)); + _244 = _258; + i8vec4 _263 = _258 + i8vec4(ssbo.i8[4]); + _244 = _263; + i8vec4 _268 = _263 + i8vec4(ubo.i8); + _244 = _268; + FragColorInt = ivec4(_268); + u8vec4 _274 = u8vec4(_246); + u8vec4 _271 = _274; + u8vec4 _279 = _274 + u8vec4(registers.u8); + _271 = _279; + u8vec4 _282 = _279 + u8vec4(216); + _271 = _282; + u8vec4 _284 = _282 + u8vec4(206); + _271 = _284; + u8vec4 _286 = _284 + u8vec4(uint8_t(10), uint8_t(20), uint8_t(30), uint8_t(40)); + _271 = _286; + u8vec4 _291 = _286 + u8vec4(ssbo.u8[4]); + _271 = _291; + u8vec4 _296 = _291 + u8vec4(ubo.u8); + _271 = _296; + FragColorUint = uvec4(_296); } diff --git a/reference/shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag b/reference/shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag similarity index 100% rename from reference/shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag rename to reference/shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag diff --git a/reference/shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp b/reference/shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp index d697ed4e..7558afaa 100644 --- a/reference/shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp +++ b/reference/shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp @@ -4,6 +4,8 @@ struct ResType int _m1; }; +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _19 : register(u0); uint SPIRV_Cross_packHalf2x16(float2 value) diff --git a/reference/shaders-hlsl/comp/access-chains.comp b/reference/shaders-hlsl/comp/access-chains.comp index 924e9191..c748200b 100644 --- a/reference/shaders-hlsl/comp/access-chains.comp +++ b/reference/shaders-hlsl/comp/access-chains.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer wo : register(u1); ByteAddressBuffer ro : register(t0); diff --git a/reference/shaders-hlsl/comp/address-buffers.comp b/reference/shaders-hlsl/comp/address-buffers.comp index a252fc8a..7f1c7975 100644 --- a/reference/shaders-hlsl/comp/address-buffers.comp +++ b/reference/shaders-hlsl/comp/address-buffers.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer WriteOnly : register(u2); ByteAddressBuffer ReadOnly : register(t0); RWByteAddressBuffer ReadWrite : register(u1); diff --git a/reference/shaders-hlsl/comp/atomic.comp b/reference/shaders-hlsl/comp/atomic.comp index 72e15bf7..e6ff891e 100644 --- a/reference/shaders-hlsl/comp/atomic.comp +++ b/reference/shaders-hlsl/comp/atomic.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer ssbo : register(u2); RWTexture2D uImage : register(u0); RWTexture2D iImage : register(u1); diff --git a/reference/shaders-hlsl/comp/globallycoherent.comp b/reference/shaders-hlsl/comp/globallycoherent.comp index 69886256..236f341e 100644 --- a/reference/shaders-hlsl/comp/globallycoherent.comp +++ b/reference/shaders-hlsl/comp/globallycoherent.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + globallycoherent RWByteAddressBuffer _29 : register(u3); ByteAddressBuffer _33 : register(t2); RWTexture2D uImageIn : register(u0); diff --git a/reference/shaders-hlsl/comp/image.comp b/reference/shaders-hlsl/comp/image.comp index c8504e63..89a99409 100644 --- a/reference/shaders-hlsl/comp/image.comp +++ b/reference/shaders-hlsl/comp/image.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWTexture2D uImageInF : register(u0); RWTexture2D uImageOutF : register(u1); RWTexture2D uImageInI : register(u2); diff --git a/reference/shaders-hlsl/comp/inverse.comp b/reference/shaders-hlsl/comp/inverse.comp index 3be954a6..f9ec89aa 100644 --- a/reference/shaders-hlsl/comp/inverse.comp +++ b/reference/shaders-hlsl/comp/inverse.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _15 : register(u0); ByteAddressBuffer _20 : register(t1); diff --git a/reference/shaders-hlsl/comp/num-workgroups-alone.comp b/reference/shaders-hlsl/comp/num-workgroups-alone.comp index dee39e3d..dc87dc84 100644 --- a/reference/shaders-hlsl/comp/num-workgroups-alone.comp +++ b/reference/shaders-hlsl/comp/num-workgroups-alone.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _10 : register(u0); cbuffer SPIRV_Cross_NumWorkgroups : register(b0) { diff --git a/reference/shaders-hlsl/comp/num-workgroups-with-builtins.comp b/reference/shaders-hlsl/comp/num-workgroups-with-builtins.comp index 1c98e5e5..2e2ad55f 100644 --- a/reference/shaders-hlsl/comp/num-workgroups-with-builtins.comp +++ b/reference/shaders-hlsl/comp/num-workgroups-with-builtins.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _10 : register(u0); cbuffer SPIRV_Cross_NumWorkgroups : register(b0) { diff --git a/reference/shaders-hlsl/comp/outer-product.comp b/reference/shaders-hlsl/comp/outer-product.comp index 71613d4f..e58c02fe 100644 --- a/reference/shaders-hlsl/comp/outer-product.comp +++ b/reference/shaders-hlsl/comp/outer-product.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _21 : register(u0); ByteAddressBuffer _26 : register(t1); diff --git a/reference/shaders-hlsl/comp/rmw-matrix.comp b/reference/shaders-hlsl/comp/rmw-matrix.comp index ed666693..30ac03f8 100644 --- a/reference/shaders-hlsl/comp/rmw-matrix.comp +++ b/reference/shaders-hlsl/comp/rmw-matrix.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _11 : register(u0); void comp_main() diff --git a/reference/shaders-hlsl/comp/rwbuffer-matrix.comp b/reference/shaders-hlsl/comp/rwbuffer-matrix.comp index e7982928..197c9a95 100644 --- a/reference/shaders-hlsl/comp/rwbuffer-matrix.comp +++ b/reference/shaders-hlsl/comp/rwbuffer-matrix.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _28 : register(u0); cbuffer UBO : register(b1) { diff --git a/reference/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp b/reference/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp index 47f2fe41..db2bbe96 100644 --- a/reference/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp +++ b/reference/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _9 : register(u0); void comp_main() diff --git a/reference/shaders-hlsl/comp/spec-constant-op-member-array.comp b/reference/shaders-hlsl/comp/spec-constant-op-member-array.comp index c4537db0..4e7c5e61 100644 --- a/reference/shaders-hlsl/comp/spec-constant-op-member-array.comp +++ b/reference/shaders-hlsl/comp/spec-constant-op-member-array.comp @@ -28,6 +28,7 @@ static const int d = (c + 50); #define SPIRV_CROSS_CONSTANT_ID_3 400 #endif static const int e = SPIRV_CROSS_CONSTANT_ID_3; +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); RWByteAddressBuffer _22 : register(u0); diff --git a/reference/shaders-hlsl/comp/ssbo-array-length.comp b/reference/shaders-hlsl/comp/ssbo-array-length.comp index 2e3df626..82657cac 100644 --- a/reference/shaders-hlsl/comp/ssbo-array-length.comp +++ b/reference/shaders-hlsl/comp/ssbo-array-length.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _11 : register(u1); void comp_main() diff --git a/reference/shaders-hlsl/comp/ssbo-array.comp b/reference/shaders-hlsl/comp/ssbo-array.comp index 90927421..dab20325 100644 --- a/reference/shaders-hlsl/comp/ssbo-array.comp +++ b/reference/shaders-hlsl/comp/ssbo-array.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer ssbo0 : register(u0); void comp_main() diff --git a/reference/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp b/reference/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp index b87574f1..0957a708 100644 --- a/reference/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp +++ b/reference/shaders-hlsl/comp/subgroups.invalid.nofxc.sm60.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + RWByteAddressBuffer _9 : register(u0, space0); static uint4 gl_SubgroupEqMask; diff --git a/reference/shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag b/reference/shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag similarity index 100% rename from reference/shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag rename to reference/shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag diff --git a/reference/shaders-msl-no-opt/comp/glsl.std450.comp b/reference/shaders-msl-no-opt/comp/glsl.std450.comp index 7ac4269d..83ac061a 100644 --- a/reference/shaders-msl-no-opt/comp/glsl.std450.comp +++ b/reference/shaders-msl-no-opt/comp/glsl.std450.comp @@ -24,6 +24,8 @@ struct ResType int _m1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + // Implementation of the GLSL radians() function template inline T radians(T d) diff --git a/reference/shaders-msl-no-opt/comp/loop.comp b/reference/shaders-msl-no-opt/comp/loop.comp index 3e6c820a..74c8b4c2 100644 --- a/reference/shaders-msl-no-opt/comp/loop.comp +++ b/reference/shaders-msl-no-opt/comp/loop.comp @@ -14,6 +14,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _177 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl-no-opt/comp/return.comp b/reference/shaders-msl-no-opt/comp/return.comp index 71fcfbe3..5e0bdcc3 100644 --- a/reference/shaders-msl-no-opt/comp/return.comp +++ b/reference/shaders-msl-no-opt/comp/return.comp @@ -8,6 +8,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _27 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x2-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-2x2-scalar.comp index 05cae5ab..bc392f3d 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x2-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x2-scalar.comp @@ -17,6 +17,8 @@ struct SSBORow float2x2 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x2 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x2-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-2x2-std140.comp index d74b5d55..e378317e 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x2-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x2-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float2x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x2 loaded = float2x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy); diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x2-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-2x2-std430.comp index 05cae5ab..bc392f3d 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x2-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x2-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float2x2 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x2 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x3-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-2x3-scalar.comp index 03a06732..e49ca8d9 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x3-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x3-scalar.comp @@ -19,6 +19,8 @@ struct SSBORow float3x2 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x3 loaded = float2x3(float3(v_29.col_major0[0]), float3(v_29.col_major0[1])); diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x3-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-2x3-std140.comp index 9ba60b47..3ef891d8 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x3-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x3-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float3x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x3 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x3-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-2x3-std430.comp index ff696ea9..6b6e1ea7 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x3-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x3-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float3x2 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x3 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x4-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-2x4-scalar.comp index 014c9139..e4725e0f 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x4-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x4-scalar.comp @@ -17,6 +17,8 @@ struct SSBORow float4x2 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x4-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-2x4-std140.comp index 50bd393b..05e5492c 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x4-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x4-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float4x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-2x4-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-2x4-std430.comp index 014c9139..e4725e0f 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-2x4-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-2x4-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float4x2 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float2x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x2-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-3x2-scalar.comp index 255e8f4f..62107187 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x2-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x2-scalar.comp @@ -19,6 +19,8 @@ struct SSBORow packed_rm_float3x2 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x2 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x2-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-3x2-std140.comp index 879d470b..96a78148 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x2-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x2-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float2x3 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x2 loaded = float3x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy, v_29.col_major0[2].xy); diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x2-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-3x2-std430.comp index 5a7b7705..8a4d00bf 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x2-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x2-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float2x3 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x2 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x3-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-3x3-scalar.comp index 2af05656..1ada47bd 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x3-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x3-scalar.comp @@ -20,6 +20,8 @@ struct SSBORow packed_rm_float3x3 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x3 loaded = float3x3(float3(v_29.col_major0[0]), float3(v_29.col_major0[1]), float3(v_29.col_major0[2])); diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x3-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-3x3-std140.comp index 223ea4e0..44f50f78 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x3-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x3-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float3x3 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x3 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x3-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-3x3-std430.comp index 223ea4e0..44f50f78 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x3-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x3-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float3x3 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x3 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x4-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-3x4-scalar.comp index d74d3dad..155c485b 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x4-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x4-scalar.comp @@ -19,6 +19,8 @@ struct SSBORow packed_rm_float3x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x4-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-3x4-std140.comp index 19881edc..a0784889 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x4-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x4-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float4x3 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-3x4-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-3x4-std430.comp index 19881edc..a0784889 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-3x4-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-3x4-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float4x3 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float3x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x2-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-4x2-scalar.comp index 627d2cef..94befba2 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x2-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x2-scalar.comp @@ -17,6 +17,8 @@ struct SSBORow float2x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x2 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x2-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-4x2-std140.comp index 62df118d..f09d1181 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x2-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x2-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float2x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x2 loaded = float4x2(v_29.col_major0[0].xy, v_29.col_major0[1].xy, v_29.col_major0[2].xy, v_29.col_major0[3].xy); diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x2-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-4x2-std430.comp index 627d2cef..94befba2 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x2-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x2-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float2x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x2 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x3-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-4x3-scalar.comp index fd5e28dd..4edf68ed 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x3-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x3-scalar.comp @@ -19,6 +19,8 @@ struct SSBORow float3x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x3 loaded = float4x3(float3(v_29.col_major0[0]), float3(v_29.col_major0[1]), float3(v_29.col_major0[2]), float3(v_29.col_major0[3])); diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x3-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-4x3-std140.comp index 8baba0e9..e19d043f 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x3-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x3-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float3x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x3 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x3-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-4x3-std430.comp index 8baba0e9..e19d043f 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x3-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x3-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float3x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x3 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x4-scalar.comp b/reference/shaders-msl-no-opt/packing/matrix-4x4-scalar.comp index 6cc48993..ff72744a 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x4-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x4-scalar.comp @@ -17,6 +17,8 @@ struct SSBORow float4x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x4-std140.comp b/reference/shaders-msl-no-opt/packing/matrix-4x4-std140.comp index 6cc48993..ff72744a 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x4-std140.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x4-std140.comp @@ -17,6 +17,8 @@ struct SSBORow float4x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-4x4-std430.comp b/reference/shaders-msl-no-opt/packing/matrix-4x4-std430.comp index 6cc48993..ff72744a 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-4x4-std430.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-4x4-std430.comp @@ -17,6 +17,8 @@ struct SSBORow float4x4 row_major1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void load_store_to_variable_col_major(device SSBOCol& v_29) { float4x4 loaded = v_29.col_major0; diff --git a/reference/shaders-msl-no-opt/packing/matrix-multiply-row-major.comp b/reference/shaders-msl-no-opt/packing/matrix-multiply-row-major.comp index 8c3d5aca..2384e364 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-multiply-row-major.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-multiply-row-major.comp @@ -11,6 +11,8 @@ struct SSBO float3 v1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _11 [[buffer(0)]]) { _11.v0 = _11.v1 * (_11.m1 * _11.m0); diff --git a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major-2.comp b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major-2.comp index 1571ebaa..3fb36e05 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major-2.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major-2.comp @@ -13,6 +13,8 @@ struct SSBO packed_float3 v1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _11 [[buffer(0)]]) { _11.v0 = (float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2])) * float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2]))) * float3(_11.v1); diff --git a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major.comp b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major.comp index 991a76bf..40f00886 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-col-major.comp @@ -11,6 +11,8 @@ struct SSBO float2 v1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _11 [[buffer(0)]]) { _11.v0 = (float2x2(_11.m0[0].xy, _11.m0[1].xy) * float2x2(_11.m1[0].xy, _11.m1[1].xy)) * _11.v1; diff --git a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major-2.comp b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major-2.comp index ab728231..7130c9a8 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major-2.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major-2.comp @@ -13,6 +13,8 @@ struct SSBO packed_float3 v1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _11 [[buffer(0)]]) { _11.v0 = float3(_11.v1) * (float3x3(float3(_11.m1[0]), float3(_11.m1[1]), float3(_11.m1[2])) * float3x3(float3(_11.m0[0]), float3(_11.m0[1]), float3(_11.m0[2]))); diff --git a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major.comp b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major.comp index d683558b..f061dd66 100644 --- a/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major.comp +++ b/reference/shaders-msl-no-opt/packing/matrix-multiply-unpacked-row-major.comp @@ -11,6 +11,8 @@ struct SSBO float2 v1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _11 [[buffer(0)]]) { _11.v0 = _11.v1 * (float2x2(_11.m1[0].xy, _11.m1[1].xy) * float2x2(_11.m0[0].xy, _11.m0[1].xy)); diff --git a/reference/shaders-msl-no-opt/packing/member-padding.comp b/reference/shaders-msl-no-opt/packing/member-padding.comp index a0939047..4f653ecd 100644 --- a/reference/shaders-msl-no-opt/packing/member-padding.comp +++ b/reference/shaders-msl-no-opt/packing/member-padding.comp @@ -11,6 +11,8 @@ struct SSBO float b; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _9 [[buffer(0)]]) { _9.a = 10.0; diff --git a/reference/shaders-msl-no-opt/packing/std140-array-of-vectors.comp b/reference/shaders-msl-no-opt/packing/std140-array-of-vectors.comp index 77da52fe..9b8f76f7 100644 --- a/reference/shaders-msl-no-opt/packing/std140-array-of-vectors.comp +++ b/reference/shaders-msl-no-opt/packing/std140-array-of-vectors.comp @@ -16,6 +16,8 @@ struct SSBO float4 v_unsized[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _29 [[buffer(0)]]) { float loaded1 = _29.v1[1].x; diff --git a/reference/shaders-msl-no-opt/packing/struct-alignment.comp b/reference/shaders-msl-no-opt/packing/struct-alignment.comp index 8baa8146..baa4ee66 100644 --- a/reference/shaders-msl-no-opt/packing/struct-alignment.comp +++ b/reference/shaders-msl-no-opt/packing/struct-alignment.comp @@ -17,6 +17,8 @@ struct SSBO Foo foo; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _12 [[buffer(0)]]) { _12.a.x = 10.0; diff --git a/reference/shaders-msl-no-opt/packing/struct-packing-array-of-scalar.comp b/reference/shaders-msl-no-opt/packing/struct-packing-array-of-scalar.comp index 7715948f..587ee4ad 100644 --- a/reference/shaders-msl-no-opt/packing/struct-packing-array-of-scalar.comp +++ b/reference/shaders-msl-no-opt/packing/struct-packing-array-of-scalar.comp @@ -13,6 +13,8 @@ struct SSBOScalar Foo v[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]]) { buffer_scalar.v[1].a[1u] = 1.0; diff --git a/reference/shaders-msl-no-opt/packing/struct-packing-recursive.comp b/reference/shaders-msl-no-opt/packing/struct-packing-recursive.comp index a7832a5d..e0652b93 100644 --- a/reference/shaders-msl-no-opt/packing/struct-packing-recursive.comp +++ b/reference/shaders-msl-no-opt/packing/struct-packing-recursive.comp @@ -24,6 +24,8 @@ struct SSBOScalar Baz baz; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]]) { buffer_scalar.baz.a.a.a[3u] = 10.0; diff --git a/reference/shaders-msl-no-opt/packing/struct-packing.comp b/reference/shaders-msl-no-opt/packing/struct-packing.comp index 115e1241..a86809fe 100644 --- a/reference/shaders-msl-no-opt/packing/struct-packing.comp +++ b/reference/shaders-msl-no-opt/packing/struct-packing.comp @@ -19,6 +19,8 @@ struct SSBOScalar Bar bar; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBOScalar& buffer_scalar [[buffer(0)]]) { buffer_scalar.foo.a[0u] = 10.0; diff --git a/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp b/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp index 5dd6113b..3c93b5d5 100644 --- a/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp +++ b/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp @@ -42,6 +42,8 @@ struct SSBO E f[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _32 [[buffer(0)]]) { _32.f[0].a = float4(2.0); diff --git a/reference/shaders-msl-no-opt/packing/struct-size-padding.comp b/reference/shaders-msl-no-opt/packing/struct-size-padding.comp index 521ac018..340ab036 100644 --- a/reference/shaders-msl-no-opt/packing/struct-size-padding.comp +++ b/reference/shaders-msl-no-opt/packing/struct-size-padding.comp @@ -42,6 +42,8 @@ struct SSBO E f[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _26 [[buffer(0)]]) { _26.f[0].a = float4(2.0); diff --git a/reference/shaders-msl/comp/access-private-workgroup-in-function.comp b/reference/shaders-msl/comp/access-private-workgroup-in-function.comp index d21c1957..7449b298 100644 --- a/reference/shaders-msl/comp/access-private-workgroup-in-function.comp +++ b/reference/shaders-msl/comp/access-private-workgroup-in-function.comp @@ -5,6 +5,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline void set_f(thread int& f) { f = 40; diff --git a/reference/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp b/reference/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp index f7757cd1..18cfd68c 100644 --- a/reference/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp +++ b/reference/shaders-msl/comp/argument-buffers-discrete.msl2.argument.discrete.comp @@ -23,6 +23,8 @@ struct SSBO2 float4 v; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + struct spvDescriptorSetBuffer0 { const device SSBO0* ssbo0 [[id(0)]]; diff --git a/reference/shaders-msl/comp/array-length.comp b/reference/shaders-msl/comp/array-length.comp index 11f0b47c..fa3aca1a 100644 --- a/reference/shaders-msl/comp/array-length.comp +++ b/reference/shaders-msl/comp/array-length.comp @@ -16,6 +16,8 @@ struct SSBO1 float bz[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline uint get_size(device SSBO& v_14, constant uint& v_14BufferSize, device SSBO1* (&ssbos)[2], constant uint* ssbosBufferSize) { return uint(int((v_14BufferSize - 16) / 16) + int((ssbosBufferSize[1] - 0) / 4)); diff --git a/reference/shaders-msl/comp/array-length.msl2.argument.discrete.comp b/reference/shaders-msl/comp/array-length.msl2.argument.discrete.comp index 6d70cbb5..bbd4cbf3 100644 --- a/reference/shaders-msl/comp/array-length.msl2.argument.discrete.comp +++ b/reference/shaders-msl/comp/array-length.msl2.argument.discrete.comp @@ -27,6 +27,8 @@ struct SSBO3 float bz[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + struct spvDescriptorSetBuffer0 { device SSBO* v_16 [[id(0)]]; diff --git a/reference/shaders-msl/comp/atomic.comp b/reference/shaders-msl/comp/atomic.comp index 04721502..fca72bfc 100644 --- a/reference/shaders-msl/comp/atomic.comp +++ b/reference/shaders-msl/comp/atomic.comp @@ -12,6 +12,8 @@ struct SSBO int i32; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& ssbo [[buffer(0)]]) { threadgroup uint shared_u32; diff --git a/reference/shaders-msl/comp/basic.comp b/reference/shaders-msl/comp/basic.comp index 36b419b7..42518f0d 100644 --- a/reference/shaders-msl/comp/basic.comp +++ b/reference/shaders-msl/comp/basic.comp @@ -21,6 +21,8 @@ struct SSBO3 uint counter; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp b/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp index 084518a5..87b0b442 100644 --- a/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp +++ b/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp @@ -21,9 +21,11 @@ struct SSBO3 uint counter; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(constant uint3& spvDispatchBase [[buffer(29)]], const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) { - gl_GlobalInvocationID += spvDispatchBase * uint3(1, 1, 1); + gl_GlobalInvocationID += spvDispatchBase * gl_WorkGroupSize; gl_WorkGroupID += spvDispatchBase; uint ident = gl_GlobalInvocationID.x; uint workgroup = gl_WorkGroupID.x; diff --git a/reference/shaders-msl/comp/basic.dynamic-buffer.msl2.comp b/reference/shaders-msl/comp/basic.dynamic-buffer.msl2.invalid.comp similarity index 100% rename from reference/shaders-msl/comp/basic.dynamic-buffer.msl2.comp rename to reference/shaders-msl/comp/basic.dynamic-buffer.msl2.invalid.comp diff --git a/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp b/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp index 170e4920..d231a029 100644 --- a/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp +++ b/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp @@ -13,6 +13,8 @@ struct SSBO1 int4 outputs[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO0& _25 [[buffer(0)]], device SSBO1& _39 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp b/reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp index 59eb961f..8db1c2f8 100644 --- a/reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp +++ b/reference/shaders-msl/comp/bitcast-16bit-2.invalid.comp @@ -18,6 +18,8 @@ struct UBO half4 const0; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO1& _21 [[buffer(0)]], device SSBO0& _29 [[buffer(1)]], constant UBO& _40 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/coherent-block.comp b/reference/shaders-msl/comp/coherent-block.comp index 580b9e3e..58bbacb7 100644 --- a/reference/shaders-msl/comp/coherent-block.comp +++ b/reference/shaders-msl/comp/coherent-block.comp @@ -8,6 +8,8 @@ struct SSBO float4 value; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(volatile device SSBO& _10 [[buffer(0)]]) { _10.value = float4(20.0); diff --git a/reference/shaders-msl/comp/coherent-image.comp b/reference/shaders-msl/comp/coherent-image.comp index c6af46b5..50904844 100644 --- a/reference/shaders-msl/comp/coherent-image.comp +++ b/reference/shaders-msl/comp/coherent-image.comp @@ -8,6 +8,8 @@ struct SSBO int4 value; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(volatile device SSBO& _10 [[buffer(0)]], texture2d uImage [[texture(0)]]) { _10.value = uImage.read(uint2(int2(10))); diff --git a/reference/shaders-msl/comp/composite-construct.comp b/reference/shaders-msl/comp/composite-construct.comp index 87fcfcd8..8d6bf982 100644 --- a/reference/shaders-msl/comp/composite-construct.comp +++ b/reference/shaders-msl/comp/composite-construct.comp @@ -21,6 +21,8 @@ struct Composite float4 b; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + constant float4 _43[2] = { float4(20.0), float4(40.0) }; template diff --git a/reference/shaders-msl/comp/copy-array-of-arrays.comp b/reference/shaders-msl/comp/copy-array-of-arrays.comp index e3b3660a..5f8b0330 100644 --- a/reference/shaders-msl/comp/copy-array-of-arrays.comp +++ b/reference/shaders-msl/comp/copy-array-of-arrays.comp @@ -12,6 +12,8 @@ struct BUF float c; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + 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 } }; diff --git a/reference/shaders-msl/comp/defer-parens.comp b/reference/shaders-msl/comp/defer-parens.comp index 76dce777..9a567fa6 100644 --- a/reference/shaders-msl/comp/defer-parens.comp +++ b/reference/shaders-msl/comp/defer-parens.comp @@ -9,6 +9,8 @@ struct SSBO int index; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _13 [[buffer(0)]]) { float4 d = _13.data; diff --git a/reference/shaders-msl/comp/dowhile.comp b/reference/shaders-msl/comp/dowhile.comp index 3482fb35..2b4de9eb 100644 --- a/reference/shaders-msl/comp/dowhile.comp +++ b/reference/shaders-msl/comp/dowhile.comp @@ -14,6 +14,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/image-cube-array-load-store.comp b/reference/shaders-msl/comp/image-cube-array-load-store.comp index ef67a326..c0b83c46 100644 --- a/reference/shaders-msl/comp/image-cube-array-load-store.comp +++ b/reference/shaders-msl/comp/image-cube-array-load-store.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(texturecube_array uImageIn [[texture(0)]], texturecube_array uImageOut [[texture(1)]]) { int3 coord = int3(9, 7, 11); diff --git a/reference/shaders-msl/comp/image.comp b/reference/shaders-msl/comp/image.comp index f3bc1455..e7c9c763 100644 --- a/reference/shaders-msl/comp/image.comp +++ b/reference/shaders-msl/comp/image.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(texture2d uImageIn [[texture(0)]], texture2d uImageOut [[texture(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { float4 v = uImageIn.read(uint2((int2(gl_GlobalInvocationID.xy) + int2(uImageIn.get_width(), uImageIn.get_height())))); diff --git a/reference/shaders-msl/comp/insert.comp b/reference/shaders-msl/comp/insert.comp index 0f56a651..e3e858a4 100644 --- a/reference/shaders-msl/comp/insert.comp +++ b/reference/shaders-msl/comp/insert.comp @@ -8,6 +8,8 @@ struct SSBO float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { float4 v; diff --git a/reference/shaders-msl/comp/int64.invalid.msl22.comp b/reference/shaders-msl/comp/int64.invalid.msl22.comp index 6eb4a8a8..dd014343 100644 --- a/reference/shaders-msl/comp/int64.invalid.msl22.comp +++ b/reference/shaders-msl/comp/int64.invalid.msl22.comp @@ -41,6 +41,8 @@ struct SSBO uint u32; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _96 [[buffer(0)]]) { SSBO0_Type ssbo_0; diff --git a/reference/shaders-msl/comp/inverse.comp b/reference/shaders-msl/comp/inverse.comp index 3ea692fa..a4ef60c4 100644 --- a/reference/shaders-msl/comp/inverse.comp +++ b/reference/shaders-msl/comp/inverse.comp @@ -19,6 +19,8 @@ struct MatrixIn float4x4 m4in; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + // Returns the determinant of a 2x2 matrix. inline float spvDet2x2(float a1, float a2, float b1, float b2) { diff --git a/reference/shaders-msl/comp/mat3.comp b/reference/shaders-msl/comp/mat3.comp index c2d9a7c8..fcb8f7a6 100644 --- a/reference/shaders-msl/comp/mat3.comp +++ b/reference/shaders-msl/comp/mat3.comp @@ -8,6 +8,8 @@ struct SSBO2 float3x3 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _22 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/mod.comp b/reference/shaders-msl/comp/mod.comp index 82d288f8..96264b05 100644 --- a/reference/shaders-msl/comp/mod.comp +++ b/reference/shaders-msl/comp/mod.comp @@ -15,6 +15,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + // Implementation of the GLSL mod() function, which is slightly different than Metal fmod() template inline Tx mod(Tx x, Ty y) diff --git a/reference/shaders-msl/comp/modf.comp b/reference/shaders-msl/comp/modf.comp index ef50a021..5a5ac3db 100644 --- a/reference/shaders-msl/comp/modf.comp +++ b/reference/shaders-msl/comp/modf.comp @@ -13,6 +13,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/outer-product.comp b/reference/shaders-msl/comp/outer-product.comp index 8e32db39..e589642d 100644 --- a/reference/shaders-msl/comp/outer-product.comp +++ b/reference/shaders-msl/comp/outer-product.comp @@ -23,6 +23,8 @@ struct ReadSSBO float4 v4; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _21 [[buffer(0)]], const device ReadSSBO& _26 [[buffer(1)]]) { _21.m22 = float2x2(_26.v2 * _26.v2.x, _26.v2 * _26.v2.y); diff --git a/reference/shaders-msl/comp/read-write-only.comp b/reference/shaders-msl/comp/read-write-only.comp index 7547b417..0cf8d8e3 100644 --- a/reference/shaders-msl/comp/read-write-only.comp +++ b/reference/shaders-msl/comp/read-write-only.comp @@ -21,6 +21,8 @@ struct SSBO1 float4 data3; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _10 [[buffer(0)]], const device SSBO0& _15 [[buffer(1)]], device SSBO1& _21 [[buffer(2)]]) { _10.data4 = _15.data0 + _21.data2; diff --git a/reference/shaders-msl/comp/rmw-matrix.comp b/reference/shaders-msl/comp/rmw-matrix.comp index 150db7ed..b53a3a75 100644 --- a/reference/shaders-msl/comp/rmw-matrix.comp +++ b/reference/shaders-msl/comp/rmw-matrix.comp @@ -13,6 +13,8 @@ struct SSBO float4x4 c1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _11 [[buffer(0)]]) { _11.a *= _11.a1; diff --git a/reference/shaders-msl/comp/rmw-opt.comp b/reference/shaders-msl/comp/rmw-opt.comp index 060f9f9c..229154fc 100644 --- a/reference/shaders-msl/comp/rmw-opt.comp +++ b/reference/shaders-msl/comp/rmw-opt.comp @@ -8,6 +8,8 @@ struct SSBO int a; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _9 [[buffer(0)]]) { _9.a += 10; diff --git a/reference/shaders-msl/comp/scalar-std450-distance-length-normalize.comp b/reference/shaders-msl/comp/scalar-std450-distance-length-normalize.comp index 312a6f94..0ae6e556 100644 --- a/reference/shaders-msl/comp/scalar-std450-distance-length-normalize.comp +++ b/reference/shaders-msl/comp/scalar-std450-distance-length-normalize.comp @@ -12,6 +12,8 @@ struct SSBO float e; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _9 [[buffer(0)]]) { _9.c = abs(_9.a - _9.b); diff --git a/reference/shaders-msl/comp/spec-constant-op-member-array.comp b/reference/shaders-msl/comp/spec-constant-op-member-array.comp index d3c8b7dc..8f54f052 100644 --- a/reference/shaders-msl/comp/spec-constant-op-member-array.comp +++ b/reference/shaders-msl/comp/spec-constant-op-member-array.comp @@ -40,6 +40,7 @@ struct SSBO constant int e_tmp [[function_constant(3)]]; constant int e = is_function_constant_defined(e_tmp) ? e_tmp : 400; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); kernel void main0(device SSBO& _22 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { diff --git a/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp b/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp index 6988febf..0cafa5f3 100644 --- a/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp +++ b/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp @@ -24,6 +24,8 @@ struct SSBO Sub_1 sub[2]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { Sub foo; diff --git a/reference/shaders-msl/comp/struct-layout.comp b/reference/shaders-msl/comp/struct-layout.comp index b6ee59f1..8f2ab2d6 100644 --- a/reference/shaders-msl/comp/struct-layout.comp +++ b/reference/shaders-msl/comp/struct-layout.comp @@ -18,6 +18,8 @@ struct SSBO Foo in_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _23 [[buffer(0)]], const device SSBO& _30 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/struct-nested.comp b/reference/shaders-msl/comp/struct-nested.comp index 86229e34..e3d04bed 100644 --- a/reference/shaders-msl/comp/struct-nested.comp +++ b/reference/shaders-msl/comp/struct-nested.comp @@ -28,6 +28,8 @@ struct dstbuffer s2_1 test[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device dstbuffer& _19 [[buffer(0)]]) { s2 testVal; diff --git a/reference/shaders-msl/comp/struct-packing.comp b/reference/shaders-msl/comp/struct-packing.comp index ad22c257..775bb348 100644 --- a/reference/shaders-msl/comp/struct-packing.comp +++ b/reference/shaders-msl/comp/struct-packing.comp @@ -118,6 +118,8 @@ struct SSBO0 float4 array[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]]) { Content_1 _60 = ssbo_140.content; diff --git a/reference/shaders-msl/comp/torture-loop.comp b/reference/shaders-msl/comp/torture-loop.comp index 1b65a3af..e92e71d6 100644 --- a/reference/shaders-msl/comp/torture-loop.comp +++ b/reference/shaders-msl/comp/torture-loop.comp @@ -14,6 +14,8 @@ struct SSBO2 float4 out_data[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint ident = gl_GlobalInvocationID.x; diff --git a/reference/shaders-msl/comp/type-alias.comp b/reference/shaders-msl/comp/type-alias.comp index a1bd9336..e419efb2 100644 --- a/reference/shaders-msl/comp/type-alias.comp +++ b/reference/shaders-msl/comp/type-alias.comp @@ -40,6 +40,8 @@ struct SSBO2 float4 outputs[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline float4 overload(thread const S0& s0) { return s0.a; diff --git a/reference/shaders-msl/comp/udiv.comp b/reference/shaders-msl/comp/udiv.comp index 32874ad7..7f7315b8 100644 --- a/reference/shaders-msl/comp/udiv.comp +++ b/reference/shaders-msl/comp/udiv.comp @@ -13,6 +13,8 @@ struct SSBO uint inputs[1]; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO2& _10 [[buffer(0)]], device SSBO& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _10.outputs[gl_GlobalInvocationID.x] = _23.inputs[gl_GlobalInvocationID.x] / 29u; diff --git a/reference/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp b/reference/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp index a37fe519..cea12980 100644 --- a/reference/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp +++ b/reference/shaders-msl/desktop-only/comp/extended-arithmetic.desktop.comp @@ -91,6 +91,8 @@ struct ResType_7 int4 _m1; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBOUint& u [[buffer(0)]], device SSBOInt& i [[buffer(1)]]) { ResType _25; diff --git a/reference/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp b/reference/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp index 2f731356..4f341227 100644 --- a/reference/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp +++ b/reference/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp @@ -121,6 +121,8 @@ struct SSBO2 packed_rm_float3x2 m2; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]]) { ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0].xy; diff --git a/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp b/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp index 4ebab8c7..2fd77131 100644 --- a/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp +++ b/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.comp @@ -10,6 +10,8 @@ struct SSBO float FragColor; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + inline uint4 spvSubgroupBallot(bool value) { simd_vote vote = simd_ballot(value); diff --git a/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp b/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp index 84fcb9c3..5c67f6e4 100644 --- a/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp +++ b/reference/shaders-msl/vulkan/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp @@ -8,6 +8,8 @@ struct SSBO float FragColor; }; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[quadgroups_per_threadgroup]], uint gl_SubgroupID [[quadgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_quadgroup]]) { _9.FragColor = float(gl_NumSubgroups); diff --git a/reference/shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.frag b/reference/shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.invalid.frag similarity index 100% rename from reference/shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.frag rename to reference/shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.invalid.frag diff --git a/reference/shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag b/reference/shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag similarity index 100% rename from reference/shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag rename to reference/shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag diff --git a/shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag b/shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag similarity index 100% rename from shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag rename to shaders-hlsl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag diff --git a/shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag b/shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag similarity index 100% rename from shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag rename to shaders-msl-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag diff --git a/shaders-msl/comp/basic.dynamic-buffer.msl2.comp b/shaders-msl/comp/basic.dynamic-buffer.msl2.invalid.comp similarity index 100% rename from shaders-msl/comp/basic.dynamic-buffer.msl2.comp rename to shaders-msl/comp/basic.dynamic-buffer.msl2.invalid.comp diff --git a/shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.frag b/shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.invalid.frag similarity index 100% rename from shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.frag rename to shaders-no-opt/asm/frag/combined-image-sampler-dxc-min16float.asm.invalid.frag diff --git a/shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag b/shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag similarity index 100% rename from shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.frag rename to shaders-no-opt/asm/frag/switch-block-case-fallthrough.asm.invalid.frag diff --git a/test_shaders.py b/test_shaders.py index 9709e2cd..a47f41b3 100755 --- a/test_shaders.py +++ b/test_shaders.py @@ -171,7 +171,7 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths): if spirv: subprocess.check_call(spirv_cmd) else: - subprocess.check_call([paths.glslang, '--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) + subprocess.check_call([paths.glslang, '--amb' ,'--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) if opt: subprocess.check_call([paths.spirv_opt, '--skip-validation', '-O', '-o', spirv_path, spirv_path]) @@ -264,7 +264,7 @@ def validate_shader_hlsl(shader, force_no_external_validation, paths): test_glslang = False if test_glslang: - subprocess.check_call([paths.glslang, '-e', 'main', '-D', '--target-env', 'vulkan1.1', '-V', shader]) + subprocess.check_call([paths.glslang, '--amb', '-e', 'main', '-D', '--target-env', 'vulkan1.1', '-V', shader]) is_no_fxc = '.nofxc.' in shader global ignore_fxc if (not ignore_fxc) and (not force_no_external_validation) and (not is_no_fxc): @@ -307,7 +307,7 @@ def cross_compile_hlsl(shader, spirv, opt, force_no_external_validation, iterati if spirv: subprocess.check_call(spirv_cmd) else: - subprocess.check_call([paths.glslang, '--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) + subprocess.check_call([paths.glslang, '--amb', '--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) if opt: subprocess.check_call([paths.spirv_opt, '--skip-validation', '-O', '-o', spirv_path, spirv_path]) @@ -339,7 +339,7 @@ def cross_compile_reflect(shader, spirv, opt, iterations, paths): if spirv: subprocess.check_call(spirv_cmd) else: - subprocess.check_call([paths.glslang, '--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) + subprocess.check_call([paths.glslang, '--amb', '--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) if opt: subprocess.check_call([paths.spirv_opt, '--skip-validation', '-O', '-o', spirv_path, spirv_path]) @@ -352,7 +352,7 @@ def cross_compile_reflect(shader, spirv, opt, iterations, paths): def validate_shader(shader, vulkan, paths): if vulkan: - subprocess.check_call([paths.glslang, '--target-env', 'vulkan1.1', '-V', shader]) + subprocess.check_call([paths.glslang, '--amb', '--target-env', 'vulkan1.1', '-V', shader]) else: subprocess.check_call([paths.glslang, shader]) @@ -370,7 +370,7 @@ def cross_compile(shader, vulkan, spirv, invalid_spirv, eliminate, is_legacy, fl if spirv: subprocess.check_call(spirv_cmd) else: - subprocess.check_call([paths.glslang, '--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) + subprocess.check_call([paths.glslang, '--amb', '--target-env', 'vulkan1.1', '-V', '-o', spirv_path, shader]) if opt and (not invalid_spirv): subprocess.check_call([paths.spirv_opt, '--skip-validation', '-O', '-o', spirv_path, spirv_path])