blob: fb4dbcd2a9be2d14454fed9e1582041d2fd1800c [file] [log] [blame]
David Neto1e19b552021-06-17 09:10:04 +00001#include <metal_stdlib>
2
3using namespace metal;
Ben Claytonf47887d2022-06-24 17:01:59 +00004
5template<typename T, size_t N>
6struct 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 Neto1e19b552021-06-17 09:10:04 +000015};
Ben Clayton8ec32a62022-02-09 23:55:51 +000016
Ben Clayton02262d82024-01-23 18:09:39 +000017#define TINT_ISOLATE_UB(VOLATILE_NAME) \
18 volatile bool VOLATILE_NAME = true; \
19 if (VOLATILE_NAME)
Ben Clayton9be037c2024-01-19 20:58:56 +000020
James Price0e22bdb2023-03-20 21:46:01 +000021struct 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 Clayton9dc42ad2024-02-05 18:01:33 +000027void 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 Neto1e19b552021-06-17 09:10:04 +000035struct S {
Ben Claytonf47887d2022-06-24 17:01:59 +000036 /* 0x0000 */ tint_array<int4, 4> arr;
David Neto1e19b552021-06-17 09:10:04 +000037};
Ben Clayton8ec32a62022-02-09 23:55:51 +000038
Ben Claytonf47887d2022-06-24 17:01:59 +000039tint_array<int4, 4> ret_arr() {
dan sinclaira01fd242023-12-05 20:23:42 +000040 tint_array<int4, 4> const tint_symbol_2 = tint_array<int4, 4>{};
Antonio Maiorano93baaae2022-03-15 15:35:13 +000041 return tint_symbol_2;
42}
43
dan sinclaira01fd242023-12-05 20:23:42 +000044S ret_struct_arr() {
45 S const tint_symbol_3 = S{};
46 return tint_symbol_3;
47}
48
Ben Clayton9dc42ad2024-02-05 18:01:33 +000049void 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 Claytonf47887d2022-06-24 17:01:59 +000050 tint_array<int4, 4> src_function = {};
dan sinclaira01fd242023-12-05 20:23:42 +000051 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 Price0e22bdb2023-03-20 21:46:01 +000053 (*(tint_private_vars)).dst = src_param;
54 (*(tint_private_vars)).dst = ret_arr();
Ben Claytonf47887d2022-06-24 17:01:59 +000055 tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
James Price0e22bdb2023-03-20 21:46:01 +000056 (*(tint_private_vars)).dst = src_let;
57 (*(tint_private_vars)).dst = src_function;
58 (*(tint_private_vars)).dst = (*(tint_private_vars)).src_private;
Ben Clayton9dc42ad2024-02-05 18:01:33 +000059 (*(tint_private_vars)).dst = *(tint_symbol_6);
dan sinclaira01fd242023-12-05 20:23:42 +000060 S const tint_symbol_1 = ret_struct_arr();
61 (*(tint_private_vars)).dst = tint_symbol_1.arr;
dan sinclaira01fd242023-12-05 20:23:42 +000062 (*(tint_private_vars)).dst = (*(tint_symbol_7)).arr;
Ben Clayton9dc42ad2024-02-05 18:01:33 +000063 (*(tint_private_vars)).dst = (*(tint_symbol_8)).arr;
Ben Claytonf47887d2022-06-24 17:01:59 +000064 tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
James Price0e22bdb2023-03-20 21:46:01 +000065 (*(tint_private_vars)).dst_nested = src_nested;
David Neto1e19b552021-06-17 09:10:04 +000066}
67
Ben Clayton9dc42ad2024-02-05 18:01:33 +000068void 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 sinclaira01fd242023-12-05 20:23:42 +000070 tint_array<int4, 4> const a = tint_array<int4, 4>{};
Ben Clayton9dc42ad2024-02-05 18:01:33 +000071 foo(a, tint_private_vars, tint_symbol_9, tint_symbol_10, tint_symbol_11);
dan sinclaira01fd242023-12-05 20:23:42 +000072}
73
Ben Clayton9dc42ad2024-02-05 18:01:33 +000074kernel 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 sinclaira01fd242023-12-05 20:23:42 +000075 thread tint_private_vars_struct tint_private_vars = {};
Ben Clayton9dc42ad2024-02-05 18:01:33 +000076 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 sinclaira01fd242023-12-05 20:23:42 +000078 return;
79}
80