Merge pull request #2108 from cdavis5e/msl-duplicate-spec-id

MSL: Deduplicate function constants.

reference/opt/shaders-msl/asm/comp/duplicate-spec-id.asm.comp

@@ -0,0 +1,20 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+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<float>(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;
+}
+

reference/opt/shaders-msl/comp/local-size-duplicate-spec-id.comp

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+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;
+}
+

reference/shaders-msl/asm/comp/duplicate-spec-id.asm.comp

@@ -0,0 +1,20 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+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<float>(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;
+}
+

reference/shaders-msl/comp/local-size-duplicate-spec-id.comp

@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+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;
+}
+

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

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

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<uint32_t, ConstantID> 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))
 				{