Merge branch 'rollDEPS' of git://github.com/zoddicus/SPIRV-Cross

checkout_glslang_spirv_tools.sh

@@ -1,8 +1,8 @@
 #!/bin/bash
 
-GLSLANG_REV=25a508cc735109cc4e382c3a1cc293a9452a41f3
+GLSLANG_REV=4b5159ea8170fa34e29f13448fddebf88e0a722a
-SPIRV_TOOLS_REV=55adf4cf707bb12c29fc12f784ebeaa29a819e9b
+SPIRV_TOOLS_REV=bbb29870b510f83f99994358179c9ea6838c3100
-SPIRV_HEADERS_REV=29c11140baaf9f7fdaa39a583672c556bf1795a1
+SPIRV_HEADERS_REV=601d738723ac381741311c6c98c36d6170be14a2
 
 if [ -z $PROTOCOL ]; then
 	PROTOCOL=git

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);
 

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);

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<uint> uImage : register(u0);
 RWTexture2D<int> iImage : register(u1);

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<float> uImageIn : register(u0);

reference/opt/shaders-hlsl/comp/image.comp

@@ -1,3 +1,5 @@
+static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u);
+
 RWTexture2D<float> uImageInF : register(u0);
 RWTexture2D<float> uImageOutF : register(u1);
 RWTexture2D<int> uImageInI : register(u2);

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);
 

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)
 {

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)
 {

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);
 

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()

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)));
+    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(_253[0]));
+    _28.Store4(0, asuint(_257[0]));
-    _28.Store4(16, asuint(_253[1]));
+    _28.Store4(16, asuint(_257[1]));
-    _28.Store4(32, asuint(_253[2]));
+    _28.Store4(32, asuint(_257[2]));
-    _28.Store4(48, asuint(_253[3]));
+    _28.Store4(48, asuint(_257[3]));
-    float2x2 _256 = asfloat(uint2x2(_28.Load(144), _28.Load(152), _28.Load(148), _28.Load(156)));
+    float2x2 _260 = asfloat(uint2x2(_28.Load(144), _28.Load(152), _28.Load(148), _28.Load(156)));
-    _28.Store2(128, asuint(_256[0]));
+    _28.Store2(128, asuint(_260[0]));
-    _28.Store2(136, asuint(_256[1]));
+    _28.Store2(136, asuint(_260[1]));
-    float2x3 _259 = asfloat(uint2x3(_28.Load(192), _28.Load(200), _28.Load(208), _28.Load(196), _28.Load(204), _28.Load(212)));
+    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(_259[0]));
+    _28.Store3(160, asuint(_263[0]));
-    _28.Store3(176, asuint(_259[1]));
+    _28.Store3(176, asuint(_263[1]));
-    float3x2 _262 = asfloat(uint3x2(_28.Load(240), _28.Load(256), _28.Load(244), _28.Load(260), _28.Load(248), _28.Load(264)));
+    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(_262[0]));
+    _28.Store2(216, asuint(_266[0]));
-    _28.Store2(224, asuint(_262[1]));
+    _28.Store2(224, asuint(_266[1]));
-    _28.Store2(232, asuint(_262[2]));
+    _28.Store2(232, asuint(_266[2]));
-    float4x4 _265 = asfloat(uint4x4(_28.Load4(0), _28.Load4(16), _28.Load4(32), _28.Load4(48)));
+    float4x4 _269 = asfloat(uint4x4(_28.Load4(0), _28.Load4(16), _28.Load4(32), _28.Load4(48)));
-    _28.Store(64, asuint(_265[0].x));
+    _28.Store(64, asuint(_269[0].x));
-    _28.Store(68, asuint(_265[1].x));
+    _28.Store(68, asuint(_269[1].x));
-    _28.Store(72, asuint(_265[2].x));
+    _28.Store(72, asuint(_269[2].x));
-    _28.Store(76, asuint(_265[3].x));
+    _28.Store(76, asuint(_269[3].x));
-    _28.Store(80, asuint(_265[0].y));
+    _28.Store(80, asuint(_269[0].y));
-    _28.Store(84, asuint(_265[1].y));
+    _28.Store(84, asuint(_269[1].y));
-    _28.Store(88, asuint(_265[2].y));
+    _28.Store(88, asuint(_269[2].y));
-    _28.Store(92, asuint(_265[3].y));
+    _28.Store(92, asuint(_269[3].y));
-    _28.Store(96, asuint(_265[0].z));
+    _28.Store(96, asuint(_269[0].z));
-    _28.Store(100, asuint(_265[1].z));
+    _28.Store(100, asuint(_269[1].z));
-    _28.Store(104, asuint(_265[2].z));
+    _28.Store(104, asuint(_269[2].z));
-    _28.Store(108, asuint(_265[3].z));
+    _28.Store(108, asuint(_269[3].z));
-    _28.Store(112, asuint(_265[0].w));
+    _28.Store(112, asuint(_269[0].w));
-    _28.Store(116, asuint(_265[1].w));
+    _28.Store(116, asuint(_269[1].w));
-    _28.Store(120, asuint(_265[2].w));
+    _28.Store(120, asuint(_269[2].w));
-    _28.Store(124, asuint(_265[3].w));
+    _28.Store(124, asuint(_269[3].w));
-    float2x2 _268 = asfloat(uint2x2(_28.Load2(128), _28.Load2(136)));
+    float2x2 _272 = asfloat(uint2x2(_28.Load2(128), _28.Load2(136)));
-    _28.Store(144, asuint(_268[0].x));
+    _28.Store(144, asuint(_272[0].x));
-    _28.Store(148, asuint(_268[1].x));
+    _28.Store(148, asuint(_272[1].x));
-    _28.Store(152, asuint(_268[0].y));
+    _28.Store(152, asuint(_272[0].y));
-    _28.Store(156, asuint(_268[1].y));
+    _28.Store(156, asuint(_272[1].y));
-    float2x3 _271 = asfloat(uint2x3(_28.Load3(160), _28.Load3(176)));
+    float2x3 _275 = asfloat(uint2x3(_28.Load3(160), _28.Load3(176)));
-    _28.Store(192, asuint(_271[0].x));
+    _28.Store(192, asuint(_275[0].x));
-    _28.Store(196, asuint(_271[1].x));
+    _28.Store(196, asuint(_275[1].x));
-    _28.Store(200, asuint(_271[0].y));
+    _28.Store(200, asuint(_275[0].y));
-    _28.Store(204, asuint(_271[1].y));
+    _28.Store(204, asuint(_275[1].y));
-    _28.Store(208, asuint(_271[0].z));
+    _28.Store(208, asuint(_275[0].z));
-    _28.Store(212, asuint(_271[1].z));
+    _28.Store(212, asuint(_275[1].z));
-    float3x2 _274 = asfloat(uint3x2(_28.Load2(216), _28.Load2(224), _28.Load2(232)));
+    float3x2 _278 = asfloat(uint3x2(_28.Load2(216), _28.Load2(224), _28.Load2(232)));
-    _28.Store(240, asuint(_274[0].x));
+    _28.Store(240, asuint(_278[0].x));
-    _28.Store(244, asuint(_274[1].x));
+    _28.Store(244, asuint(_278[1].x));
-    _28.Store(248, asuint(_274[2].x));
+    _28.Store(248, asuint(_278[2].x));
-    _28.Store(256, asuint(_274[0].y));
+    _28.Store(256, asuint(_278[0].y));
-    _28.Store(260, asuint(_274[1].y));
+    _28.Store(260, asuint(_278[1].y));
-    _28.Store(264, asuint(_274[2].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));

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()

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);
 

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()

