MSL: Declare gl_WorkGroupSize constant with [[maybe_unused]].

Avoids ugly warnings on nearly every compute shader.
We could do analysis to detect whether we need to emit this constant,
but it's a bit tedious to figure out if an OpConstantComponent is
actually used by opcodes, so just make it simple.

reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp

@@ -12,7 +12,7 @@ constant uint _5_tmp [[function_constant(10)]];
 constant uint _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 9u;
 constant uint _6_tmp [[function_constant(12)]];
 constant uint _6 = is_function_constant_defined(_6_tmp) ? _6_tmp : 4u;
-constant uint3 gl_WorkGroupSize = uint3(_5, 20u, _6);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_5, 20u, _6);
 
 kernel void main0(device SSBO& _4 [[buffer(0)]])
 {

reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp

@@ -12,7 +12,7 @@ constant uint _3_tmp [[function_constant(0)]];
 constant uint _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 1u;
 constant uint _4_tmp [[function_constant(2)]];
 constant uint _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 3u;
-constant uint3 gl_WorkGroupSize = uint3(_3, 2u, _4);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_3, 2u, _4);
 
 kernel void main0(device _6& _8 [[buffer(0)]], device _6& _9 [[buffer(1)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {

reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp

@@ -8,7 +8,7 @@ struct cb1_struct
     float4 _m0[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(16u, 16u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 16u, 1u);
 
 kernel void main0(constant cb1_struct& cb0_1 [[buffer(0)]], texture2d<float, access::write> u0 [[texture(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {

reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp

@@ -8,7 +8,7 @@ struct cb1_struct
     float4 _m0[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(16u, 16u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 16u, 1u);
 
 kernel void main0(constant cb1_struct& cb0_1 [[buffer(0)]], texture2d<float, access::write> u0 [[texture(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {

reference/opt/shaders-msl/comp/barriers.comp

@@ -3,7 +3,7 @@
 
 using namespace metal;
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 kernel void main0()
 {

reference/opt/shaders-msl/comp/builtins.comp

@@ -3,7 +3,7 @@
 
 using namespace metal;
 
-constant uint3 gl_WorkGroupSize = uint3(8u, 4u, 2u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 4u, 2u);
 
 kernel void main0(uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {

reference/opt/shaders-msl/comp/composite-array-initialization.comp

@@ -25,7 +25,7 @@ struct SSBO
     Data outdata[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
 
 constant Data_1 _25[2] = { Data_1{ 1.0, 2.0 }, Data_1{ 3.0, 4.0 } };
 

reference/opt/shaders-msl/comp/culling.comp

@@ -21,7 +21,7 @@ struct SSBO3
     uint count;
 };
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {

reference/opt/shaders-msl/comp/packing-test-1.comp

@@ -19,7 +19,7 @@ struct Buffer1
     float buf1[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(32u, 1u, 1u);
 
 kernel void main0(device Buffer0& _15 [[buffer(1)]], device Buffer1& _34 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {

reference/opt/shaders-msl/comp/packing-test-2.comp

@@ -19,7 +19,7 @@ struct Buffer1
     float buf1[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(32u, 1u, 1u);
 
 kernel void main0(device Buffer0& _14 [[buffer(1)]], device Buffer1& _24 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {

reference/opt/shaders-msl/comp/shared-array-of-arrays.comp

@@ -8,7 +8,7 @@ struct SSBO
     float out_data[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 4u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 4u, 1u);
 
 kernel void main0(device SSBO& _67 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {

reference/opt/shaders-msl/comp/shared.comp

@@ -13,7 +13,7 @@ struct SSBO2
     float out_data[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
 {

reference/opt/shaders-msl/comp/spec-constant-work-group-size.comp

@@ -18,7 +18,7 @@ struct SSBO
 constant uint _21 = (uint(a) + 0u);
 constant uint _22_tmp [[function_constant(10)]];
 constant uint _22 = is_function_constant_defined(_22_tmp) ? _22_tmp : 1u;
-constant uint3 gl_WorkGroupSize = uint3(_22, 20u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_22, 20u, 1u);
 constant uint _27 = gl_WorkGroupSize.x;
 constant uint _28 = (_21 + _27);
 constant uint _29 = gl_WorkGroupSize.y;

reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp

@@ -12,7 +12,7 @@ constant uint _5_tmp [[function_constant(10)]];
 constant uint _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 9u;
 constant uint _6_tmp [[function_constant(12)]];
 constant uint _6 = is_function_constant_defined(_6_tmp) ? _6_tmp : 4u;
-constant uint3 gl_WorkGroupSize = uint3(_5, 20u, _6);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_5, 20u, _6);
 
 kernel void main0(device SSBO& _4 [[buffer(0)]])
 {

reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp

@@ -12,7 +12,7 @@ constant uint _3_tmp [[function_constant(0)]];
 constant uint _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 1u;
 constant uint _4_tmp [[function_constant(2)]];
 constant uint _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 3u;
-constant uint3 gl_WorkGroupSize = uint3(_3, 2u, _4);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_3, 2u, _4);
 
 kernel void main0(device _6& _8 [[buffer(0)]], device _6& _9 [[buffer(1)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {

reference/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp

@@ -10,7 +10,7 @@ struct cb1_struct
     float4 _m0[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(16u, 16u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 16u, 1u);
 
 int2 get_texcoord(thread const int2& base, thread const int2& index, thread uint3& gl_LocalInvocationID)
 {

reference/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp

@@ -8,7 +8,7 @@ struct cb1_struct
     float4 _m0[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(16u, 16u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 16u, 1u);
 
 kernel void main0(constant cb1_struct& cb0_1 [[buffer(0)]], texture2d<float, access::write> u0 [[texture(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {

reference/shaders-msl/comp/barriers.comp

@@ -5,7 +5,7 @@
 
 using namespace metal;
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 void barrier_shared()
 {

reference/shaders-msl/comp/builtins.comp

@@ -3,7 +3,7 @@
 
 using namespace metal;
 
-constant uint3 gl_WorkGroupSize = uint3(8u, 4u, 2u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 4u, 2u);
 
 kernel void main0(uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {

reference/shaders-msl/comp/composite-array-initialization.comp

@@ -25,7 +25,7 @@ struct SSBO
     Data outdata[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
 
 constant Data_1 _25[2] = { Data_1{ 1.0, 2.0 }, Data_1{ 3.0, 4.0 } };
 

reference/shaders-msl/comp/culling.comp

@@ -21,7 +21,7 @@ struct SSBO3
     uint count;
 };
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {

reference/shaders-msl/comp/packing-test-1.comp

@@ -25,7 +25,7 @@ struct Buffer1
     float buf1[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(32u, 1u, 1u);
 
 kernel void main0(device Buffer0& _15 [[buffer(1)]], device Buffer1& _34 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {

reference/shaders-msl/comp/packing-test-2.comp

@@ -19,7 +19,7 @@ struct Buffer1
     float buf1[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(32u, 1u, 1u);
 
 kernel void main0(device Buffer0& _14 [[buffer(1)]], device Buffer1& _24 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {

reference/shaders-msl/comp/shared-array-of-arrays.comp

@@ -10,7 +10,7 @@ struct SSBO
     float out_data[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 4u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 4u, 1u);
 
 void work(threadgroup float (&foo)[4][4], thread uint3& gl_LocalInvocationID, thread uint& gl_LocalInvocationIndex, device SSBO& v_67, thread uint3& gl_GlobalInvocationID)
 {

reference/shaders-msl/comp/shared.comp

@@ -13,7 +13,7 @@ struct SSBO2
     float out_data[1];
 };
 
-constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
 {

reference/shaders-msl/comp/spec-constant-work-group-size.comp

@@ -18,7 +18,7 @@ struct SSBO
 constant uint _21 = (uint(a) + 0u);
 constant uint _22_tmp [[function_constant(10)]];
 constant uint _22 = is_function_constant_defined(_22_tmp) ? _22_tmp : 1u;
-constant uint3 gl_WorkGroupSize = uint3(_22, 20u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_22, 20u, 1u);
 constant uint _27 = gl_WorkGroupSize.x;
 constant uint _28 = (_21 + _27);
 constant uint _29 = gl_WorkGroupSize.y;

spirv_msl.cpp

@@ -2952,7 +2952,7 @@ void CompilerMSL::emit_specialization_constants_and_structs()
 				// TODO: This can be expressed as a [[threads_per_threadgroup]] input semantic, but we need to know
 				// the work group size at compile time in SPIR-V, and [[threads_per_threadgroup]] would need to be passed around as a global.
 				// The work group size may be a specialization constant.
-				statement("constant uint3 ", builtin_to_glsl(BuiltInWorkgroupSize, StorageClassWorkgroup), " = ",
+				statement("constant uint3 ", builtin_to_glsl(BuiltInWorkgroupSize, StorageClassWorkgroup), " [[maybe_unused]] = ",
 				          constant_expression(get<SPIRConstant>(workgroup_size_id)), ";");
 				emitted = true;
 			}