| #include <metal_stdlib> |
| |
| using namespace metal; |
| struct tint_array_wrapper { |
| /* 0x0000 */ int arr[8]; |
| }; |
| struct In2 { |
| /* 0x0000 */ tint_array_wrapper data_in2; |
| }; |
| struct tint_array_wrapper_1 { |
| /* 0x0000 */ int arr[1024]; |
| }; |
| struct Out0 { |
| /* 0x0000 */ tint_array_wrapper_1 data_out0; |
| }; |
| struct tint_array_wrapper_2 { |
| /* 0x0000 */ int arr[512]; |
| }; |
| struct In1 { |
| /* 0x0000 */ tint_array_wrapper_2 data_in1; |
| }; |
| struct In0 { |
| /* 0x0000 */ tint_array_wrapper_2 data_in0; |
| }; |
| |
| void main_1(const device In2& x_13, const device In1& x_17, device Out0& x_15, const device In0& x_19, thread uint3* const tint_symbol_1) { |
| uint base_index_in = 0u; |
| uint base_index_out = 0u; |
| int index_in0 = 0; |
| int index_in1 = 0; |
| int index_out0 = 0; |
| int index_out1 = 0; |
| int condition_index = 0; |
| int i = 0; |
| int temp0 = 0; |
| int temp1 = 0; |
| uint const x_58 = (*(tint_symbol_1)).x; |
| base_index_in = (128u * x_58); |
| uint const x_61 = (*(tint_symbol_1)).x; |
| base_index_out = (256u * x_61); |
| index_in0 = 127; |
| index_in1 = 383; |
| index_out0 = 255; |
| index_out1 = 383; |
| condition_index = 0; |
| i = 0; |
| while (true) { |
| int const x_67 = i; |
| if ((x_67 < 256)) { |
| } else { |
| break; |
| } |
| int const x_70 = condition_index; |
| int const x_72 = x_13.data_in2.arr[x_70]; |
| if ((x_72 == 0)) { |
| uint const x_77 = base_index_out; |
| int const x_78 = index_out0; |
| uint const x_81 = base_index_in; |
| int const x_82 = index_in0; |
| int const x_86 = x_17.data_in1.arr[(x_81 + as_type<uint>(x_82))]; |
| x_15.data_out0.arr[(x_77 + as_type<uint>(x_78))] = x_86; |
| int const x_88 = index_out0; |
| index_out0 = as_type<int>((as_type<uint>(x_88) - as_type<uint>(1))); |
| int const x_90 = index_in1; |
| index_in1 = as_type<int>((as_type<uint>(x_90) - as_type<uint>(1))); |
| } else { |
| uint const x_92 = base_index_out; |
| int const x_93 = index_out1; |
| uint const x_96 = base_index_in; |
| int const x_97 = index_in1; |
| int const x_101 = x_19.data_in0.arr[(x_96 + as_type<uint>(x_97))]; |
| x_15.data_out0.arr[(x_92 + as_type<uint>(x_93))] = x_101; |
| int const x_103 = index_out1; |
| index_out1 = as_type<int>((as_type<uint>(x_103) - as_type<uint>(1))); |
| int const x_105 = index_in1; |
| index_in1 = as_type<int>((as_type<uint>(x_105) - as_type<uint>(1))); |
| } |
| int const x_107 = condition_index; |
| int const x_110 = x_13.data_in2.arr[as_type<int>((as_type<uint>(x_107) + as_type<uint>(1)))]; |
| int const x_111 = condition_index; |
| condition_index = as_type<int>((as_type<uint>(x_111) + as_type<uint>(x_110))); |
| int const x_113 = index_in0; |
| temp0 = x_113; |
| int const x_114 = index_in1; |
| index_in0 = x_114; |
| int const x_115 = temp0; |
| index_in1 = x_115; |
| int const x_116 = index_out0; |
| temp1 = x_116; |
| int const x_117 = index_out1; |
| index_out0 = x_117; |
| int const x_118 = temp1; |
| index_out1 = x_118; |
| { |
| int const x_119 = i; |
| i = as_type<int>((as_type<uint>(x_119) + as_type<uint>(1))); |
| } |
| } |
| return; |
| } |
| |
| void tint_symbol_inner(const device In2& x_13, const device In1& x_17, device Out0& x_15, const device In0& x_19, uint3 gl_WorkGroupID_param, thread uint3* const tint_symbol_2) { |
| *(tint_symbol_2) = gl_WorkGroupID_param; |
| main_1(x_13, x_17, x_15, x_19, tint_symbol_2); |
| } |
| |
| kernel void tint_symbol(uint3 gl_WorkGroupID_param [[threadgroup_position_in_grid]], const device In2& x_13 [[buffer(1)]], const device In1& x_17 [[buffer(2)]], device Out0& x_15 [[buffer(0)]], const device In0& x_19 [[buffer(3)]]) { |
| thread uint3 tint_symbol_3 = 0u; |
| tint_symbol_inner(x_13, x_17, x_15, x_19, gl_WorkGroupID_param, &(tint_symbol_3)); |
| return; |
| } |
| |