diff --git a/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp b/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp index 1cf531c4..1e2880f4 100644 --- a/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp +++ b/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)]]) { diff --git a/reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp b/reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp index 8c56e7af..473298c2 100644 --- a/reference/opt/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp +++ b/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]]) { diff --git a/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp b/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp index 7b7e9120..c69204ed 100644 --- a/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp +++ b/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 u0 [[texture(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) { diff --git a/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp b/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp index 03794023..4738c355 100644 --- a/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp +++ b/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 u0 [[texture(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) { diff --git a/reference/opt/shaders-msl/comp/barriers.comp b/reference/opt/shaders-msl/comp/barriers.comp index 23a19914..8822d080 100644 --- a/reference/opt/shaders-msl/comp/barriers.comp +++ b/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() { diff --git a/reference/opt/shaders-msl/comp/builtins.comp b/reference/opt/shaders-msl/comp/builtins.comp index 189bb1b7..cbb76c9c 100644 --- a/reference/opt/shaders-msl/comp/builtins.comp +++ b/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]]) { diff --git a/reference/opt/shaders-msl/comp/composite-array-initialization.comp b/reference/opt/shaders-msl/comp/composite-array-initialization.comp index 8b1faa10..55f79667 100644 --- a/reference/opt/shaders-msl/comp/composite-array-initialization.comp +++ b/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 } }; diff --git a/reference/opt/shaders-msl/comp/culling.comp b/reference/opt/shaders-msl/comp/culling.comp index edc01a87..95ffff83 100644 --- a/reference/opt/shaders-msl/comp/culling.comp +++ b/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]]) { diff --git a/reference/opt/shaders-msl/comp/packing-test-1.comp b/reference/opt/shaders-msl/comp/packing-test-1.comp index 24133543..13101773 100644 --- a/reference/opt/shaders-msl/comp/packing-test-1.comp +++ b/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]]) { diff --git a/reference/opt/shaders-msl/comp/packing-test-2.comp b/reference/opt/shaders-msl/comp/packing-test-2.comp index 460f1276..b3b89d76 100644 --- a/reference/opt/shaders-msl/comp/packing-test-2.comp +++ b/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]]) { diff --git a/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp b/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp index 86d77e59..ddb8c962 100644 --- a/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp +++ b/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]]) { diff --git a/reference/opt/shaders-msl/comp/shared.comp b/reference/opt/shaders-msl/comp/shared.comp index ca8420f7..029ab539 100644 --- a/reference/opt/shaders-msl/comp/shared.comp +++ b/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]]) { diff --git a/reference/opt/shaders-msl/comp/spec-constant-work-group-size.comp b/reference/opt/shaders-msl/comp/spec-constant-work-group-size.comp index b875e1bb..bb796ab9 100644 --- a/reference/opt/shaders-msl/comp/spec-constant-work-group-size.comp +++ b/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; diff --git a/reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp b/reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp index 1cf531c4..1e2880f4 100644 --- a/reference/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp +++ b/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)]]) { diff --git a/reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp b/reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp index 8c56e7af..473298c2 100644 --- a/reference/shaders-msl/asm/comp/storage-buffer-basic.invalid.asm.comp +++ b/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]]) { diff --git a/reference/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp b/reference/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp index a16d019c..8a9e4450 100644 --- a/reference/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp +++ b/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) { diff --git a/reference/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp b/reference/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp index 0ece12c1..17766250 100644 --- a/reference/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp +++ b/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 u0 [[texture(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) { diff --git a/reference/shaders-msl/comp/barriers.comp b/reference/shaders-msl/comp/barriers.comp index 3b8474c7..fdaf8747 100644 --- a/reference/shaders-msl/comp/barriers.comp +++ b/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() { diff --git a/reference/shaders-msl/comp/builtins.comp b/reference/shaders-msl/comp/builtins.comp index 4330d578..4f03b099 100644 --- a/reference/shaders-msl/comp/builtins.comp +++ b/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]]) { diff --git a/reference/shaders-msl/comp/composite-array-initialization.comp b/reference/shaders-msl/comp/composite-array-initialization.comp index 37a36b7f..9d35115e 100644 --- a/reference/shaders-msl/comp/composite-array-initialization.comp +++ b/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 } }; diff --git a/reference/shaders-msl/comp/culling.comp b/reference/shaders-msl/comp/culling.comp index e6d8f44a..32acf599 100644 --- a/reference/shaders-msl/comp/culling.comp +++ b/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]]) { diff --git a/reference/shaders-msl/comp/packing-test-1.comp b/reference/shaders-msl/comp/packing-test-1.comp index 06160ebc..e9b10736 100644 --- a/reference/shaders-msl/comp/packing-test-1.comp +++ b/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]]) { diff --git a/reference/shaders-msl/comp/packing-test-2.comp b/reference/shaders-msl/comp/packing-test-2.comp index cdf909ff..61a4970c 100644 --- a/reference/shaders-msl/comp/packing-test-2.comp +++ b/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]]) { diff --git a/reference/shaders-msl/comp/shared-array-of-arrays.comp b/reference/shaders-msl/comp/shared-array-of-arrays.comp index 9a09e297..7acb0ab8 100644 --- a/reference/shaders-msl/comp/shared-array-of-arrays.comp +++ b/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) { diff --git a/reference/shaders-msl/comp/shared.comp b/reference/shaders-msl/comp/shared.comp index 3648639c..c9ffde15 100644 --- a/reference/shaders-msl/comp/shared.comp +++ b/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]]) { diff --git a/reference/shaders-msl/comp/spec-constant-work-group-size.comp b/reference/shaders-msl/comp/spec-constant-work-group-size.comp index b875e1bb..bb796ab9 100644 --- a/reference/shaders-msl/comp/spec-constant-work-group-size.comp +++ b/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; diff --git a/spirv_msl.cpp b/spirv_msl.cpp index e302a77b..b8957366 100644 --- a/spirv_msl.cpp +++ b/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(workgroup_size_id)), ";"); emitted = true; }