1#include <metal_stdlib> 2 3using namespace metal; 4 5template <typename A, typename T> 6vec<T, 2> atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { 7 T prev_value = compare; 8 bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed); 9 return {prev_value, matched}; 10} 11 12void atomicCompareExchangeWeak_89ea3b(threadgroup atomic_int* const tint_symbol) { 13 int2 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1); 14} 15 16void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) { 17 { 18 atomic_store_explicit(tint_symbol_1, int(), memory_order_relaxed); 19 } 20 threadgroup_barrier(mem_flags::mem_threadgroup); 21 atomicCompareExchangeWeak_89ea3b(tint_symbol_1); 22} 23 24kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) { 25 threadgroup atomic_int tint_symbol_2; 26 compute_main_inner(local_invocation_index, &(tint_symbol_2)); 27 return; 28} 29 30