mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-08-05 19:55:37 +00:00
08f4b557fc
Also fixed implementation of this atomic in GLSL. It was emitting code that would not compile because, as for HLSL, we must pass in the variable directly to atomic funcs, not via an in/out arg to a function. Bug: tint:1185 Change-Id: Id0e9f99d6368717511ef3a94473634c512e10cb8 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/91881 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Ben Clayton <bclayton@google.com> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
34 lines
1.1 KiB
Plaintext
34 lines
1.1 KiB
Plaintext
#include <metal_stdlib>
|
|
|
|
using namespace metal;
|
|
|
|
struct atomic_compare_exchange_resultu32 {
|
|
uint old_value;
|
|
bool exchanged;
|
|
};
|
|
template <typename A, typename T>
|
|
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
|
|
T old_value = compare;
|
|
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
|
|
return {old_value, exchanged};
|
|
}
|
|
|
|
void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) {
|
|
atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u);
|
|
}
|
|
|
|
void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {
|
|
{
|
|
atomic_store_explicit(tint_symbol_1, uint(), memory_order_relaxed);
|
|
}
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
|
atomicCompareExchangeWeak_83580d(tint_symbol_1);
|
|
}
|
|
|
|
kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
|
threadgroup atomic_uint tint_symbol_2;
|
|
compute_main_inner(local_invocation_index, &(tint_symbol_2));
|
|
return;
|
|
}
|
|
|