From 5a5be7f9b9f5a140decfd652f81616e728936e48 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Mon, 17 Jan 2022 15:29:13 +0100 Subject: [PATCH] MSL: Handle signed atomic min/max. C++ deduces this based on the pointer type, so cast to atomic_uint/int if we have to. --- .../asm/comp/atomic-min-max-sign.asm.comp | 28 ++++++++++ .../asm/comp/atomic-min-max-sign.asm.comp | 56 +++++++++++++++++++ spirv_msl.cpp | 34 ++++++++--- spirv_msl.hpp | 4 +- 4 files changed, 111 insertions(+), 11 deletions(-) create mode 100644 reference/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp create mode 100644 shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp 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 +#include +#include + +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);