blob: 3ae8c42e3df15c0c7dbc4d24640bd24017e2a8b8 [file] [log] [blame]
James Price128980f2023-01-06 02:25:06 +00001#include <metal_stdlib>
2
3using namespace metal;
Ben Clayton9be037c2024-01-19 20:58:56 +00004
Ben Clayton02262d82024-01-23 18:09:39 +00005#define TINT_ISOLATE_UB(VOLATILE_NAME) \
David Neto012551a2024-10-15 16:00:17 +00006 {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
Ben Clayton9be037c2024-01-19 20:58:56 +00007
James Price128980f2023-01-06 02:25:06 +00008int 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
15void foo(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
16 {
17 int i = 0;
David Neto012551a2024-10-15 16:00:17 +000018 while(true) {
19 TINT_ISOLATE_UB(tint_volatile_false);
James Price128980f2023-01-06 02:25:06 +000020 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