James Price | 128980f | 2023-01-06 02:25:06 +0000 | [diff] [blame] | 1 | #include <metal_stdlib> |
| 2 | |
| 3 | using namespace metal; |
Ben Clayton | 9be037c | 2024-01-19 20:58:56 +0000 | [diff] [blame] | 4 | |
Ben Clayton | 02262d8 | 2024-01-23 18:09:39 +0000 | [diff] [blame] | 5 | #define TINT_ISOLATE_UB(VOLATILE_NAME) \ |
David Neto | 012551a | 2024-10-15 16:00:17 +0000 | [diff] [blame] | 6 | {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;} |
Ben Clayton | 9be037c | 2024-01-19 20:58:56 +0000 | [diff] [blame] | 7 | |
James Price | 128980f | 2023-01-06 02:25:06 +0000 | [diff] [blame] | 8 | int tint_workgroupUniformLoad(threadgroup int* const p) { |
| 9 | threadgroup_barrier(mem_flags::mem_threadgroup); |
| 10 | int const result = *(p); |
| 11 | threadgroup_barrier(mem_flags::mem_threadgroup); |
| 12 | return result; |
| 13 | } |
| 14 | |
| 15 | void foo(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) { |
| 16 | { |
| 17 | int i = 0; |
David Neto | 012551a | 2024-10-15 16:00:17 +0000 | [diff] [blame] | 18 | while(true) { |
| 19 | TINT_ISOLATE_UB(tint_volatile_false); |
James Price | 128980f | 2023-01-06 02:25:06 +0000 | [diff] [blame] | 20 | int const tint_symbol = i; |
| 21 | int const tint_symbol_1 = tint_workgroupUniformLoad(tint_symbol_4); |
| 22 | if (!((tint_symbol < tint_symbol_1))) { |
| 23 | break; |
| 24 | } |
| 25 | { |
| 26 | } |
| 27 | { |
| 28 | int const tint_symbol_2 = i; |
| 29 | int const tint_symbol_3 = tint_workgroupUniformLoad(tint_symbol_5); |
| 30 | i = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(tint_symbol_3))); |
| 31 | } |
| 32 | } |
| 33 | } |
| 34 | } |
| 35 | |