From a1ed41fe2427dfa2b5d0139444ceb77ad16a5a73 Mon Sep 17 00:00:00 2001 From: Jaroslaw Rzepecki Date: Fri, 13 Oct 2017 11:13:58 +0100 Subject: IVGCVSW-601: support for asymetric padding in cl conv and depthwise conv Change-Id: I5c6c95091ae77dba96459c0640f9f6167a988c8c Reviewed-on: http://mpd-gerrit.cambridge.arm.com/91700 Reviewed-by: Anthony Barbier Tested-by: Kaizen --- src/core/CL/cl_kernels/convolution_layer.cl | 14 +++++++------- src/core/CL/cl_kernels/depthwise_convolution.cl | 12 ++++++------ 2 files changed, 13 insertions(+), 13 deletions(-) (limited to 'src/core/CL/cl_kernels') diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl index 9e9d0b0ccc..e3018461e3 100644 --- a/src/core/CL/cl_kernels/convolution_layer.cl +++ b/src/core/CL/cl_kernels/convolution_layer.cl @@ -97,7 +97,7 @@ __kernel void reshape_to_columns( } } -#if defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) +#if defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) /** This kernel performs a reshaping of the input tensor to a tensor used to perform convolution using GEMM. * * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float @@ -134,8 +134,8 @@ __kernel void im2col_generic( const int batch = get_global_id(2) / filter_depth; // the batch // Calculate input indeces - const int xi = xc * STRIDE_X - PAD_X; - const int yi = yc * STRIDE_Y - PAD_Y; + const int xi = xc * STRIDE_X - PAD_LEFT; + const int yi = yc * STRIDE_Y - PAD_TOP; // Calculate output indeces const int xo = ch * KERNEL_WIDTH * KERNEL_HEIGHT; @@ -149,9 +149,9 @@ __kernel void im2col_generic( { for(int x = xi, x_e = xi + KERNEL_WIDTH; x < x_e; ++x, ++output_ptr) { -#if PAD_X == 0 && PAD_Y == 0 +#if PAD_LEFT == 0 && PAD_TOP == 0 && PAD_RIGHT == 0 && PAD_BOTTOM == 0 *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y)); -#else // PAD_X == 0 && PAD_Y == 0 +#else // PAD_LEFT == 0 && PAD_TOP == 0 && PAD_RIGHT == 0 && PAD_BOTTOM == 0 if(x < 0 || x >= SRC_WIDTH || y < 0 || y >= SRC_HEIGHT) { *output_ptr = 0; @@ -160,7 +160,7 @@ __kernel void im2col_generic( { *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y)); } -#endif // PAD_X == 0 && PAD_Y == 0 +#endif // PAD_LEFT == 0 && PAD_TOP == 0 && PAD_RIGHT == 0 && PAD_BOTTOM == 0 } } @@ -245,7 +245,7 @@ __kernel void im2col_kernel3x3_padx0_pady0( } #endif // HAS_BIAS } -#endif //defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) +#endif //defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) #if defined(WIDTH_OUTPUT) /** This kernel performs a reshaping of the output of the convolution layer. diff --git a/src/core/CL/cl_kernels/depthwise_convolution.cl b/src/core/CL/cl_kernels/depthwise_convolution.cl index 9c2c3a5b37..081a4e6c44 100644 --- a/src/core/CL/cl_kernels/depthwise_convolution.cl +++ b/src/core/CL/cl_kernels/depthwise_convolution.cl @@ -226,11 +226,11 @@ __kernel void depthwise_weights_reshape(TENSOR3D_DECLARATION(src), IMAGE_DECLARA } #endif //defined(SRC_WIDTH) && defined(DATA_TYPE) -#if defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(DATA_TYPE) +#if defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(DATA_TYPE) /** This kernel performs a reshaping of the input tensor to a tensor used to perform depthwise convolution using vector to matrix multiplication. * * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float - * @note The convolution information must be passed at compile time using -DSTRIDE_X, -DSTRIDE_Y, -DPAD_X, -DPAD_Y, -DKERNEL_WIDHT, -DKERNEL_HEIGHT, -DSRC_WIDTH, -DSRC_HEIGHT + * @note The convolution information must be passed at compile time using -DSTRIDE_X, -DSTRIDE_Y, -DPAD_LEFT, -DPAD_TOP, -DPAD_RIGHT, -DPAD_BOTTOM, -DKERNEL_WIDHT, -DKERNEL_HEIGHT, -DSRC_WIDTH, -DSRC_HEIGHT * * @param[in] src_ptr Pointer to the source tensor. Supported data types: QS8/QS16/F16/F32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) @@ -255,11 +255,11 @@ __kernel void depthwise_im2col(TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(d Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); const int src_pixel_linear = get_global_id(1) * STRIDE_X; - const int full_length = SRC_WIDTH + 2 * PAD_X; + const int full_length = SRC_WIDTH + PAD_LEFT + PAD_RIGHT; const int max_initial_x = STRIDE_X * (((full_length - KERNEL_WIDTH) / STRIDE_X) + 1); - const int src_x = -PAD_X + src_pixel_linear % max_initial_x; - const int src_y = -PAD_Y + src_pixel_linear / max_initial_x * STRIDE_Y; + const int src_x = -PAD_LEFT + src_pixel_linear % max_initial_x; + const int src_y = -PAD_TOP + src_pixel_linear / max_initial_x * STRIDE_Y; const int src_z = get_global_id(2); __global uchar *input_ptr = src_ptr + src_offset_first_element_in_bytes + src_z * src_stride_z; @@ -281,7 +281,7 @@ __kernel void depthwise_im2col(TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(d } } -#endif //defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(SRC_WIDTH) && defined(DATA_TYPE) +#endif //defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(SRC_WIDTH) && defined(DATA_TYPE) #if defined(CONV_WIDTH) && defined(CONV_HEIGHT) && defined(DATA_TYPE) -- cgit v1.2.1