blob: 393cfd11160cf3be9548bda6eeb937645a1db0cf [file] [log] [blame]
#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(2)]], const device In1& x_17 [[buffer(1)]], device Out0& x_15 [[buffer(3)]], const device In0& x_19 [[buffer(0)]]) {
thread uint3 tint_symbol_3 = 0u;
tint_symbol_inner(x_13, x_17, x_15, x_19, gl_WorkGroupID_param, &(tint_symbol_3));
return;
}