reference/opt/shaders-hlsl/comp/ssbo-array.comp

@@ -1,3 +1,5 @@
+static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u);
+
 void comp_main()
 {
 }

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;

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()
 {
 }

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)]];

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[] =

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)]];

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;

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];

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)
     {

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<uint>(as_type<half2>(_25.inputs[gl_GlobalInvocationID.x].xy) + half2(half(1.0))));

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<short2>(_29.inputs[gl_GlobalInvocationID.x].x) + as_type<short2>(_40.const0.xy);

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);

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<int> uImage [[texture(0)]])
 {
     _10.value = uImage.read(uint2(int2(10)));

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<typename T, uint A>
 inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
 {

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;

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;

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;
+    float4 _59;
-    int _58;
+    int _60;
-    _58 = 0;
+    _60 = 0;
-    _57 = _28.in_data[gl_GlobalInvocationID.x];
+    _59 = _28.in_data[gl_GlobalInvocationID.x];
     float4 _42;
     for (;;)
     {
-        _42 = _28.mvp * _57;
+        _42 = _28.mvp * _59;
-        int _44 = _58 + 1;
+        int _44 = _60 + 1;
         if (_44 < 16)
         {
-            _58 = _44;
+            _60 = _44;
-            _57 = _42;
+            _59 = _42;
         }
         else
         {

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<float> uImageIn [[texture(0)]], texturecube_array<float, access::write> 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);

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<float> uImageIn [[texture(0)]], texture2d<float, access::write> uImageOut [[texture(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     int2 _23 = int2(gl_GlobalInvocationID.xy);

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;
+    float4 _46 = _53;
-    _45.x = 10.0;
+    _46.x = 10.0;
-    float4 _47 = _45;
+    float4 _48 = _46;
-    _47.y = 30.0;
+    _48.y = 30.0;
-    float4 _49 = _47;
+    float4 _50 = _48;
-    _49.z = 70.0;
+    _50.z = 70.0;
-    float4 _51 = _49;
+    float4 _52 = _50;
-    _51.w = 90.0;
+    _52.w = 90.0;
-    _27.out_data[gl_GlobalInvocationID.x] = _51;
+    _27.out_data[gl_GlobalInvocationID.x] = _52;
     _27.out_data[gl_GlobalInvocationID.x].y = 20.0;
 }
 

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;
+    long4 _139;
-    ulong4 _141;
+    ulong4 _143;
-    _137 = abs((_137 + long4(30l, 40l, 50l, 60l)) + long4(_141 + ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul)));
+    _139 = abs((_139 + long4(30l, 40l, 50l, 60l)) + long4(_143 + ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul)));
-    _141 += ulong4(long4(999999999999999999l, 8888888888888888l, 77777777777777777l, 6666666666666666l));
+    _143 += ulong4(long4(999999999999999999l, 8888888888888888l, 77777777777777777l, 6666666666666666l));
-    ulong _109 = ulong(_162);
+    ulong _109 = ulong(_164);
-    _96.s32 = int(uint(((ulong(_137.x) + _141.y) + _109) + _109));
+    _96.s32 = int(uint(((ulong(_139.x) + _143.y) + _109) + _109));
-    _96.u32 = uint(((ulong(_137.y) + _141.z) + ulong(_162 + 1l)) + _109);
+    _96.u32 = uint(((ulong(_139.y) + _143.z) + ulong(_164 + 1l)) + _109);
 }
 

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)
 {

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));

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<typename Tx, typename Ty>
 inline Tx mod(Tx x, Ty y)

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;

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);

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;

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;

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;

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);

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]])
 {

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];
+    float _155[2];
-    _153[0] = _27.sub[gl_WorkGroupID.x].f[0].x;
+    _155[0] = _27.sub[gl_WorkGroupID.x].f[0].x;
-    _153[1] = _27.sub[gl_WorkGroupID.x].f[1].x;
+    _155[1] = _27.sub[gl_WorkGroupID.x].f[1].x;
-    float2 _154[2];
+    float2 _156[2];
-    _154[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy;
+    _156[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy;
-    _154[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy;
+    _156[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy;
-    float3 _155[2];
+    float3 _157[2];
-    _155[0] = _27.sub[gl_WorkGroupID.x].f3[0];
+    _157[0] = _27.sub[gl_WorkGroupID.x].f3[0];
-    _155[1] = _27.sub[gl_WorkGroupID.x].f3[1];
+    _157[1] = _27.sub[gl_WorkGroupID.x].f3[1];
-    float4 _156[2];
+    float4 _158[2];
-    _156[0] = _27.sub[gl_WorkGroupID.x].f4[0];
+    _158[0] = _27.sub[gl_WorkGroupID.x].f4[0];
-    _156[1] = _27.sub[gl_WorkGroupID.x].f4[1];
+    _158[1] = _27.sub[gl_WorkGroupID.x].f4[1];
-    _153[gl_GlobalInvocationID.x] += 1.0;
+    _155[gl_GlobalInvocationID.x] += 1.0;
-    _154[gl_GlobalInvocationID.x] += float2(2.0);
+    _156[gl_GlobalInvocationID.x] += float2(2.0);
-    _155[gl_GlobalInvocationID.x] += float3(3.0);
+    _157[gl_GlobalInvocationID.x] += float3(3.0);
-    _156[gl_GlobalInvocationID.x] += float4(4.0);
+    _158[gl_GlobalInvocationID.x] += float4(4.0);
-    _27.sub[gl_WorkGroupID.x].f[0].x = _153[0];
+    _27.sub[gl_WorkGroupID.x].f[0].x = _155[0];
-    _27.sub[gl_WorkGroupID.x].f[1].x = _153[1];
+    _27.sub[gl_WorkGroupID.x].f[1].x = _155[1];
-    _27.sub[gl_WorkGroupID.x].f2[0].xy = _154[0];
+    _27.sub[gl_WorkGroupID.x].f2[0].xy = _156[0];
-    _27.sub[gl_WorkGroupID.x].f2[1].xy = _154[1];
+    _27.sub[gl_WorkGroupID.x].f2[1].xy = _156[1];
-    _27.sub[gl_WorkGroupID.x].f3[0] = _155[0];
+    _27.sub[gl_WorkGroupID.x].f3[0] = _157[0];
-    _27.sub[gl_WorkGroupID.x].f3[1] = _155[1];
+    _27.sub[gl_WorkGroupID.x].f3[1] = _157[1];
-    _27.sub[gl_WorkGroupID.x].f4[0] = _156[0];
+    _27.sub[gl_WorkGroupID.x].f4[0] = _158[0];
-    _27.sub[gl_WorkGroupID.x].f4[1] = _156[1];
+    _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);
 }

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;

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;

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;

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;
+    float4 _101;
-    _99 = _24.in_data[gl_GlobalInvocationID.x];
+    _101 = _24.in_data[gl_GlobalInvocationID.x];
-    for (int _93 = 0; (_93 + 1) < 10; )
+    for (int _95 = 0; (_95 + 1) < 10; )
     {
-        _99 *= 2.0;
+        _101 *= 2.0;
-        _93 += 2;
+        _95 += 2;
         continue;
     }
-    float4 _98;
+    float4 _100;
-    _98 = _99;
+    _100 = _101;
-    float4 _103;
+    float4 _105;
-    for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++)
+    for (uint _96 = 0u; _96 < 16u; _100 = _105, _96++)
     {
-        _103 = _98;
+        _105 = _100;
-        for (uint _100 = 0u; _100 < 30u; )
+        for (uint _102 = 0u; _102 < 30u; )
         {
-            _103 = _24.mvp * _103;
+            _105 = _24.mvp * _105;
-            _100++;
+            _102++;
             continue;
         }
     }
-    _89.out_data[gl_GlobalInvocationID.x] = _98;
+    _89.out_data[gl_GlobalInvocationID.x] = _100;
 }
 

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;

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;

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;

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<char2>(_196);
+    char2 _201 = as_type<char2>(10);
-    char4 _199 = as_type<char4>(_197);
+    char2 _198 = _201;
-    _196 = as_type<short>(_198);
+    char4 _199 = as_type<char4>(20);
+    _196 = as_type<short>(_201);
     _197 = as_type<int>(_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<uchar2>(_220);
+    uchar2 _225 = as_type<uchar2>(10u);
-    uchar4 _223 = as_type<uchar4>(_221);
+    uchar2 _222 = _225;
-    _220 = as_type<ushort>(_222);
+    uchar4 _223 = as_type<uchar4>(20u);
+    _220 = as_type<ushort>(_225);
     _221 = as_type<uint>(_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);
+    char4 _251 = _246 + char4(registers.i8);
-    _244 += char4(-40);
+    _244 = _251;
-    _244 += char4(-50);
+    char4 _254 = _251 + char4(-40);
-    _244 += char4(char(10), char(20), char(30), char(40));
+    _244 = _254;
-    _244 += char4(ssbo.i8[4]);
+    char4 _256 = _254 + char4(-50);
-    _244 += char4(ubo.i8);
+    _244 = _256;
-    out.FragColorInt = int4(_244);
+    char4 _258 = _256 + char4(char(10), char(20), char(30), char(40));
-    uchar4 _271 = uchar4(_246);
+    _244 = _258;
-    _271 += uchar4(registers.u8);
+    char4 _263 = _258 + char4(ssbo.i8[4]);
-    _271 += uchar4(216);
+    _244 = _263;
-    _271 += uchar4(206);
+    char4 _268 = _263 + char4(ubo.i8);
-    _271 += uchar4(uchar(10), uchar(20), uchar(30), uchar(40));
+    _244 = _268;
-    _271 += uchar4(ssbo.u8[4]);
+    out.FragColorInt = int4(_268);
-    _271 += uchar4(ubo.u8);
+    uchar4 _274 = uchar4(_246);
-    out.FragColorUint = uint4(_271);
+    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;
 }
 

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;

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);

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);

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;
 }
 

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;
+    float _184;
-    _180 = _183;
+    _184 = _187;
-    for (int _179 = 0; _179 < 20; )
+    for (int _183 = 0; _183 < 20; )
     {
-        _180 += 10.0;
+        _184 += 10.0;
-        _179++;
+        _183++;
         continue;
     }
