James Price | de73431 | 2024-05-16 22:20:44 +0000 | [diff] [blame] | 1 | #include <metal_stdlib> |
| 2 | using namespace metal; |
James Price | 48be7e8 | 2024-06-19 22:38:07 +0000 | [diff] [blame] | 3 | |
James Price | de73431 | 2024-05-16 22:20:44 +0000 | [diff] [blame] | 4 | template<typename T, size_t N> |
| 5 | struct tint_array { |
| 6 | const constant T& operator[](size_t i) const constant { return elements[i]; } |
| 7 | device T& operator[](size_t i) device { return elements[i]; } |
| 8 | const device T& operator[](size_t i) const device { return elements[i]; } |
| 9 | thread T& operator[](size_t i) thread { return elements[i]; } |
| 10 | const thread T& operator[](size_t i) const thread { return elements[i]; } |
| 11 | threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } |
| 12 | const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } |
| 13 | T elements[N]; |
| 14 | }; |
dan sinclair | f1f381a | 2023-11-22 09:44:15 +0000 | [diff] [blame] | 15 | |
James Price | de73431 | 2024-05-16 22:20:44 +0000 | [diff] [blame] | 16 | struct S { |
| 17 | tint_array<int, 4> arr; |
| 18 | }; |
James Price | 48be7e8 | 2024-06-19 22:38:07 +0000 | [diff] [blame] | 19 | |
James Price | de73431 | 2024-05-16 22:20:44 +0000 | [diff] [blame] | 20 | struct tint_module_vars_struct { |
| 21 | device int* s; |
| 22 | }; |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 23 | |
James Price | de73431 | 2024-05-16 22:20:44 +0000 | [diff] [blame] | 24 | int foo() { |
| 25 | tint_array<int, 4> const src = tint_array<int, 4>{}; |
| 26 | tint_array<int, 4> dst = {}; |
| 27 | S dst_struct = {}; |
| 28 | tint_array<tint_array<int, 4>, 2> dst_array = {}; |
| 29 | thread tint_array<int, 4>* const dst_ptr = (&dst); |
| 30 | thread S* const dst_struct_ptr = (&dst_struct); |
| 31 | thread tint_array<tint_array<int, 4>, 2>* const dst_array_ptr = (&dst_array); |
| 32 | dst_struct.arr = src; |
| 33 | dst_array[1] = src; |
| 34 | (*dst_ptr) = src; |
| 35 | (*dst_struct_ptr).arr = src; |
| 36 | (*dst_array_ptr)[0] = src; |
James Price | 2d24558 | 2024-09-04 16:04:34 +0000 | [diff] [blame] | 37 | return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>((*dst_ptr)[0]) + as_type<uint>((*dst_struct_ptr).arr[0])))) + as_type<uint>((*dst_array_ptr)[0][0]))); |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 38 | } |
James Price | 48be7e8 | 2024-06-19 22:38:07 +0000 | [diff] [blame] | 39 | |
James Price | de73431 | 2024-05-16 22:20:44 +0000 | [diff] [blame] | 40 | kernel void tint_symbol(device int* s [[buffer(0)]]) { |
James Price | 47fddc2 | 2024-05-17 19:29:38 +0000 | [diff] [blame] | 41 | tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s}; |
James Price | de73431 | 2024-05-16 22:20:44 +0000 | [diff] [blame] | 42 | (*tint_module_vars.s) = foo(); |
dan sinclair | a01fd24 | 2023-12-05 20:23:42 +0000 | [diff] [blame] | 43 | } |