From c1d3bcdc00cf13a5dc4ede795933ab7e20da4158 Mon Sep 17 00:00:00 2001 From: yejianwu Date: Wed, 4 Apr 2018 15:58:10 +0800 Subject: [PATCH] move macro to inline function for check_out_of_range_for_image2d --- mace/kernels/opencl/cl/activation.cl | 3 +-- mace/kernels/opencl/cl/addn.cl | 2 -- mace/kernels/opencl/cl/batch_norm.cl | 3 +-- mace/kernels/opencl/cl/bias_add.cl | 3 +-- mace/kernels/opencl/cl/buffer_to_image.cl | 16 +--------------- mace/kernels/opencl/cl/channel_shuffle.cl | 4 +--- mace/kernels/opencl/cl/common.h | 2 ++ mace/kernels/opencl/cl/concat.cl | 6 ++---- mace/kernels/opencl/cl/conv_2d.cl | 8 -------- mace/kernels/opencl/cl/conv_2d_1x1.cl | 8 -------- mace/kernels/opencl/cl/conv_2d_3x3.cl | 10 ---------- mace/kernels/opencl/cl/cwise.cl | 3 +-- mace/kernels/opencl/cl/depth_to_space.cl | 6 ++---- mace/kernels/opencl/cl/depthwise_conv2d.cl | 16 ---------------- mace/kernels/opencl/cl/eltwise.cl | 3 +-- mace/kernels/opencl/cl/fully_connected.cl | 6 ++---- mace/kernels/opencl/cl/matmul.cl | 8 -------- mace/kernels/opencl/cl/pooling.cl | 2 -- mace/kernels/opencl/cl/resize_bilinear.cl | 3 +-- mace/kernels/opencl/cl/slice.cl | 5 +---- mace/kernels/opencl/cl/softmax.cl | 3 +-- mace/kernels/opencl/cl/space_to_batch.cl | 6 ++---- mace/kernels/opencl/cl/winograd_transform.cl | 12 ++---------- 23 files changed, 22 insertions(+), 116 deletions(-) diff --git a/mace/kernels/opencl/cl/activation.cl b/mace/kernels/opencl/cl/activation.cl index a3aa0be7..4ddb8fd0 100644 --- a/mace/kernels/opencl/cl/activation.cl +++ b/mace/kernels/opencl/cl/activation.cl @@ -31,8 +31,7 @@ __kernel void activation(KERNEL_ERROR_PARAMS DATA_TYPE4 out = do_activation(in, relux_max_limit); #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, hb, kernel_error); -#endif + WRITE_IMAGET(output, (int2)(pos, hb), out); } diff --git a/mace/kernels/opencl/cl/addn.cl b/mace/kernels/opencl/cl/addn.cl index 7d2d4323..444b57e5 100644 --- a/mace/kernels/opencl/cl/addn.cl +++ b/mace/kernels/opencl/cl/addn.cl @@ -32,9 +32,7 @@ __kernel void addn(KERNEL_ERROR_PARAMS out = out + in3; #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, hb, kernel_error); -#endif 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 f3ba9acc..7e46889e 100644 --- a/mace/kernels/opencl/cl/batch_norm.cl +++ b/mace/kernels/opencl/cl/batch_norm.cl @@ -49,8 +49,7 @@ __kernel void batch_norm(KERNEL_ERROR_PARAMS out = do_activation(out, relux_max_limit); #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, hb, kernel_error); -#endif + 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 77cf7e7e..b86738d6 100644 --- a/mace/kernels/opencl/cl/bias_add.cl +++ b/mace/kernels/opencl/cl/bias_add.cl @@ -24,8 +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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, hb, kernel_error); -#endif + 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 4a4b068e..e311c467 100644 --- a/mace/kernels/opencl/cl/buffer_to_image.cl +++ b/mace/kernels/opencl/cl/buffer_to_image.cl @@ -50,9 +50,7 @@ __kernel void filter_buffer_to_image(KERNEL_ERROR_PARAMS } int2 coord = (int2)(w, h); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, h, kernel_error); -#endif WRITE_IMAGET(output, coord, values); } @@ -160,9 +158,7 @@ __kernel void dw_filter_buffer_to_image(KERNEL_ERROR_PARAMS } int2 coord = (int2)(w, h); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, h, kernel_error); -#endif WRITE_IMAGET(output, coord, values); } @@ -205,9 +201,7 @@ __kernel void in_out_buffer_to_image(KERNEL_ERROR_PARAMS values = vload4(0, input + offset); } int2 coord = (int2)(w, h); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, h, kernel_error); -#endif WRITE_IMAGET(output, coord, values); } @@ -284,9 +278,7 @@ __kernel void arg_buffer_to_image(KERNEL_ERROR_PARAMS values = vload4(0, input + offset); } int2 coord = (int2)(w, h); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, h, kernel_error); -#endif WRITE_IMAGET(output, coord, values); } @@ -364,9 +356,7 @@ __kernel void in_out_height_buffer_to_image(KERNEL_ERROR_PARAMS values.x = *(input + offset); } int2 coord = (int2)(w, h); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, h, kernel_error); -#endif WRITE_IMAGET(output, coord, values); } @@ -448,9 +438,7 @@ __kernel void in_out_width_buffer_to_image(KERNEL_ERROR_PARAMS values.x = *(input + offset); } int2 coord = (int2)(w, h); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, h, kernel_error); -#endif WRITE_IMAGET(output, coord, values); } @@ -525,9 +513,7 @@ __kernel void winograd_filter_buffer_to_image(KERNEL_ERROR_PARAMS int2 coord = (int2)(w, h); -#ifdef OUT_OF_RANGE_CHECK - check_out_of_range_for_image2d(output, coord.x, coord.y + out_channels * 15, kernel_error); -#endif + 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 a73ab8c5..5a1fb5cd 100644 --- a/mace/kernels/opencl/cl/channel_shuffle.cl +++ b/mace/kernels/opencl/cl/channel_shuffle.cl @@ -51,9 +51,7 @@ __kernel void channel_shuffle(KERNEL_ERROR_PARAMS int out_x = mad24(mad24(group_chan_blk_idx, groups, g_blk), width, width_idx); -#ifdef OUT_OF_RANGE_CHECK - check_out_of_range_for_image2d(output, out_x + groups_blks_width * 3, hb_idx, kernel_error); -#endif + 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 b3054c94..b038bb7d 100644 --- a/mace/kernels/opencl/cl/common.h +++ b/mace/kernels/opencl/cl/common.h @@ -78,10 +78,12 @@ inline void check_out_of_range_for_image2d(__write_only image2d_t image, __private const int x, __private const int y, global char *kernel_error) { +#ifdef OUT_OF_RANGE_CHECK int2 image_dim = get_image_dim(image); if (x >= image_dim.x || y >= image_dim.y) { *kernel_error = '1'; } +#endif } #endif // MACE_KERNELS_OPENCL_CL_COMMON_H_ diff --git a/mace/kernels/opencl/cl/concat.cl b/mace/kernels/opencl/cl/concat.cl index b1d5fb52..31b7eb47 100644 --- a/mace/kernels/opencl/cl/concat.cl +++ b/mace/kernels/opencl/cl/concat.cl @@ -82,9 +82,8 @@ __kernel void concat_channel(KERNEL_ERROR_PARAMS const int pos = mad24(chan_blk_idx, width, width_idx); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); -#endif + WRITE_IMAGET(output, (int2)(pos, hb_idx), data); } @@ -115,9 +114,8 @@ __kernel void concat_channel_multi(KERNEL_ERROR_PARAMS const int pos = mad24(chan_blk_idx + chan_blk_offset, width, width_idx); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); -#endif + 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 03a2b47e..2bf65bc4 100644 --- a/mace/kernels/opencl/cl/conv_2d.cl +++ b/mace/kernels/opencl/cl/conv_2d.cl @@ -129,30 +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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif 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 0a748925..4ee8e0a5 100644 --- a/mace/kernels/opencl/cl/conv_2d_1x1.cl +++ b/mace/kernels/opencl/cl/conv_2d_1x1.cl @@ -105,29 +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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); -#endif 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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); -#endif 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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); -#endif 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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + out_x_idx, out_hb, kernel_error); -#endif 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 a6d07b8b..6fbd0d2e 100644 --- a/mace/kernels/opencl/cl/conv_2d_3x3.cl +++ b/mace/kernels/opencl/cl/conv_2d_3x3.cl @@ -136,45 +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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out3); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif 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 f5d96647..ebaca716 100644 --- a/mace/kernels/opencl/cl/cwise.cl +++ b/mace/kernels/opencl/cl/cwise.cl @@ -44,8 +44,7 @@ __kernel void cwise(KERNEL_ERROR_PARAMS out.w = fabs(in0.w); #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, hb, kernel_error); -#endif + 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 f6e82b13..2cbb6c7a 100644 --- a/mace/kernels/opencl/cl/depth_to_space.cl +++ b/mace/kernels/opencl/cl/depth_to_space.cl @@ -37,9 +37,8 @@ __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)); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_pos, out_hb, kernel_error); -#endif + WRITE_IMAGET(output, (int2)(out_pos, out_hb), in_data); } @@ -80,8 +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)); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_pos, out_hb, kernel_error); -#endif + 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 932e37da..7015bb71 100644 --- a/mace/kernels/opencl/cl/depthwise_conv2d.cl +++ b/mace/kernels/opencl/cl/depthwise_conv2d.cl @@ -123,30 +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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += out_w_blks; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out3); } @@ -261,29 +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; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out0); w += 1; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out1); w += 1; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif WRITE_IMAGET(output, (int2)(out_x_base + w, out_hb), out2); w += 1; if (w >= out_width) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_x_base + w, out_hb, kernel_error); -#endif 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 81ba7d87..a171bf07 100644 --- a/mace/kernels/opencl/cl/eltwise.cl +++ b/mace/kernels/opencl/cl/eltwise.cl @@ -37,8 +37,7 @@ __kernel void eltwise(KERNEL_ERROR_PARAMS out = in0 - in1; #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, w, hb, kernel_error); -#endif + 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 b913986a..086d5e4f 100644 --- a/mace/kernels/opencl/cl/fully_connected.cl +++ b/mace/kernels/opencl/cl/fully_connected.cl @@ -60,9 +60,8 @@ __kernel void fully_connected(KERNEL_ERROR_PARAMS result = do_activation(result, relux_max_limit); #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_blk_idx, batch_idx, kernel_error); -#endif + WRITE_IMAGET(output, (int2)(out_blk_idx, batch_idx), result); } @@ -154,9 +153,8 @@ __kernel void fully_connected_width(KERNEL_ERROR_PARAMS result = do_activation(result, relux_max_limit); #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_blk_idx, batch_idx, kernel_error); -#endif + 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 1fa2b46f..8641b4f5 100644 --- a/mace/kernels/opencl/cl/matmul.cl +++ b/mace/kernels/opencl/cl/matmul.cl @@ -48,26 +48,18 @@ __kernel void matmul(KERNEL_ERROR_PARAMS c3 += (DATA_TYPE4)(dot(a0, b3), dot(a1, b3), dot(a2, b3), dot(a3, b3)); } -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(C, gx, gy, kernel_error); -#endif WRITE_IMAGET(C, (int2)(gx, gy), c0); if ((gx + 1) >= N) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(C, gx + 1, gy, kernel_error); -#endif WRITE_IMAGET(C, (int2)(gx + 1, gy), c1); if ((gx + 2) >= N) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(C, gx + 2, gy, kernel_error); -#endif WRITE_IMAGET(C, (int2)(gx + 2, gy), c2); if ((gx + 3) >= N) return; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(C, gx + 3, gy, kernel_error); -#endif 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 c3abeb01..3ab5f62d 100644 --- a/mace/kernels/opencl/cl/pooling.cl +++ b/mace/kernels/opencl/cl/pooling.cl @@ -96,8 +96,6 @@ __kernel void pooling(KERNEL_ERROR_PARAMS #endif const int pos = mad24(out_chan_idx, out_width, out_width_idx); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, out_hb_idx, kernel_error); -#endif 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 091fb617..fb069c42 100644 --- a/mace/kernels/opencl/cl/resize_bilinear.cl +++ b/mace/kernels/opencl/cl/resize_bilinear.cl @@ -58,9 +58,8 @@ __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); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_w_offset + w, out_h_offset + h, kernel_error); -#endif + 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 0116fd91..c78e6dc8 100644 --- a/mace/kernels/opencl/cl/slice.cl +++ b/mace/kernels/opencl/cl/slice.cl @@ -24,9 +24,6 @@ __kernel void slice(KERNEL_ERROR_PARAMS width, width_idx), hb_idx)); const int pos = mad24(chan_blk_idx, width, width_idx); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); -#endif - WRITE_IMAGET(output, - (int2)(pos, hb_idx), data); + 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 ae434eb9..ad554b54 100644 --- a/mace/kernels/opencl/cl/softmax.cl +++ b/mace/kernels/opencl/cl/softmax.cl @@ -85,8 +85,7 @@ __kernel void softmax(KERNEL_ERROR_PARAMS data = native_exp(data) / sum; } -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, pos, hb_idx, kernel_error); -#endif + 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 dcf92778..46764a08 100644 --- a/mace/kernels/opencl/cl/space_to_batch.cl +++ b/mace/kernels/opencl/cl/space_to_batch.cl @@ -46,9 +46,8 @@ __kernel void space_to_batch(KERNEL_ERROR_PARAMS int2 batch_coord = (int2)(mul24(chan_idx, batch_width) + batch_w_idx, batch_hb_idx); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(batch_data, batch_coord.x, batch_coord.y, kernel_error); -#endif + WRITE_IMAGET(batch_data, batch_coord, value); } @@ -94,9 +93,8 @@ __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); -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(space_data, space_coord.x, space_coord.y, kernel_error); -#endif + 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 aac11e97..23d4bc6f 100644 --- a/mace/kernels/opencl/cl/winograd_transform.cl +++ b/mace/kernels/opencl/cl/winograd_transform.cl @@ -94,9 +94,8 @@ __kernel void winograd_transform_2x2(KERNEL_ERROR_PARAMS input3[2] = tv3[2] - tv3[1]; input3[3] = tv3[1] - tv3[3]; -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, out_width_idx, chan_blk_idx + chan_blk_idx * 15, kernel_error); -#endif + #pragma unroll for (short i = 0; i < 4; ++i) { WRITE_IMAGET(output, (int2)(out_width_idx, chan_blk_idx), input0[i]); @@ -213,30 +212,23 @@ __kernel void winograd_inverse_transform_2x2(KERNEL_ERROR_PARAMS in1[1] = do_activation(in1[1], relux_max_limit); #endif -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, coord_x, coord_y, kernel_error); -#endif + WRITE_IMAGET(output, (int2)(coord_x, coord_y), in0[0]); t = 0; if (out_width_idx + 1 < out_width) { -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, coord_x + 1, coord_y, kernel_error); -#endif WRITE_IMAGET(output, (int2)(coord_x + 1, coord_y), in0[1]); t += 1; } if (out_height_idx + 1 < out_height) { -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, coord_x, coord_y + 1, kernel_error); -#endif WRITE_IMAGET(output, (int2)(coord_x, coord_y + 1), in1[0]); t += 1; } if (t == 2) { -#ifdef OUT_OF_RANGE_CHECK check_out_of_range_for_image2d(output, coord_x + 1, coord_y + 1, kernel_error); -#endif WRITE_IMAGET(output, (int2)(coord_x + 1, coord_y + 1), in1[1]); } -- GitLab