David Neto | 1e19b55 | 2021-06-17 09:10:04 +0000 | [diff] [blame] | 1 | #include <metal_stdlib> |
| 2 | |
| 3 | using namespace metal; |
Ben Clayton | f47887d | 2022-06-24 17:01:59 +0000 | [diff] [blame] | 4 | |
| 5 | template<typename T, size_t N> |
| 6 | struct tint_array { |
| 7 | const constant T& operator[](size_t i) const constant { return elements[i]; } |
| 8 | device T& operator[](size_t i) device { return elements[i]; } |
| 9 | const device T& operator[](size_t i) const device { return elements[i]; } |
| 10 | thread T& operator[](size_t i) thread { return elements[i]; } |
| 11 | const thread T& operator[](size_t i) const thread { return elements[i]; } |
| 12 | threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } |
| 13 | const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } |
| 14 | T elements[N]; |
David Neto | 1e19b55 | 2021-06-17 09:10:04 +0000 | [diff] [blame] | 15 | }; |
Ben Clayton | 8ec32a6 | 2022-02-09 23:55:51 +0000 | [diff] [blame] | 16 | |
Ben Clayton | 02262d8 | 2024-01-23 18:09:39 +0000 | [diff] [blame] | 17 | #define TINT_ISOLATE_UB(VOLATILE_NAME) \ |
| 18 | volatile bool VOLATILE_NAME = true; \ |
| 19 | if (VOLATILE_NAME) |
Ben Clayton | 9be037c | 2024-01-19 20:58:56 +0000 | [diff] [blame] | 20 | |
James Price | 0e22bdb | 2023-03-20 21:46:01 +0000 | [diff] [blame] | 21 | struct tint_private_vars_struct { |
| 22 | tint_array<int4, 4> src_private; |
| 23 | tint_array<int4, 4> dst; |
| 24 | tint_array<tint_array<tint_array<int, 2>, 3>, 4> dst_nested; |
| 25 | }; |
| 26 | |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 27 | void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<int4, 4>* const tint_symbol_5) { |
| 28 | TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) { |
| 29 | uint const i = idx; |
| 30 | (*(tint_symbol_5))[i] = int4(0); |
| 31 | } |
| 32 | threadgroup_barrier(mem_flags::mem_threadgroup); |
| 33 | } |
| 34 | |
David Neto | 1e19b55 | 2021-06-17 09:10:04 +0000 | [diff] [blame] | 35 | struct S { |
Ben Clayton | f47887d | 2022-06-24 17:01:59 +0000 | [diff] [blame] | 36 | /* 0x0000 */ tint_array<int4, 4> arr; |
David Neto | 1e19b55 | 2021-06-17 09:10:04 +0000 | [diff] [blame] | 37 | }; |
Ben Clayton | 8ec32a6 | 2022-02-09 23:55:51 +0000 | [diff] [blame] | 38 | |
Ben Clayton | f47887d | 2022-06-24 17:01:59 +0000 | [diff] [blame] | 39 | tint_array<int4, 4> ret_arr() { |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 40 | tint_array<int4, 4> const tint_symbol_2 = tint_array<int4, 4>{}; |
Antonio Maiorano | 93baaae | 2022-03-15 15:35:13 +0000 | [diff] [blame] | 41 | return tint_symbol_2; |
| 42 | } |
| 43 | |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 44 | S ret_struct_arr() { |
| 45 | S const tint_symbol_3 = S{}; |
| 46 | return tint_symbol_3; |
| 47 | } |
| 48 | |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 49 | void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8) { |
Ben Clayton | f47887d | 2022-06-24 17:01:59 +0000 | [diff] [blame] | 50 | tint_array<int4, 4> src_function = {}; |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 51 | tint_array<int4, 4> const tint_symbol_4 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)}; |
| 52 | (*(tint_private_vars)).dst = tint_symbol_4; |
James Price | 0e22bdb | 2023-03-20 21:46:01 +0000 | [diff] [blame] | 53 | (*(tint_private_vars)).dst = src_param; |
| 54 | (*(tint_private_vars)).dst = ret_arr(); |
Ben Clayton | f47887d | 2022-06-24 17:01:59 +0000 | [diff] [blame] | 55 | tint_array<int4, 4> const src_let = tint_array<int4, 4>{}; |
James Price | 0e22bdb | 2023-03-20 21:46:01 +0000 | [diff] [blame] | 56 | (*(tint_private_vars)).dst = src_let; |
| 57 | (*(tint_private_vars)).dst = src_function; |
| 58 | (*(tint_private_vars)).dst = (*(tint_private_vars)).src_private; |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 59 | (*(tint_private_vars)).dst = *(tint_symbol_6); |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 60 | S const tint_symbol_1 = ret_struct_arr(); |
| 61 | (*(tint_private_vars)).dst = tint_symbol_1.arr; |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 62 | (*(tint_private_vars)).dst = (*(tint_symbol_7)).arr; |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 63 | (*(tint_private_vars)).dst = (*(tint_symbol_8)).arr; |
Ben Clayton | f47887d | 2022-06-24 17:01:59 +0000 | [diff] [blame] | 64 | tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {}; |
James Price | 0e22bdb | 2023-03-20 21:46:01 +0000 | [diff] [blame] | 65 | (*(tint_private_vars)).dst_nested = src_nested; |
David Neto | 1e19b55 | 2021-06-17 09:10:04 +0000 | [diff] [blame] | 66 | } |
| 67 | |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 68 | void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_9, const constant S* const tint_symbol_10, device S* const tint_symbol_11) { |
| 69 | tint_zero_workgroup_memory(local_invocation_index, tint_symbol_9); |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 70 | tint_array<int4, 4> const a = tint_array<int4, 4>{}; |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 71 | foo(a, tint_private_vars, tint_symbol_9, tint_symbol_10, tint_symbol_11); |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 72 | } |
| 73 | |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 74 | kernel void tint_symbol(const constant S* tint_symbol_13 [[buffer(0)]], device S* tint_symbol_14 [[buffer(1)]], uint local_invocation_index [[thread_index_in_threadgroup]]) { |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 75 | thread tint_private_vars_struct tint_private_vars = {}; |
Ben Clayton | 9dc42ad | 2024-02-05 18:01:33 +0000 | [diff] [blame] | 76 | threadgroup tint_array<int4, 4> tint_symbol_12; |
| 77 | tint_symbol_inner(local_invocation_index, &(tint_private_vars), &(tint_symbol_12), tint_symbol_13, tint_symbol_14); |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 78 | return; |
| 79 | } |
| 80 | |