MSL: Handle signed atomic min/max.
C++ deduces this based on the pointer type, so cast to atomic_uint/int
if we have to.
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 0000000..3fdf46b
--- /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 0000000..832a273
--- /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 4abb437..3828fc4 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -7760,7 +7760,7 @@
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 @@
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 @@
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 @@
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 @@
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 @@
}
// 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 @@
}
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 @@
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 @@
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 f01ccea..aa6de74 100644
--- a/spirv_msl.hpp
+++ b/spirv_msl.hpp
@@ -920,8 +920,8 @@
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);
