According to the official Metal Feature Set Tables, the Apple9 family introduces full 64-bit atomics such as atomic_compare_exchange_weak_explicit() and others. However when trying to use it the compiler throws an error "no matching function call...". Example:
kernel void test_64bit_cas(device atomic_ulong* atomic_var [[buffer(0)]]) {
ulong expected = 0;
ulong desired = 1;
// This fails to compile despite Apple9 hardware supporting 64-bit buffer atomics
atomic_compare_exchange_weak_explicit(atomic_var, &expected, desired, memory_order_relaxed, memory_order_relaxed);
}
Am i doing something wrong ?