AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

MSL: Handle signed atomic min/max.

Browse Source 
C++ deduces this based on the pointer type, so cast to atomic_uint/int
if we have to.
This commit is contained in:
Hans-Kristian Arntzen 2022-01-17 15:29:13 +01:00
parent 7dd974b9db
commit 5a5be7f9b9
4 changed files with 111 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

28
reference/shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp Normal file
View File

@ -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));
}

56
shaders-msl-no-opt/asm/comp/atomic-min-max-sign.asm.comp Normal file
View File

@ -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

34
spirv_msl.cpp
View File

@ -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

4
spirv_msl.hpp
View File

@ -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);