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_*`.

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