Ben Clayton | f2ec7f3 | 2021-06-29 11:53:15 +0000 | [diff] [blame] | 1 | #include <metal_stdlib> |
Ben Clayton | 6a77236 | 2021-06-18 18:56:13 +0000 | [diff] [blame] | 2 | |
Ben Clayton | f2ec7f3 | 2021-06-29 11:53:15 +0000 | [diff] [blame] | 3 | using namespace metal; |
Ben Clayton | 77a90cb | 2023-03-06 18:25:08 +0000 | [diff] [blame] | 4 | void atomicExchange_e114ba(threadgroup atomic_int* const tint_symbol, device int* const tint_symbol_1) { |
James Price | a5d73ce | 2021-08-04 22:15:28 +0000 | [diff] [blame] | 5 | int res = atomic_exchange_explicit(tint_symbol, 1, memory_order_relaxed); |
Ben Clayton | 77a90cb | 2023-03-06 18:25:08 +0000 | [diff] [blame] | 6 | *(tint_symbol_1) = res; |
James Price | a5d73ce | 2021-08-04 22:15:28 +0000 | [diff] [blame] | 7 | } |
| 8 | |
Ben Clayton | 77a90cb | 2023-03-06 18:25:08 +0000 | [diff] [blame] | 9 | void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_2, device int* const tint_symbol_3) { |
James Price | a5d73ce | 2021-08-04 22:15:28 +0000 | [diff] [blame] | 10 | { |
Ben Clayton | 77a90cb | 2023-03-06 18:25:08 +0000 | [diff] [blame] | 11 | atomic_store_explicit(tint_symbol_2, 0, memory_order_relaxed); |
James Price | a5d73ce | 2021-08-04 22:15:28 +0000 | [diff] [blame] | 12 | } |
| 13 | threadgroup_barrier(mem_flags::mem_threadgroup); |
Ben Clayton | 77a90cb | 2023-03-06 18:25:08 +0000 | [diff] [blame] | 14 | atomicExchange_e114ba(tint_symbol_2, tint_symbol_3); |
Ben Clayton | e6d171a | 2021-06-18 18:56:13 +0000 | [diff] [blame] | 15 | } |
| 16 | |
Ben Clayton | 77a90cb | 2023-03-06 18:25:08 +0000 | [diff] [blame] | 17 | kernel void compute_main(device int* tint_symbol_5 [[buffer(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) { |
| 18 | threadgroup atomic_int tint_symbol_4; |
| 19 | compute_main_inner(local_invocation_index, &(tint_symbol_4), tint_symbol_5); |
Ben Clayton | f2ec7f3 | 2021-06-29 11:53:15 +0000 | [diff] [blame] | 20 | return; |
Ben Clayton | e6d171a | 2021-06-18 18:56:13 +0000 | [diff] [blame] | 21 | } |
| 22 | |