MSL: Handle signed atomic min/max.

C++ deduces this based on the pointer type, so cast to atomic_uint/int
if we have to.

4 files changed, 111 insertions, 11 deletions

diff --git a/reference/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp b/reference/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp
new file mode 100644
index 00000000..3fdf46bb
--- /dev/null
+++ b/reference/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp
@@ -0,0 +1,28 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    uint a;
+    int b;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device SSBO& _4 [[buffer(0)]])
+{
+    uint _26 = atomic_fetch_max_explicit((device atomic_uint*)&_4.a, 1u, memory_order_relaxed);
+    uint _27 = uint(atomic_fetch_min_explicit((device atomic_int*)&_4.a, int(1u), memory_order_relaxed));
+    uint _28 = atomic_fetch_min_explicit((device atomic_uint*)&_4.a, 4294967295u, memory_order_relaxed);
+    uint _29 = uint(atomic_fetch_max_explicit((device atomic_int*)&_4.a, int(4294967295u), memory_order_relaxed));
+    int _30 = atomic_fetch_max_explicit((device atomic_int*)&_4.b, -3, memory_order_relaxed);
+    int _31 = int(atomic_fetch_min_explicit((device atomic_uint*)&_4.b, uint(-3), memory_order_relaxed));
+    int _32 = atomic_fetch_min_explicit((device atomic_int*)&_4.b, 4, memory_order_relaxed);
+    int _33 = int(atomic_fetch_max_explicit((device atomic_uint*)&_4.b, uint(4), memory_order_relaxed));
+}
+
diff --git a/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp b/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp
new file mode 100644
index 00000000..832a2735
--- /dev/null
+++ b/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp
@@ -0,0 +1,56 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 10
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "a"
+               OpMemberName %SSBO 1 "b"
+               OpName %_ ""
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpMemberDecorate %SSBO 1 Offset 4
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+        %int = OpTypeInt 32 1
+       %SSBO = OpTypeStruct %uint %int
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+      %int_0 = OpConstant %int 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%uint_4294967295 = OpConstant %uint 4294967295
+      %int_1 = OpConstant %int 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+     %int_n3 = OpConstant %int -3
+      %int_4 = OpConstant %int 4
+     %v3uint = OpTypeVector %uint 3
+%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %13 = OpAccessChain %_ptr_Uniform_uint %_ %int_0
+         %18 = OpAccessChain %_ptr_Uniform_uint %_ %int_0
+         %22 = OpAccessChain %_ptr_Uniform_int %_ %int_1
+         %25 = OpAccessChain %_ptr_Uniform_int %_ %int_1
+         %30 = OpAtomicUMax %uint %13 %uint_1 %uint_0 %uint_1
+         %31 = OpAtomicSMin %uint %13 %uint_1 %uint_0 %uint_1
+         %32 = OpAtomicUMin %uint %18 %uint_1 %uint_0 %uint_4294967295
+         %33 = OpAtomicSMax %uint %18 %uint_1 %uint_0 %uint_4294967295
+         %34 = OpAtomicSMax %int %22 %uint_1 %uint_0 %int_n3
+         %35 = OpAtomicUMin %int %22 %uint_1 %uint_0 %int_n3
+         %36 = OpAtomicSMin %int %25 %uint_1 %uint_0 %int_4
+         %37 = OpAtomicUMax %int %25 %uint_1 %uint_0 %int_4
+               OpReturn
+               OpFunctionEnd
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index 4abb437f..f3d33d5f 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -7760,7 +7760,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		uint32_t ptr = ops[2];
 		uint32_t mem_sem = ops[4];
 		uint32_t val = ops[5];
-		emit_atomic_func_op(result_type, id, "atomic_exchange_explicit", mem_sem, mem_sem, false, ptr, val);
+		emit_atomic_func_op(result_type, id, "atomic_exchange_explicit", opcode, mem_sem, mem_sem, false, ptr, val);
 		break;
 	}
 
@@ -7773,7 +7773,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		uint32_t mem_sem_fail = ops[5];
 		uint32_t val = ops[6];
 		uint32_t comp = ops[7];
-		emit_atomic_func_op(result_type, id, "atomic_compare_exchange_weak_explicit", mem_sem_pass, mem_sem_fail, true,
+		emit_atomic_func_op(result_type, id, "atomic_compare_exchange_weak_explicit", opcode,
+		                    mem_sem_pass, mem_sem_fail, true,
 		                    ptr, comp, true, false, val);
 		break;
 	}
