Permalink
Find file Copy path
Fetching contributors…
Cannot retrieve contributors at this time
131 lines (126 sloc) 43.4 KB
// kernel for packing data
__kernel void default_function__kernel0(__global void* restrict data_vec, __global float* restrict data) {
for (int vh = 0; vh < 3; ++vh) {
((__global float*)data_vec)[(((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) * 256) + ((int)get_group_id(0))) * 3) + vh) * 6)] = ((((((1 - vh) <= ((int)get_group_id(2))) && (((int)get_group_id(2)) < (57 - vh))) && (1 <= ((int)get_group_id(1)))) && (((int)get_group_id(1)) < 15)) ? data[((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) + (((int)get_group_id(0)) * 784)) + (vh * 14)) * 4) + -57)] : 0.000000e+00f);
((__global float*)data_vec)[((((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) * 256) + ((int)get_group_id(0))) * 3) + vh) * 6) + 1)] = ((((((1 - vh) <= ((int)get_group_id(2))) && (((int)get_group_id(2)) < (57 - vh))) && (0 <= ((int)get_group_id(1)))) && (((int)get_group_id(1)) < 14)) ? data[((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) + (((int)get_group_id(0)) * 784)) + (vh * 14)) * 4) + -56)] : 0.000000e+00f);
((__global float*)data_vec)[((((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) * 256) + ((int)get_group_id(0))) * 3) + vh) * 6) + 2)] = ((((((1 - vh) <= ((int)get_group_id(2))) && (((int)get_group_id(2)) < (57 - vh))) && (0 <= ((int)get_group_id(1)))) && (((int)get_group_id(1)) < 14)) ? data[((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) + (((int)get_group_id(0)) * 784)) + (vh * 14)) * 4) + -55)] : 0.000000e+00f);
((__global float*)data_vec)[((((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) * 256) + ((int)get_group_id(0))) * 3) + vh) * 6) + 3)] = ((((((1 - vh) <= ((int)get_group_id(2))) && (((int)get_group_id(2)) < (57 - vh))) && (0 <= ((int)get_group_id(1)))) && (((int)get_group_id(1)) < 14)) ? data[((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) + (((int)get_group_id(0)) * 784)) + (vh * 14)) * 4) + -54)] : 0.000000e+00f);
((__global float*)data_vec)[((((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) * 256) + ((int)get_group_id(0))) * 3) + vh) * 6) + 4)] = ((((((1 - vh) <= ((int)get_group_id(2))) && (((int)get_group_id(2)) < (57 - vh))) && (0 <= ((int)get_group_id(1)))) && (((int)get_group_id(1)) < 14)) ? data[((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) + (((int)get_group_id(0)) * 784)) + (vh * 14)) * 4) + -53)] : 0.000000e+00f);
((__global float*)data_vec)[((((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) * 256) + ((int)get_group_id(0))) * 3) + vh) * 6) + 5)] = ((((((1 - vh) <= ((int)get_group_id(2))) && (((int)get_group_id(2)) < (57 - vh))) && (-1 <= ((int)get_group_id(1)))) && (((int)get_group_id(1)) < 13)) ? data[((((((((int)get_group_id(2)) * 14) + ((int)get_group_id(1))) + (((int)get_group_id(0)) * 784)) + (vh * 14)) * 4) + -52)] : 0.000000e+00f);
}
}
// kernel for packing filter
__kernel void default_function__kernel1(__global void* restrict kernel_vec, __global float* restrict weight) {
float4 _1;
int4 _2 = (int4)(((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9))+(2304*0), ((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9))+(2304*1), ((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9))+(2304*2), ((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9))+(2304*3));
_1.s0 = weight[_2.s0];
_1.s1 = weight[_2.s1];
_1.s2 = weight[_2.s2];
_1.s3 = weight[_2.s3];
vstore4(_1, 0, (__global float*)kernel_vec + (((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36));
float4 _3;
int4 _4 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 1))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 1))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 1))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 1))+(2304*3));
_3.s0 = weight[_4.s0];
_3.s1 = weight[_4.s1];
_3.s2 = weight[_4.s2];
_3.s3 = weight[_4.s3];
vstore4(_3, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 4));
float4 _5;
int4 _6 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 2))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 2))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 2))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 2))+(2304*3));
_5.s0 = weight[_6.s0];
_5.s1 = weight[_6.s1];
_5.s2 = weight[_6.s2];
_5.s3 = weight[_6.s3];
vstore4(_5, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 8));
float4 _7;
int4 _8 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 3))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 3))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 3))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 3))+(2304*3));
_7.s0 = weight[_8.s0];
_7.s1 = weight[_8.s1];
_7.s2 = weight[_8.s2];
_7.s3 = weight[_8.s3];
vstore4(_7, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 12));
float4 _9;
int4 _10 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 4))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 4))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 4))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 4))+(2304*3));
_9.s0 = weight[_10.s0];
_9.s1 = weight[_10.s1];
_9.s2 = weight[_10.s2];
_9.s3 = weight[_10.s3];
vstore4(_9, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 16));
float4 _11;
int4 _12 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 5))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 5))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 5))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 5))+(2304*3));
_11.s0 = weight[_12.s0];
_11.s1 = weight[_12.s1];
_11.s2 = weight[_12.s2];
_11.s3 = weight[_12.s3];
vstore4(_11, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 20));
float4 _13;
int4 _14 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 6))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 6))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 6))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 6))+(2304*3));
_13.s0 = weight[_14.s0];
_13.s1 = weight[_14.s1];
_13.s2 = weight[_14.s2];
_13.s3 = weight[_14.s3];
vstore4(_13, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 24));
float4 _15;
int4 _16 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 7))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 7))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 7))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 7))+(2304*3));
_15.s0 = weight[_16.s0];
_15.s1 = weight[_16.s1];
_15.s2 = weight[_16.s2];
_15.s3 = weight[_16.s3];
vstore4(_15, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 28));
float4 _17;
int4 _18 = (int4)((((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 8))+(2304*0), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 8))+(2304*1), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 8))+(2304*2), (((((((int)get_group_id(1)) * 1024) + ((int)get_group_id(0))) * 9) + 8))+(2304*3));
_17.s0 = weight[_18.s0];
_17.s1 = weight[_18.s1];
_17.s2 = weight[_18.s2];
_17.s3 = weight[_18.s3];
vstore4(_17, 0, (__global float*)kernel_vec + ((((((int)get_group_id(1)) * 256) + ((int)get_group_id(0))) * 36) + 32));
}
// kernel for convolution
__kernel void default_function__kernel2(__global void* restrict conv, __global void* restrict data_vec, __global void* restrict kernel_vec) {
vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4(((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f, 0.000000e+00f)), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
for (int ci = 0; ci < 256; ++ci) {
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[(((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18)], ((__global float*)data_vec)[(((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18)], ((__global float*)data_vec)[(((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18)], ((__global float*)data_vec)[(((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18)])) * vload4(0, (__global float*)kernel_vec + (((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)])) * vload4(0, (__global float*)kernel_vec + (((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)])) * vload4(0, (__global float*)kernel_vec + (((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)])) * vload4(0, (__global float*)kernel_vec + (((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 1)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 4)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 4)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 4)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 4)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 2)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 8)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 3)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 8)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 4)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 8)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 5)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 5)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 5)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 5)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 8)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 6)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 6)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 6)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 6)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 12)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 12)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 12)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 12)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 7)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 16)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 16)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 16)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 16)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 8)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 20)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 9)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 20)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 10)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 20)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 11)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 11)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 11)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 11)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 20)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 12)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 12)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 12)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 12)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 24)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 24)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 24)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 24)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 13)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 28)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 28)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 28)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 28)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
vstore4((vload4(0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 14)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 32)))), 0, (__global float*)conv + (((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 15)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 32)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 4));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 16)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 32)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 8));
vstore4((vload4(0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12)) + (((float4)(((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 17)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 17)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 17)], ((__global float*)data_vec)[((((((((int)get_group_id(1)) * 14) + ((int)get_group_id(0))) * 256) + ci) * 18) + 17)])) * vload4(0, (__global float*)kernel_vec + ((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 256) + ci) * 36) + 32)))), 0, (__global float*)conv + ((((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 14) + ((int)get_group_id(0))) * 16) + 12));
}
}
// kernel for unpacking the output
__kernel void default_function__kernel3(__global float* restrict output_unpack, __global void* restrict conv) {
output_unpack[((((((((int)get_group_id(2)) * 8) + ((int)get_local_id(2))) * 56) + ((int)get_group_id(1))) * 56) + ((int)get_group_id(0)))] = ((__global float*)conv)[((((((((int)get_group_id(2)) * 2) + (((int)get_local_id(2)) / 4)) * 12544) + (((int)get_local_id(2)) % 4)) + (((int)get_group_id(1)) * 224)) + (((int)get_group_id(0)) * 4))];
}