| Ben Clayton | 6992f51 | 2022-12-01 18:49:09 +0000 | [diff] [blame] | 1 | #include <metal_stdlib> |
| 2 | |
| 3 | using namespace metal; |
| Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 4 | void tint_zero_workgroup_memory(uint local_idx, threadgroup int* const tint_symbol_1) { |
| Antonio Maiorano | f5abb82 | 2024-02-28 01:01:02 +0000 | [diff] [blame] | 5 | if ((local_idx < 1u)) { |
| Ben Clayton | 6992f51 | 2022-12-01 18:49:09 +0000 | [diff] [blame] | 6 | *(tint_symbol_1) = 0; |
| 7 | } |
| 8 | threadgroup_barrier(mem_flags::mem_threadgroup); |
| Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 9 | } |
| 10 | |
| 11 | int func(threadgroup int* const pointer) { |
| 12 | return *(pointer); |
| 13 | } |
| 14 | |
| 15 | void tint_symbol_inner(uint local_invocation_index, threadgroup int* const tint_symbol_2) { |
| 16 | tint_zero_workgroup_memory(local_invocation_index, tint_symbol_2); |
| 17 | int const r = func(tint_symbol_2); |
| Ben Clayton | 6992f51 | 2022-12-01 18:49:09 +0000 | [diff] [blame] | 18 | } |
| 19 | |
| 20 | kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) { |
| Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 21 | threadgroup int tint_symbol_3; |
| 22 | tint_symbol_inner(local_invocation_index, &(tint_symbol_3)); |
| Ben Clayton | 6992f51 | 2022-12-01 18:49:09 +0000 | [diff] [blame] | 23 | return; |
| 24 | } |
| 25 | |