-    _11.data = _180;
+    _11.data = _184;
 }
 

reference/opt/shaders/comp/dowhile.comp

@@ -14,19 +14,19 @@ layout(binding = 1, std430) writeonly buffer SSBO2
 
 void main()
 {
-    vec4 _57;
+    vec4 _59;
-    int _58;
+    int _60;
-    _58 = 0;
+    _60 = 0;
-    _57 = _28.in_data[gl_GlobalInvocationID.x];
+    _59 = _28.in_data[gl_GlobalInvocationID.x];
     vec4 _42;
     for (;;)
     {
-        _42 = _28.mvp * _57;
+        _42 = _28.mvp * _59;
-        int _44 = _58 + 1;
+        int _44 = _60 + 1;
         if (_44 < 16)
         {
-            _58 = _44;
+            _60 = _44;
-            _57 = _42;
+            _59 = _42;
         }
         else
         {

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;
+    vec4 _46 = _53;
-    _45.x = 10.0;
+    _46.x = 10.0;
-    vec4 _47 = _45;
+    vec4 _48 = _46;
-    _47.y = 30.0;
+    _48.y = 30.0;
-    vec4 _49 = _47;
+    vec4 _50 = _48;
-    _49.z = 70.0;
+    _50.z = 70.0;
-    vec4 _51 = _49;
+    vec4 _52 = _50;
-    _51.w = 90.0;
+    _52.w = 90.0;
-    _27.out_data[gl_GlobalInvocationID.x] = _51;
+    _27.out_data[gl_GlobalInvocationID.x] = _52;
     _27.out_data[gl_GlobalInvocationID.x].y = 20.0;
 }
 

reference/opt/shaders/comp/torture-loop.comp

@@ -14,27 +14,27 @@ layout(binding = 1, std430) writeonly buffer SSBO2
 
 void main()
 {
-    vec4 _99;
+    vec4 _101;
-    _99 = _24.in_data[gl_GlobalInvocationID.x];
+    _101 = _24.in_data[gl_GlobalInvocationID.x];
-    for (int _93 = 0; (_93 + 1) < 10; )
+    for (int _95 = 0; (_95 + 1) < 10; )
     {
-        _99 *= 2.0;
+        _101 *= 2.0;
-        _93 += 2;
+        _95 += 2;
         continue;
     }
-    vec4 _98;
+    vec4 _100;
-    _98 = _99;
+    _100 = _101;
-    vec4 _103;
+    vec4 _105;
-    for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++)
+    for (uint _96 = 0u; _96 < 16u; _100 = _105, _96++)
     {
-        _103 = _98;
+        _105 = _100;
-        for (uint _100 = 0u; _100 < 30u; )
+        for (uint _102 = 0u; _102 < 30u; )
         {
-            _103 = _24.mvp * _103;
+            _105 = _24.mvp * _105;
-            _100++;
+            _102++;
             continue;
         }
     }
-    _89.out_data[gl_GlobalInvocationID.x] = _98;
+    _89.out_data[gl_GlobalInvocationID.x] = _100;
 }
 

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;
+    _112.next.value = _50.head1.value + _50.head2.value;
-    Node param = n.next;
+    _50.head1.value = 20;
-    Node param_1 = _50.head1;
+    _50.head1.value = _50.head2.value * 10;
-    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);
 }
 

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);
+    i8vec2 _201 = unpack8(10s);
-    i8vec4 _199 = unpack8(_197);
+    i8vec2 _198 = _201;
-    _196 = pack16(_198);
+    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);
+    u8vec2 _225 = unpack8(10us);
-    u8vec4 _223 = unpack8(_221);
+    u8vec2 _222 = _225;
-    _220 = pack16(_222);
+    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);
+    i8vec4 _251 = _246 + i8vec4(registers.i8);
-    _244 += i8vec4(-40);
+    _244 = _251;
-    _244 += i8vec4(-50);
+    i8vec4 _254 = _251 + i8vec4(-40);
-    _244 += i8vec4(int8_t(10), int8_t(20), int8_t(30), int8_t(40));
+    _244 = _254;
-    _244 += i8vec4(ssbo.i8[4]);
+    i8vec4 _256 = _254 + i8vec4(-50);
-    _244 += i8vec4(ubo.i8);
+    _244 = _256;
-    FragColorInt = ivec4(_244);
+    i8vec4 _258 = _256 + i8vec4(int8_t(10), int8_t(20), int8_t(30), int8_t(40));
-    u8vec4 _271 = u8vec4(_246);
+    _244 = _258;
-    _271 += u8vec4(registers.u8);
+    i8vec4 _263 = _258 + i8vec4(ssbo.i8[4]);
-    _271 += u8vec4(216);
+    _244 = _263;
-    _271 += u8vec4(206);
+    i8vec4 _268 = _263 + i8vec4(ubo.i8);
-    _271 += u8vec4(uint8_t(10), uint8_t(20), uint8_t(30), uint8_t(40));
+    _244 = _268;
-    _271 += u8vec4(ssbo.u8[4]);
+    FragColorInt = ivec4(_268);
-    _271 += u8vec4(ubo.u8);
+    u8vec4 _274 = u8vec4(_246);
-    FragColorUint = uvec4(_271);
+    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);
 }
 

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)

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);
 

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);

reference/shaders-hlsl/comp/atomic.comp

@@ -1,3 +1,5 @@
+static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u);
+
 RWByteAddressBuffer ssbo : register(u2);
 RWTexture2D<uint> uImage : register(u0);
 RWTexture2D<int> iImage : register(u1);

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<float> uImageIn : register(u0);

reference/shaders-hlsl/comp/image.comp

@@ -1,3 +1,5 @@
+static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u);
+
 RWTexture2D<float> uImageInF : register(u0);
 RWTexture2D<float> uImageOutF : register(u1);
 RWTexture2D<int> uImageInI : register(u2);

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);
 

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)
 {

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)
 {

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);
 

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()

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)
 {

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()

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);
 

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()

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()

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;

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<typename T>
 inline T radians(T d)

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;

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;

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;

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);

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;

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]));

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;

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;

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;

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;

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;

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;

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);

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;