From 41007cdc7d9c3a1f9b61e307d149f94a68fa3c9c Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Thu, 16 Feb 2023 14:56:38 -0800 Subject: [PATCH] MSL: Deduplicate function constants. It is possible in SPIR-V to declare multiple specialization constants with the same constant ID. The most common cause of this in GLSL is defining a spec constant, then declaring the workgroup size to use that spec constant by its ID. But, MSL forbids defining multiple function constants with the same function constant ID. So, we must only emit one definition of the actual function constant (with the `[[function_constant(id)]]` attribute); but we can point the other variables at this one definition. Fixes three tests in the Vulkan CTS under `dEQP-VK.compute.basic.max_local_size_*`. --- .../asm/comp/duplicate-spec-id.asm.comp | 20 +++++++ .../comp/local-size-duplicate-spec-id.comp | 23 ++++++++ .../asm/comp/duplicate-spec-id.asm.comp | 20 +++++++ .../comp/local-size-duplicate-spec-id.comp | 26 +++++++++ .../asm/comp/duplicate-spec-id.asm.comp | 54 +++++++++++++++++++ .../comp/local-size-duplicate-spec-id.comp | 15 ++++++ spirv_msl.cpp | 24 +++++++-- 7 files changed, 177 insertions(+), 5 deletions(-) create mode 100644 reference/opt/shaders-msl/asm/comp/duplicate-spec-id.asm.comp create mode 100644 reference/opt/shaders-msl/comp/local-size-duplicate-spec-id.comp create mode 100644 reference/shaders-msl/asm/comp/duplicate-spec-id.asm.comp create mode 100644 reference/shaders-msl/comp/local-size-duplicate-spec-id.comp create mode 100644 shaders-msl/asm/comp/duplicate-spec-id.asm.comp create mode 100644 shaders-msl/comp/local-size-duplicate-spec-id.comp diff --git a/reference/opt/shaders-msl/asm/comp/duplicate-spec-id.asm.comp b/reference/opt/shaders-msl/asm/comp/duplicate-spec-id.asm.comp new file mode 100644 index 00000000..90bebdd6 --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/duplicate-spec-id.asm.comp @@ -0,0 +1,20 @@ +#include +#include + +using namespace metal; + +struct StorageBuffer +{ + float values[1]; +}; + +constant int foo_tmp [[function_constant(0)]]; +constant int foo = is_function_constant_defined(foo_tmp) ? foo_tmp : 1; +constant float bar = is_function_constant_defined(foo_tmp) ? as_type(foo_tmp) : 2.0; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +kernel void main0(device StorageBuffer& ssbo [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]]) +{ + ssbo.values[gl_LocalInvocationIndex] = float(foo) + bar; +} + diff --git a/reference/opt/shaders-msl/comp/local-size-duplicate-spec-id.comp b/reference/opt/shaders-msl/comp/local-size-duplicate-spec-id.comp new file mode 100644 index 00000000..19a56fc9 --- /dev/null +++ b/reference/opt/shaders-msl/comp/local-size-duplicate-spec-id.comp @@ -0,0 +1,23 @@ +#include +#include + +using namespace metal; + +struct StorageBuffer +{ + uint values[1]; +}; + +constant uint _22_tmp [[function_constant(0)]]; +constant uint _22 = is_function_constant_defined(_22_tmp) ? _22_tmp : 1u; +constant uint _23_tmp [[function_constant(1)]]; +constant uint _23 = is_function_constant_defined(_23_tmp) ? _23_tmp : 1u; +constant uint _24_tmp [[function_constant(2)]]; +constant uint _24 = is_function_constant_defined(_24_tmp) ? _24_tmp : 1u; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_22, _23, _24); + +kernel void main0(device StorageBuffer& ssbo [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]]) +{ + ssbo.values[gl_LocalInvocationIndex] = 1u; +} + diff --git a/reference/shaders-msl/asm/comp/duplicate-spec-id.asm.comp b/reference/shaders-msl/asm/comp/duplicate-spec-id.asm.comp new file mode 100644 index 00000000..90bebdd6 --- /dev/null +++ b/reference/shaders-msl/asm/comp/duplicate-spec-id.asm.comp @@ -0,0 +1,20 @@ +#include +#include + +using namespace metal; + +struct StorageBuffer +{ + float values[1]; +}; + +constant int foo_tmp [[function_constant(0)]]; +constant int foo = is_function_constant_defined(foo_tmp) ? foo_tmp : 1; +constant float bar = is_function_constant_defined(foo_tmp) ? as_type(foo_tmp) : 2.0; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +kernel void main0(device StorageBuffer& ssbo [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]]) +{ + ssbo.values[gl_LocalInvocationIndex] = float(foo) + bar; +} + diff --git a/reference/shaders-msl/comp/local-size-duplicate-spec-id.comp b/reference/shaders-msl/comp/local-size-duplicate-spec-id.comp new file mode 100644 index 00000000..9e8e901a --- /dev/null +++ b/reference/shaders-msl/comp/local-size-duplicate-spec-id.comp @@ -0,0 +1,26 @@ +#include +#include + +using namespace metal; + +struct StorageBuffer +{ + uint values[1]; +}; + +constant int local_size_x_val_tmp [[function_constant(0)]]; +constant int local_size_x_val = is_function_constant_defined(local_size_x_val_tmp) ? local_size_x_val_tmp : 1; +constant int local_size_y_val_tmp [[function_constant(1)]]; +constant int local_size_y_val = is_function_constant_defined(local_size_y_val_tmp) ? local_size_y_val_tmp : 1; +constant int local_size_z_val_tmp [[function_constant(2)]]; +constant int local_size_z_val = is_function_constant_defined(local_size_z_val_tmp) ? local_size_z_val_tmp : 1; +constant uint _22 = is_function_constant_defined(local_size_x_val_tmp) ? uint(local_size_x_val_tmp) : 1u; +constant uint _23 = is_function_constant_defined(local_size_y_val_tmp) ? uint(local_size_y_val_tmp) : 1u; +constant uint _24 = is_function_constant_defined(local_size_z_val_tmp) ? uint(local_size_z_val_tmp) : 1u; +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_22, _23, _24); + +kernel void main0(device StorageBuffer& ssbo [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]]) +{ + ssbo.values[gl_LocalInvocationIndex] = 1u; +} + diff --git a/shaders-msl/asm/comp/duplicate-spec-id.asm.comp b/shaders-msl/asm/comp/duplicate-spec-id.asm.comp new file mode 100644 index 00000000..4a5aa3d8 --- /dev/null +++ b/shaders-msl/asm/comp/duplicate-spec-id.asm.comp @@ -0,0 +1,54 @@ +; SPIR-V +; Version: 1.3 +; Generator: Khronos Glslang Reference Front End; 11 +; Bound: 26 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_LocalInvocationIndex + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %main "main" + OpName %StorageBuffer "StorageBuffer" + OpMemberName %StorageBuffer 0 "values" + OpName %ssbo "ssbo" + OpName %gl_LocalInvocationIndex "gl_LocalInvocationIndex" + OpName %foo "foo" + OpName %bar "bar" + OpDecorate %_runtimearr_float ArrayStride 4 + OpMemberDecorate %StorageBuffer 0 Offset 0 + OpDecorate %StorageBuffer Block + OpDecorate %ssbo DescriptorSet 0 + OpDecorate %ssbo Binding 0 + OpDecorate %gl_LocalInvocationIndex BuiltIn LocalInvocationIndex + OpDecorate %foo SpecId 0 + OpDecorate %bar SpecId 0 + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 +%_runtimearr_float = OpTypeRuntimeArray %float +%StorageBuffer = OpTypeStruct %_runtimearr_float +%_ptr_StorageBuffer_StorageBuffer = OpTypePointer StorageBuffer %StorageBuffer + %ssbo = OpVariable %_ptr_StorageBuffer_StorageBuffer StorageBuffer + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%gl_LocalInvocationIndex = OpVariable %_ptr_Input_uint Input + %foo = OpSpecConstant %int 1 + %bar = OpSpecConstant %float 2 +%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 +%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %main = OpFunction %void None %3 + %5 = OpLabel + %16 = OpLoad %uint %gl_LocalInvocationIndex + %18 = OpConvertSToF %float %foo + %20 = OpFAdd %float %18 %bar + %22 = OpAccessChain %_ptr_StorageBuffer_float %ssbo %int_0 %16 + OpStore %22 %20 + OpReturn + OpFunctionEnd diff --git a/shaders-msl/comp/local-size-duplicate-spec-id.comp b/shaders-msl/comp/local-size-duplicate-spec-id.comp new file mode 100644 index 00000000..060858b9 --- /dev/null +++ b/shaders-msl/comp/local-size-duplicate-spec-id.comp @@ -0,0 +1,15 @@ +#version 450 + +layout(constant_id=0) const int local_size_x_val = 1; +layout(constant_id=1) const int local_size_y_val = 1; +layout(constant_id=2) const int local_size_z_val = 1; + +layout(local_size_x_id=0, local_size_y_id=1, local_size_z_id=2) in; + +layout(set=0, binding=0) buffer StorageBuffer { + uint values[]; +} ssbo; + +void main() { + ssbo.values[gl_LocalInvocationIndex] = 1u; +} diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 5f4c0add..4652d963 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -7344,6 +7344,15 @@ void CompilerMSL::emit_specialization_constants_and_structs() emitted = false; declared_structs.clear(); + // It is possible to have multiple spec constants that use the same spec constant ID. + // The most common cause of this is defining spec constants in GLSL while also declaring + // the workgroup size to use those spec constants. But, Metal forbids declaring more than + // one variable with the same function constant ID. + // In this case, we must only declare one variable with the [[function_constant(id)]] + // attribute, and use its initializer to initialize all the spec constants with + // that ID. + std::unordered_map unique_func_constants; + for (auto &id_ : ir.ids_for_constant_undef_or_type) { auto &id = ir.ids[id_]; @@ -7367,7 +7376,11 @@ void CompilerMSL::emit_specialization_constants_and_structs() string sc_type_name = type_to_glsl(type); add_resource_name(c.self); string sc_name = to_name(c.self); - string sc_tmp_name = sc_name + "_tmp"; + uint32_t constant_id = get_decoration(c.self, DecorationSpecId); + if (!unique_func_constants.count(constant_id)) + unique_func_constants.insert(make_pair(constant_id, c.self)); + SPIRType::BaseType sc_tmp_type = expression_type(unique_func_constants[constant_id]).basetype; + string sc_tmp_name = to_name(unique_func_constants[constant_id]) + "_tmp"; // Function constants are only supported in MSL 1.2 and later. // If we don't support it just declare the "default" directly. @@ -7377,12 +7390,13 @@ void CompilerMSL::emit_specialization_constants_and_structs() if (msl_options.supports_msl_version(1, 2) && has_decoration(c.self, DecorationSpecId) && !c.is_used_as_array_length) { - uint32_t constant_id = get_decoration(c.self, DecorationSpecId); // Only scalar, non-composite values can be function constants. - statement("constant ", sc_type_name, " ", sc_tmp_name, " [[function_constant(", constant_id, - ")]];"); + if (unique_func_constants[constant_id] == c.self) + statement("constant ", sc_type_name, " ", sc_tmp_name, " [[function_constant(", constant_id, + ")]];"); statement("constant ", sc_type_name, " ", sc_name, " = is_function_constant_defined(", sc_tmp_name, - ") ? ", sc_tmp_name, " : ", constant_expression(c), ";"); + ") ? ", bitcast_expression(type, sc_tmp_type, sc_tmp_name), " : ", constant_expression(c), + ";"); } else if (has_decoration(c.self, DecorationSpecId)) {