@@ -7787,7 +7788,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		uint32_t id = ops[1];
 		uint32_t ptr = ops[2];
 		uint32_t mem_sem = ops[4];
-		emit_atomic_func_op(result_type, id, "atomic_load_explicit", mem_sem, mem_sem, false, ptr, 0);
+		emit_atomic_func_op(result_type, id, "atomic_load_explicit", opcode, mem_sem, mem_sem, false, ptr, 0);
 		break;
 	}
 
@@ -7798,7 +7799,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		uint32_t ptr = ops[0];
 		uint32_t mem_sem = ops[2];
 		uint32_t val = ops[3];
-		emit_atomic_func_op(result_type, id, "atomic_store_explicit", mem_sem, mem_sem, false, ptr, val);
+		emit_atomic_func_op(result_type, id, "atomic_store_explicit", opcode, mem_sem, mem_sem, false, ptr, val);
 		break;
 	}
 
@@ -7810,7 +7811,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		uint32_t ptr = ops[2];                                                                                   \
 		uint32_t mem_sem = ops[4];                                                                               \
 		uint32_t val = valsrc;                                                                                   \
-		emit_atomic_func_op(result_type, id, "atomic_fetch_" #op "_explicit", mem_sem, mem_sem, false, ptr, val, \
+		emit_atomic_func_op(result_type, id, "atomic_fetch_" #op "_explicit", opcode,                            \
+		                    mem_sem, mem_sem, false, ptr, val,                                                   \
 		                    false, valconst);                                                                    \
 	} while (false)
 
@@ -8796,13 +8798,22 @@ bool CompilerMSL::maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs)
 }
 
 // Emits one of the atomic functions. In MSL, the atomic functions operate on pointers
-void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, uint32_t mem_order_1,
-                                      uint32_t mem_order_2, bool has_mem_order_2, uint32_t obj, uint32_t op1,
+void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, Op opcode,
+                                      uint32_t mem_order_1, uint32_t mem_order_2, bool has_mem_order_2, uint32_t obj, uint32_t op1,
                                       bool op1_is_pointer, bool op1_is_literal, uint32_t op2)
 {
 	string exp = string(op) + "(";
 
 	auto &type = get_pointee_type(expression_type(obj));
+	auto expected_type = type.basetype;
+	if (opcode == OpAtomicUMax || opcode == OpAtomicUMin)
+		expected_type = to_unsigned_basetype(type.width);
+	else if (opcode == OpAtomicSMax || opcode == OpAtomicSMin)
+		expected_type = to_signed_basetype(type.width);
+
+	auto remapped_type = type;
+	remapped_type.basetype = expected_type;
+
 	exp += "(";
 	auto *var = maybe_get_backing_variable(obj);
 	if (!var)
@@ -8820,7 +8831,9 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
 	}
 
 	exp += " atomic_";
-	exp += type_to_glsl(type);
+	// For signed and unsigned min/max, we can signal this through the pointer type.
+	// There is no other way, since C++ does not have explicit signage for atomics.
+	exp += type_to_glsl(remapped_type);
 	exp += "*)";
 
 	exp += "&";
@@ -8863,7 +8876,7 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
 			if (op1_is_literal)
 				exp += join(", ", op1);
 			else
-				exp += ", " + to_expression(op1);
+				exp += ", " + bitcast_expression(expected_type, op1);
 		}
 		if (op2)
 			exp += ", " + to_expression(op2);
@@ -8874,6 +8887,9 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
 
 		exp += ")";
 
+		if (expected_type != type.basetype)
+			exp = bitcast_expression(type, expected_type, exp);
+
 		if (strcmp(op, "atomic_store_explicit") != 0)
 			emit_op(result_type, result_id, exp, false);
 		else
diff --git a/spirv_msl.hpp b/spirv_msl.hpp
index f01cceaf..e065519c 100644
--- a/spirv_msl.hpp
+++ b/spirv_msl.hpp
@@ -920,8 +920,8 @@ protected:
 	std::string get_tess_factor_struct_name();
 	SPIRType &get_uint_type();
 	uint32_t get_uint_type_id();
-	void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, uint32_t mem_order_1,
-	                         uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
+	void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, spv::Op opcode,
+	                         uint32_t mem_order_1, uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
 	                         bool op1_is_pointer = false, bool op1_is_literal = false, uint32_t op2 = 0);
 	const char *get_memory_order(uint32_t spv_mem_sem);
 	void add_pragma_line(const std::string &line);