I wrote a conv2d schedule for opencl device, after tunning,i got two config. The first config worked well, but the second config caused "*** NDRANGE_KERNEL executed abnormally *** ".
I dumped the OpenCL source code:
The first one:
__kernel void fuse_conv2d_clip_39_kernel0(__global float* restrict input0, __global float* restrict input1, __global float* restrict tensor, __global float* restrict input2) {
float compute[24];
__local float pad_temp_shared[224];
__local float input1_shared[48];
float pad_temp_shared_local[1];
float input1_shared_local[24];
for (int ff_init = 0; ff_init < 24; ++ff_init) {
compute[ff_init] = 0.000000e+00f;
}
for (int rc_outer = 0; rc_outer < 12; ++rc_outer) {
barrier(CLK_LOCAL_MEM_FENCE);
for (int ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner = 0; ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner < 2; ++ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) {
pad_temp_shared[(((((int)get_local_id(1)) * 16) + ((((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 8) * 8)) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) % 8))] = input0[((((((((int)get_group_id(1)) * 784) + (((int)get_group_id(0)) * 8)) + (rc_outer * 6272)) + ((((((int)get_local_id(1)) * 2) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 8)) / 14) * 3136)) + ((((((int)get_local_id(1)) * 2) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 8)) % 14) * 56)) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) % 8))];
}
if ((((int)get_local_id(1)) * 2) < (24 - (((int)get_local_id(0)) / 2))) {
if ((((int)get_group_id(2)) * 24) < ((144 - (((int)get_local_id(1)) * 2)) - (((int)get_local_id(0)) / 2))) {
input1_shared[(((((int)get_local_id(1)) * 4) + ((((int)get_local_id(0)) / 2) * 2)) + (((int)get_local_id(0)) % 2))] = input1[(((((((int)get_group_id(2)) * 576) + (rc_outer * 2)) + (((int)get_local_id(1)) * 48)) + ((((int)get_local_id(0)) / 2) * 24)) + (((int)get_local_id(0)) % 2))];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int rc_inner = 0; rc_inner < 2; ++rc_inner) {
pad_temp_shared_local[0] = pad_temp_shared[(((((int)get_local_id(1)) * 8) + ((int)get_local_id(0))) + (rc_inner * 112))];
for (int ax0 = 0; ax0 < 24; ++ax0) {
input1_shared_local[ax0] = input1_shared[(rc_inner + (ax0 * 2))];
}
for (int ff = 0; ff < 24; ++ff) {
compute[ff] = (compute[ff] + (pad_temp_shared_local[0] * input1_shared_local[ff]));
}
}
}
for (int ax1_inner_inner = 0; ax1_inner_inner < 24; ++ax1_inner_inner) {
tensor[((((((((int)get_group_id(2)) * 75264) + (((int)get_group_id(1)) * 784)) + (((int)get_group_id(0)) * 8)) + (((int)get_local_id(1)) * 56)) + ((int)get_local_id(0))) + (ax1_inner_inner * 3136))] = max(min((compute[ax1_inner_inner] + input2[((((int)get_group_id(2)) * 24) + ax1_inner_inner)]), 6.000000e+00f), 0.000000e+00f);
}
}
The second one:
__kernel void fuse_conv2d_clip_39_kernel0(__global float* restrict input0, __global float* restrict input1, __global float* restrict tensor, __global float* restrict input2) {
float compute[24];
__local float pad_temp_shared[64];
__local float input1_shared[48];
float pad_temp_shared_local[1];
float input1_shared_local[24];
for (int ff_init = 0; ff_init < 24; ++ff_init) {
compute[ff_init] = 0.000000e+00f;
}
for (int rc_outer = 0; rc_outer < 12; ++rc_outer) {
barrier(CLK_LOCAL_MEM_FENCE);
for (int ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner = 0; ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner < 2; ++ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) {
pad_temp_shared[((((((((int)get_local_id(1)) * 2) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 8)) / 4) * 32) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) % 8)) + ((((((int)get_local_id(1)) * 2) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 8)) % 4) * 8))] = input0[((((((((int)get_group_id(1)) * 224) + (((int)get_group_id(0)) * 8)) + (rc_outer * 6272)) + ((((((int)get_local_id(1)) * 2) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 8)) / 4) * 3136)) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) % 8)) + ((((((int)get_local_id(1)) * 2) + (((((int)get_local_id(0)) * 2) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 8)) % 4) * 56))];
}
for (int ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1 = 0; ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1 < 2; ++ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1) {
if ((((int)get_local_id(1)) * 6) < (24 - ((int)get_local_id(0)))) {
if (((((int)get_group_id(2)) * 24) + (((int)get_local_id(1)) * 6)) < (144 - ((int)get_local_id(0)))) {
input1_shared[(((((int)get_local_id(1)) * 12) + (((int)get_local_id(0)) * 2)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1)] = input1[(((((((int)get_group_id(2)) * 576) + (rc_outer * 2)) + (((int)get_local_id(1)) * 144)) + (((int)get_local_id(0)) * 24)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1)];
}
}
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int rc_inner = 0; rc_inner < 2; ++rc_inner) {
pad_temp_shared_local[0] = pad_temp_shared[(((((int)get_local_id(1)) * 8) + ((int)get_local_id(0))) + (rc_inner * 32))];
for (int ax0 = 0; ax0 < 24; ++ax0) {
input1_shared_local[ax0] = input1_shared[(rc_inner + (ax0 * 2))];
}
for (int ff = 0; ff < 24; ++ff) {
compute[ff] = (compute[ff] + (pad_temp_shared_local[0] * input1_shared_local[ff]));
}
}
}
for (int ax1_inner_inner = 0; ax1_inner_inner < 24; ++ax1_inner_inner) {
tensor[((((((((int)get_group_id(2)) * 75264) + (((int)get_group_id(1)) * 224)) + (((int)get_group_id(0)) * 8)) + (((int)get_local_id(1)) * 56)) + ((int)get_local_id(0))) + (ax1_inner_inner * 3136))] = max(min((compute[ax1_inner_inner] + input2[((((int)get_group_id(2)) * 24) + ax1_inner_inner)]), 6.000000e+00f), 0.000000e+00f);
}
}
Seems the error causes by the NDRANGE exceed the limit, but i am not sure, could someone explain it, thanks very much!