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

Merge pull request #932 from billhollings/master

Browse Source 
MSL: Fix infinite CAS loop on atomic_compare_exchange_weak_explicit().
This commit is contained in:
Hans-Kristian Arntzen 2019-04-06 13:20:43 +02:00 committed by GitHub
commit 133ea8fd82
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
3 changed files with 15 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
reference/opt/shaders-msl/comp/atomic.comp
View File

@ -27,7 +27,7 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_32 = 10u;
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_uint*)&ssbo.u32, &_32, 2u, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_uint*)&ssbo.u32, &_32, 2u, memory_order_relaxed, memory_order_relaxed) && _32 == 10u);
int _36 = atomic_fetch_add_explicit((volatile device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _38 = atomic_fetch_or_explicit((volatile device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _40 = atomic_fetch_xor_explicit((volatile device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
@ -39,7 +39,7 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_52 = 10;
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_int*)&ssbo.i32, &_52, 2, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_int*)&ssbo.i32, &_52, 2, memory_order_relaxed, memory_order_relaxed) && _52 == 10);
shared_u32 = 10u;
shared_i32 = 10;
uint _57 = atomic_fetch_add_explicit((volatile threadgroup atomic_uint*)&shared_u32, 1u, memory_order_relaxed);
@ -53,7 +53,7 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_64 = 10u;
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_uint*)&shared_u32, &_64, 2u, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_uint*)&shared_u32, &_64, 2u, memory_order_relaxed, memory_order_relaxed) && _64 == 10u);
int _65 = atomic_fetch_add_explicit((volatile threadgroup atomic_int*)&shared_i32, 1, memory_order_relaxed);
int _66 = atomic_fetch_or_explicit((volatile threadgroup atomic_int*)&shared_i32, 1, memory_order_relaxed);
int _67 = atomic_fetch_xor_explicit((volatile threadgroup atomic_int*)&shared_i32, 1, memory_order_relaxed);
@ -65,6 +65,6 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_72 = 10;
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_int*)&shared_i32, &_72, 2, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_int*)&shared_i32, &_72, 2, memory_order_relaxed, memory_order_relaxed) && _72 == 10);
}

8
reference/shaders-msl/comp/atomic.comp
View File

@ -27,7 +27,7 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_32 = 10u;
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_uint*)&ssbo.u32, &_32, 2u, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_uint*)&ssbo.u32, &_32, 2u, memory_order_relaxed, memory_order_relaxed) && _32 == 10u);
int _36 = atomic_fetch_add_explicit((volatile device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _38 = atomic_fetch_or_explicit((volatile device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
int _40 = atomic_fetch_xor_explicit((volatile device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
@ -39,7 +39,7 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_52 = 10;
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_int*)&ssbo.i32, &_52, 2, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile device atomic_int*)&ssbo.i32, &_52, 2, memory_order_relaxed, memory_order_relaxed) && _52 == 10);
shared_u32 = 10u;
shared_i32 = 10;
uint _57 = atomic_fetch_add_explicit((volatile threadgroup atomic_uint*)&shared_u32, 1u, memory_order_relaxed);
@ -53,7 +53,7 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_64 = 10u;
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_uint*)&shared_u32, &_64, 2u, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_uint*)&shared_u32, &_64, 2u, memory_order_relaxed, memory_order_relaxed) && _64 == 10u);
int _65 = atomic_fetch_add_explicit((volatile threadgroup atomic_int*)&shared_i32, 1, memory_order_relaxed);
int _66 = atomic_fetch_or_explicit((volatile threadgroup atomic_int*)&shared_i32, 1, memory_order_relaxed);
int _67 = atomic_fetch_xor_explicit((volatile threadgroup atomic_int*)&shared_i32, 1, memory_order_relaxed);
@ -65,6 +65,6 @@ kernel void main0(device SSBO& ssbo [[buffer(2)]])
do
{
_72 = 10;
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_int*)&shared_i32, &_72, 2, memory_order_relaxed, memory_order_relaxed));
} while (!atomic_compare_exchange_weak_explicit((volatile threadgroup atomic_int*)&shared_i32, &_72, 2, memory_order_relaxed, memory_order_relaxed) && _72 == 10);
}

10
spirv_msl.cpp
View File

@ -4043,13 +4043,17 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
exp += get_memory_order(mem_order_2);
exp += ")";
// MSL only supports the weak atomic compare exchange,
// so emit a CAS loop here.
// MSL only supports the weak atomic compare exchange, so emit a CAS loop here.
// The MSL function returns false if the atomic write fails OR the comparison test fails,
// so we must validate that it wasn't the comparison test that failed before continuing
// the CAS loop, otherwise it will loop infinitely, with the comparison test always failing.
// The function updates the comparitor value from the memory value, so the additional
// comparison test evaluates the memory value against the expected value.
statement(variable_decl(type, to_name(result_id)), ";");
statement("do");
begin_scope();
statement(to_name(result_id), " = ", to_expression(op1), ";");
end_scope_decl(join("while (!", exp, ")"));
end_scope_decl(join("while (!", exp, " && ", to_name(result_id), " == ", to_enclosed_expression(op1), ")"));
set<SPIRExpression>(result_id, to_name(result_id), result_type, true);
}
else