| #include <metal_stdlib> | |
| using namespace metal; | |
| struct S { | |
| /* 0x0000 */ array<int4, 4> arr; | |
| }; | |
| struct tint_module_vars_struct { | |
| thread array<int4, 4>* src_private; | |
| threadgroup array<int4, 4>* src_workgroup; | |
| const constant S* src_uniform; | |
| device S* src_storage; | |
| }; | |
| struct tint_symbol_1 { | |
| array<int4, 4> tint_symbol; | |
| }; | |
| array<int4, 4> ret_arr() { | |
| return array<int4, 4>{}; | |
| } | |
| S ret_struct_arr() { | |
| return S{}; | |
| } | |
| void foo(array<int4, 4> src_param, tint_module_vars_struct tint_module_vars) { | |
| array<int4, 4> src_function = {}; | |
| array<int4, 4> dst = {}; | |
| dst = array<int4, 4>{int4(1), int4(2), int4(3), int4(3)}; | |
| dst = src_param; | |
| dst = ret_arr(); | |
| array<int4, 4> const src_let = array<int4, 4>{}; | |
| dst = src_let; | |
| dst = src_function; | |
| dst = (*tint_module_vars.src_private); | |
| dst = (*tint_module_vars.src_workgroup); | |
| dst = ret_struct_arr().arr; | |
| dst = (*tint_module_vars.src_uniform).arr; | |
| dst = (*tint_module_vars.src_storage).arr; | |
| array<array<array<int, 2>, 3>, 4> dst_nested = {}; | |
| array<array<array<int, 2>, 3>, 4> src_nested = {}; | |
| dst_nested = src_nested; | |
| } | |
| void main_inner(uint tint_local_index, tint_module_vars_struct tint_module_vars) { | |
| { | |
| uint idx = tint_local_index; | |
| while(true) { | |
| uint const v = idx; | |
| if ((v >= 4u)) { | |
| break; | |
| } | |
| (*tint_module_vars.src_workgroup)[v] = int4(0); | |
| { | |
| idx = (idx + 1u); | |
| } | |
| } | |
| } | |
| (threadgroup_barrier(mem_flags::mem_threadgroup)); | |
| array<int4, 4> const val = array<int4, 4>{}; | |
| (foo(val, tint_module_vars)); | |
| } | |
| [[max_total_threads_per_threadgroup(1)]] | |
| kernel void v_1(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v_2 [[threadgroup(0)]], const constant S* src_uniform [[buffer(0)]], device S* src_storage [[buffer(1)]]) { | |
| thread array<int4, 4> src_private = {}; | |
| tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.src_private=(&src_private), .src_workgroup=(&(*v_2).tint_symbol), .src_uniform=src_uniform, .src_storage=src_storage}; | |
| (main_inner(tint_local_index, tint_module_vars)); | |
| } |