diff --git a/mace/kernels/opencl/cl/activation.cl b/mace/kernels/opencl/cl/activation.cl index 4ddb8fd03b6a897c42dd1941145a000f1c54c1b4..d5d7dc4a12f634b50d22b2a7f11e4adb92cd05a4 100644 --- a/mace/kernels/opencl/cl/activation.cl +++ b/mace/kernels/opencl/cl/activation.cl @@ -31,7 +31,7 @@ __kernel void activation(KERNEL_ERROR_PARAMS DATA_TYPE4 out = do_activation(in, relux_max_limit); #endif - check_out_of_range_for_image2d(output, pos, hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, hb, kernel_error); WRITE_IMAGET(output, (int2)(pos, hb), out); } diff --git a/mace/kernels/opencl/cl/addn.cl b/mace/kernels/opencl/cl/addn.cl index 444b57e591a6189b305ef24d7c6eb8372fcc3a39..52c8ab6d2b6ef1415da7fe27d3a54014158bb931 100644 --- a/mace/kernels/opencl/cl/addn.cl +++ b/mace/kernels/opencl/cl/addn.cl @@ -32,7 +32,7 @@ __kernel void addn(KERNEL_ERROR_PARAMS out = out + in3; #endif - check_out_of_range_for_image2d(output, w, hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, hb, kernel_error); WRITE_IMAGET(output, (int2)(w, hb), out); } diff --git a/mace/kernels/opencl/cl/batch_norm.cl b/mace/kernels/opencl/cl/batch_norm.cl index 7e46889e57f484e8b9f167bf8a4b5d16d32beddc..8303b3ffd3f58a936a3b5287fb6f3b646d8032c2 100644 --- a/mace/kernels/opencl/cl/batch_norm.cl +++ b/mace/kernels/opencl/cl/batch_norm.cl @@ -49,7 +49,7 @@ __kernel void batch_norm(KERNEL_ERROR_PARAMS out = do_activation(out, relux_max_limit); #endif - check_out_of_range_for_image2d(output, pos, hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, hb, kernel_error); WRITE_IMAGET(output, (int2)(pos, hb), out); } diff --git a/mace/kernels/opencl/cl/bias_add.cl b/mace/kernels/opencl/cl/bias_add.cl index b86738d68e346feb7d06d8c365180c72c287abd6..a09de1fb1aad297fcaf723af8143ba7519e11891 100644 --- a/mace/kernels/opencl/cl/bias_add.cl +++ b/mace/kernels/opencl/cl/bias_add.cl @@ -24,7 +24,7 @@ __kernel void bias_add(KERNEL_ERROR_PARAMS DATA_TYPE4 bias_value = READ_IMAGET(bias, SAMPLER, (int2)(ch_blk, 0)); DATA_TYPE4 out = in + bias_value; - check_out_of_range_for_image2d(output, pos, hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, hb, kernel_error); WRITE_IMAGET(output, (int2)(pos, hb), out); } diff --git a/mace/kernels/opencl/cl/buffer_to_image.cl b/mace/kernels/opencl/cl/buffer_to_image.cl index e311c467ce10ed9ec65199f0ff3ad5201fdd4791..32c6fc6ce792ebdd03dcff23a07ac4b1f6aed72a 100644 --- a/mace/kernels/opencl/cl/buffer_to_image.cl +++ b/mace/kernels/opencl/cl/buffer_to_image.cl @@ -50,7 +50,7 @@ __kernel void filter_buffer_to_image(KERNEL_ERROR_PARAMS } int2 coord = (int2)(w, h); - check_out_of_range_for_image2d(output, w, h, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, h, kernel_error); WRITE_IMAGET(output, coord, values); } @@ -158,7 +158,7 @@ __kernel void dw_filter_buffer_to_image(KERNEL_ERROR_PARAMS } int2 coord = (int2)(w, h); - check_out_of_range_for_image2d(output, w, h, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, h, kernel_error); WRITE_IMAGET(output, coord, values); } @@ -201,7 +201,7 @@ __kernel void in_out_buffer_to_image(KERNEL_ERROR_PARAMS values = vload4(0, input + offset); } int2 coord = (int2)(w, h); - check_out_of_range_for_image2d(output, w, h, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, h, kernel_error); WRITE_IMAGET(output, coord, values); } @@ -278,7 +278,7 @@ __kernel void arg_buffer_to_image(KERNEL_ERROR_PARAMS values = vload4(0, input + offset); } int2 coord = (int2)(w, h); - check_out_of_range_for_image2d(output, w, h, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, h, kernel_error); WRITE_IMAGET(output, coord, values); } @@ -356,7 +356,7 @@ __kernel void in_out_height_buffer_to_image(KERNEL_ERROR_PARAMS values.x = *(input + offset); } int2 coord = (int2)(w, h); - check_out_of_range_for_image2d(output, w, h, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, h, kernel_error); WRITE_IMAGET(output, coord, values); } @@ -438,7 +438,7 @@ __kernel void in_out_width_buffer_to_image(KERNEL_ERROR_PARAMS values.x = *(input + offset); } int2 coord = (int2)(w, h); - check_out_of_range_for_image2d(output, w, h, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, h, kernel_error); WRITE_IMAGET(output, coord, values); } @@ -513,7 +513,7 @@ __kernel void winograd_filter_buffer_to_image(KERNEL_ERROR_PARAMS int2 coord = (int2)(w, h); - check_out_of_range_for_image2d(output, coord.x, coord.y + out_channels * 15, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, coord.x, coord.y + out_channels * 15, kernel_error); #pragma unroll for (short i = 0; i < 4; ++i) { diff --git a/mace/kernels/opencl/cl/channel_shuffle.cl b/mace/kernels/opencl/cl/channel_shuffle.cl index 5a1fb5cde38911918a7e68fa1ae2250990b7f010..1674cc1db49db146adf7a258bb570c3bda5fba30 100644 --- a/mace/kernels/opencl/cl/channel_shuffle.cl +++ b/mace/kernels/opencl/cl/channel_shuffle.cl @@ -51,7 +51,7 @@ __kernel void channel_shuffle(KERNEL_ERROR_PARAMS int out_x = mad24(mad24(group_chan_blk_idx, groups, g_blk), width, width_idx); - check_out_of_range_for_image2d(output, out_x + groups_blks_width * 3, hb_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x + groups_blks_width * 3, hb_idx, kernel_error); WRITE_IMAGET(output, (int2)(out_x, hb_idx), out_chan_data0); out_x += groups_blks_width; diff --git a/mace/kernels/opencl/cl/common.h b/mace/kernels/opencl/cl/common.h index b038bb7d17a12b5531a5444990d6d88e16e3f8e2..905987b4f56e98d2f5160db079e33b65fd611b1f 100644 --- a/mace/kernels/opencl/cl/common.h +++ b/mace/kernels/opencl/cl/common.h @@ -47,6 +47,13 @@ #endif +#ifdef OUT_OF_RANGE_CHECK +#define CHECK_OUT_OF_RANGE_FOR_IMAGE2D(image, x, y, kernel_error) \ + check_out_of_range_for_image2d(image, x, y, kernel_error) +#else +#define CHECK_OUT_OF_RANGE_FOR_IMAGE2D(image, x, y, kernel_error) +#endif + __constant sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; diff --git a/mace/kernels/opencl/cl/concat.cl b/mace/kernels/opencl/cl/concat.cl index 31b7eb47f06829e065815c445c4efa7e36143931..76ab1192477f9b46451ddb35dbc5cdd99d6dab29 100644 --- a/mace/kernels/opencl/cl/concat.cl +++ b/mace/kernels/opencl/cl/concat.cl @@ -82,7 +82,7 @@ __kernel void concat_channel(KERNEL_ERROR_PARAMS const int pos = mad24(chan_blk_idx, width, width_idx); - check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, hb_idx, kernel_error); WRITE_IMAGET(output, (int2)(pos, hb_idx), data); } @@ -114,7 +114,7 @@ __kernel void concat_channel_multi(KERNEL_ERROR_PARAMS const int pos = mad24(chan_blk_idx + chan_blk_offset, width, width_idx); - check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, hb_idx, kernel_error); WRITE_IMAGET(output, (int2)(pos, hb_idx), data); } diff --git a/mace/kernels/opencl/cl/conv_2d.cl b/mace/kernels/opencl/cl/conv_2d.cl index 2bf65bc4c0001a83f88cccaab635d290c202c2c2..5b27b1e9337cefd5093c0026c13172eabff58bf8 100644 --- a/mace/kernels/opencl/cl/conv_2d.cl +++ b/mace/kernels/opencl/cl/conv_2d.cl @@ -129,22 +129,22 @@ __kernel void conv_2d(KERNEL_ERROR_PARAMS const int out_x_base = mul24(out_ch_blk, out_width); int w = out_w_blk; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out3); } diff --git a/mace/kernels/opencl/cl/conv_2d_1x1.cl b/mace/kernels/opencl/cl/conv_2d_1x1.cl index 4ee8e0a5e08240c48836dac532bed017b141500d..48274184f3619fa83ae37bb828729b05a8fd2b20 100644 --- a/mace/kernels/opencl/cl/conv_2d_1x1.cl +++ b/mace/kernels/opencl/cl/conv_2d_1x1.cl @@ -105,21 +105,21 @@ __kernel void conv_2d_1x1(KERNEL_ERROR_PARAMS const int out_x_base = mul24(out_ch_blk, width); int out_x_idx = out_w_blk; - check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + out_x_idx, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + out_x_idx, out_hb), out0); out_x_idx += out_w_blks; if (out_x_idx >= width) return; - check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + out_x_idx, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + out_x_idx, out_hb), out1); out_x_idx += out_w_blks; if (out_x_idx >= width) return; - check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + out_x_idx, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + out_x_idx, out_hb), out2); out_x_idx += out_w_blks; if (out_x_idx >= width) return; - check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + out_x_idx, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + out_x_idx, out_hb), out3); } diff --git a/mace/kernels/opencl/cl/conv_2d_3x3.cl b/mace/kernels/opencl/cl/conv_2d_3x3.cl index 6fbd0d2ea9834dc407284c3e0c128d31ded7eace..f034bd57d626a0399bd35e5980334cf86f824342 100644 --- a/mace/kernels/opencl/cl/conv_2d_3x3.cl +++ b/mace/kernels/opencl/cl/conv_2d_3x3.cl @@ -136,35 +136,35 @@ __kernel void conv_2d_3x3(KERNEL_ERROR_PARAMS const int out_x_base = mul24(out_ch_blk, out_width); int w = out_w_blk; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out3); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out4); diff --git a/mace/kernels/opencl/cl/cwise.cl b/mace/kernels/opencl/cl/cwise.cl index ebaca71643f9a41a12f793d85e3cb4bfb5c4b359..f02dadf45fa3c63bf0d39348108a2c4775b1c6b1 100644 --- a/mace/kernels/opencl/cl/cwise.cl +++ b/mace/kernels/opencl/cl/cwise.cl @@ -44,7 +44,7 @@ __kernel void cwise(KERNEL_ERROR_PARAMS out.w = fabs(in0.w); #endif - check_out_of_range_for_image2d(output, w, hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, hb, kernel_error); WRITE_IMAGET(output, (int2)(w, hb), out); } diff --git a/mace/kernels/opencl/cl/depth_to_space.cl b/mace/kernels/opencl/cl/depth_to_space.cl index 2cbb6c7aaaa613edf128dc2f3e318d116ba84972..6b2cdcc64737582ddf52afe1e9c233f9e9e7c02e 100644 --- a/mace/kernels/opencl/cl/depth_to_space.cl +++ b/mace/kernels/opencl/cl/depth_to_space.cl @@ -37,7 +37,7 @@ __kernel void depth_to_space(KERNEL_ERROR_PARAMS const int in_pos = mad24(in_d, input_width, in_w); DATA_TYPE4 in_data = READ_IMAGET(input, SAMPLER, (int2)(in_pos, in_hb)); - check_out_of_range_for_image2d(output, out_pos, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_pos, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_pos, out_hb), in_data); } @@ -79,7 +79,7 @@ __kernel void space_to_depth(KERNEL_ERROR_PARAMS const int out_pos = mad24(out_d, output_width, out_w); DATA_TYPE4 in_data = READ_IMAGET(input, SAMPLER, (int2)(in_pos, hb)); - check_out_of_range_for_image2d(output, out_pos, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_pos, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_pos, out_hb), in_data); } diff --git a/mace/kernels/opencl/cl/depthwise_conv2d.cl b/mace/kernels/opencl/cl/depthwise_conv2d.cl index 7015bb711dbfb216fb4b4e7f80ce24ee68ab38b7..bfdc790098e328f2b62999d5d8e2c85fbe631e68 100644 --- a/mace/kernels/opencl/cl/depthwise_conv2d.cl +++ b/mace/kernels/opencl/cl/depthwise_conv2d.cl @@ -123,22 +123,22 @@ __kernel void depthwise_conv2d(KERNEL_ERROR_PARAMS const short out_x_base = mul24(out_ch_blk, out_width); short w = out_w_blk; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += out_w_blks; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out3); } @@ -253,21 +253,21 @@ __kernel void depthwise_conv2d_s1(KERNEL_ERROR_PARAMS const short out_x_base = mul24(out_ch_blk, out_width); short w = out_w_blk; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += 1; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += 1; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += 1; if (w >= out_width) return; - check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_x_base + w, out_hb, kernel_error); WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out3); } diff --git a/mace/kernels/opencl/cl/eltwise.cl b/mace/kernels/opencl/cl/eltwise.cl index a171bf071273409aa42c15774abdf97bba43ec23..923131e87947c9200b4d83389bbeddefc6d95a9b 100644 --- a/mace/kernels/opencl/cl/eltwise.cl +++ b/mace/kernels/opencl/cl/eltwise.cl @@ -37,7 +37,7 @@ __kernel void eltwise(KERNEL_ERROR_PARAMS out = in0 - in1; #endif - check_out_of_range_for_image2d(output, w, hb, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, w, hb, kernel_error); WRITE_IMAGET(output, (int2)(w, hb), out); } diff --git a/mace/kernels/opencl/cl/fully_connected.cl b/mace/kernels/opencl/cl/fully_connected.cl index 086d5e4ffea90d0874bca08ec983f1f0ef9af1a7..cbacd43949003725d98b6df6c597e74b1c331d14 100644 --- a/mace/kernels/opencl/cl/fully_connected.cl +++ b/mace/kernels/opencl/cl/fully_connected.cl @@ -60,7 +60,7 @@ __kernel void fully_connected(KERNEL_ERROR_PARAMS result = do_activation(result, relux_max_limit); #endif - check_out_of_range_for_image2d(output, out_blk_idx, batch_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_blk_idx, batch_idx, kernel_error); WRITE_IMAGET(output, (int2)(out_blk_idx, batch_idx), result); } @@ -153,7 +153,7 @@ __kernel void fully_connected_width(KERNEL_ERROR_PARAMS result = do_activation(result, relux_max_limit); #endif - check_out_of_range_for_image2d(output, out_blk_idx, batch_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_blk_idx, batch_idx, kernel_error); WRITE_IMAGET(output, (int2)(out_blk_idx, batch_idx), result); } diff --git a/mace/kernels/opencl/cl/matmul.cl b/mace/kernels/opencl/cl/matmul.cl index 8641b4f5a52693d44c490c464e2dfec4ac404716..099aac0a24f7b9200967c3d0164774d39121553c 100644 --- a/mace/kernels/opencl/cl/matmul.cl +++ b/mace/kernels/opencl/cl/matmul.cl @@ -48,18 +48,18 @@ __kernel void matmul(KERNEL_ERROR_PARAMS c3 += (DATA_TYPE4)(dot(a0, b3), dot(a1, b3), dot(a2, b3), dot(a3, b3)); } - check_out_of_range_for_image2d(C, gx, gy, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(C, gx, gy, kernel_error); WRITE_IMAGET(C, (int2)(gx, gy), c0); if ((gx + 1) >= N) return; - check_out_of_range_for_image2d(C, gx + 1, gy, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(C, gx + 1, gy, kernel_error); WRITE_IMAGET(C, (int2)(gx + 1, gy), c1); if ((gx + 2) >= N) return; - check_out_of_range_for_image2d(C, gx + 2, gy, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(C, gx + 2, gy, kernel_error); WRITE_IMAGET(C, (int2)(gx + 2, gy), c2); if ((gx + 3) >= N) return; - check_out_of_range_for_image2d(C, gx + 3, gy, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(C, gx + 3, gy, kernel_error); WRITE_IMAGET(C, (int2)(gx + 3, gy), c3); } diff --git a/mace/kernels/opencl/cl/pooling.cl b/mace/kernels/opencl/cl/pooling.cl index 3ab5f62de2eaa7f856835c2fd1abe4c885e69981..4d5beef642361a987b0d1e3b71bfe855980bc1d7 100644 --- a/mace/kernels/opencl/cl/pooling.cl +++ b/mace/kernels/opencl/cl/pooling.cl @@ -96,6 +96,6 @@ __kernel void pooling(KERNEL_ERROR_PARAMS #endif const int pos = mad24(out_chan_idx, out_width, out_width_idx); - check_out_of_range_for_image2d(output, pos, out_hb_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, out_hb_idx, kernel_error); WRITE_IMAGET(output, (int2)(pos, out_hb_idx), res); } diff --git a/mace/kernels/opencl/cl/resize_bilinear.cl b/mace/kernels/opencl/cl/resize_bilinear.cl index fb069c42fc5ce8215bcfd7a606df41a167fc4e65..4e30e630902f555410f208181a3976acf62fd793 100644 --- a/mace/kernels/opencl/cl/resize_bilinear.cl +++ b/mace/kernels/opencl/cl/resize_bilinear.cl @@ -58,7 +58,7 @@ __kernel void resize_bilinear_nocache(KERNEL_ERROR_PARAMS const int out_w_offset = mul24(ch_blk, out_width); const int out_h_offset = mul24(b, out_height); - check_out_of_range_for_image2d(output, out_w_offset + w, out_h_offset + h, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_w_offset + w, out_h_offset + h, kernel_error); WRITE_IMAGET(output, (int2)(out_w_offset + w, out_h_offset + h), out); } diff --git a/mace/kernels/opencl/cl/slice.cl b/mace/kernels/opencl/cl/slice.cl index c78e6dc826a8740f982bb1edb5a5fa15836a11ad..723f5ba1b50703d29f11585bf12e3eb692660282 100644 --- a/mace/kernels/opencl/cl/slice.cl +++ b/mace/kernels/opencl/cl/slice.cl @@ -24,6 +24,6 @@ __kernel void slice(KERNEL_ERROR_PARAMS width, width_idx), hb_idx)); const int pos = mad24(chan_blk_idx, width, width_idx); - check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, hb_idx, kernel_error); WRITE_IMAGET(output, (int2)(pos, hb_idx), data); } diff --git a/mace/kernels/opencl/cl/softmax.cl b/mace/kernels/opencl/cl/softmax.cl index ad554b542ca0e11b09ccf0a1e9475aeff0f27ef4..1bc1dc7e5d5967d99b15dc03cd721f09c40b9333 100644 --- a/mace/kernels/opencl/cl/softmax.cl +++ b/mace/kernels/opencl/cl/softmax.cl @@ -85,7 +85,7 @@ __kernel void softmax(KERNEL_ERROR_PARAMS data = native_exp(data) / sum; } - check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, pos, hb_idx, kernel_error); WRITE_IMAGET(output, (int2)(pos, hb_idx), data); } diff --git a/mace/kernels/opencl/cl/space_to_batch.cl b/mace/kernels/opencl/cl/space_to_batch.cl index 46764a08fc2f7d6fcd813eb5daf48f1527ea7ed0..a8031a04ad58ff147e5b9927f007914b055e0691 100644 --- a/mace/kernels/opencl/cl/space_to_batch.cl +++ b/mace/kernels/opencl/cl/space_to_batch.cl @@ -46,7 +46,7 @@ __kernel void space_to_batch(KERNEL_ERROR_PARAMS int2 batch_coord = (int2)(mul24(chan_idx, batch_width) + batch_w_idx, batch_hb_idx); - check_out_of_range_for_image2d(batch_data, batch_coord.x, batch_coord.y, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(batch_data, batch_coord.x, batch_coord.y, kernel_error); WRITE_IMAGET(batch_data, batch_coord, value); } @@ -93,7 +93,7 @@ __kernel void batch_to_space(KERNEL_ERROR_PARAMS int2 space_coord = (int2)(mul24(chan_idx, space_width) + space_w_idx, space_b_idx * space_height + space_h_idx); - check_out_of_range_for_image2d(space_data, space_coord.x, space_coord.y, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(space_data, space_coord.x, space_coord.y, kernel_error); WRITE_IMAGET(space_data, space_coord, value); } diff --git a/mace/kernels/opencl/cl/winograd_transform.cl b/mace/kernels/opencl/cl/winograd_transform.cl index 23d4bc6fd2de3e42150c2e607e541b37aee8b0bd..329454afe64f1476a097923a5498042b412d110d 100644 --- a/mace/kernels/opencl/cl/winograd_transform.cl +++ b/mace/kernels/opencl/cl/winograd_transform.cl @@ -94,7 +94,7 @@ __kernel void winograd_transform_2x2(KERNEL_ERROR_PARAMS input3[2] = tv3[2] - tv3[1]; input3[3] = tv3[1] - tv3[3]; - check_out_of_range_for_image2d(output, out_width_idx, chan_blk_idx + chan_blk_idx * 15, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, out_width_idx, chan_blk_idx + chan_blk_idx * 15, kernel_error); #pragma unroll for (short i = 0; i < 4; ++i) { @@ -212,23 +212,23 @@ __kernel void winograd_inverse_transform_2x2(KERNEL_ERROR_PARAMS in1[1] = do_activation(in1[1], relux_max_limit); #endif - check_out_of_range_for_image2d(output, coord_x, coord_y, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, coord_x, coord_y, kernel_error); WRITE_IMAGET(output, (int2)(coord_x, coord_y), in0[0]); t = 0; if (out_width_idx + 1 < out_width) { - check_out_of_range_for_image2d(output, coord_x + 1, coord_y, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, coord_x + 1, coord_y, kernel_error); WRITE_IMAGET(output, (int2)(coord_x + 1, coord_y), in0[1]); t += 1; } if (out_height_idx + 1 < out_height) { - check_out_of_range_for_image2d(output, coord_x, coord_y + 1, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, coord_x, coord_y + 1, kernel_error); WRITE_IMAGET(output, (int2)(coord_x, coord_y + 1), in1[0]); t += 1; } if (t == 2) { - check_out_of_range_for_image2d(output, coord_x + 1, coord_y + 1, kernel_error); + CHECK_OUT_OF_RANGE_FOR_IMAGE2D(output, coord_x + 1, coord_y + 1, kernel_error); WRITE_IMAGET(output, (int2)(coord_x + 1, coord_y + 1), in1[1]); } diff --git a/tools/tuning_run.sh b/tools/tuning_run.sh index 28741d58dd05c53b9c4f9f6a1ad2e9079e790ee2..c4e8dbe99e403b6d067b4a802022f341e6a49141 100644 --- a/tools/tuning_run.sh +++ b/tools/tuning_run.sh @@ -70,7 +70,6 @@ else ADB_CMD_STR="LD_LIBRARY_PATH=${PHONE_DATA_DIR} \ MACE_TUNING=${tuning_flag} \ - MACE_OUT_OF_RANGE_CHECK="1" \ MACE_CPP_MIN_VLOG_LEVEL=$VLOG_LEVEL \ MACE_RUN_PARAMETER_PATH=${PHONE_DATA_DIR}/mace_run.config \ MACE_KERNEL_PATH=$KERNEL_DIR \