diff options
Diffstat (limited to 'src/core/CL/cl_kernels/depthwise_convolution.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/depthwise_convolution.cl | 151 |
1 files changed, 149 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/depthwise_convolution.cl b/src/core/CL/cl_kernels/depthwise_convolution.cl index fb4a0fc157..1b2f5cccaa 100644 --- a/src/core/CL/cl_kernels/depthwise_convolution.cl +++ b/src/core/CL/cl_kernels/depthwise_convolution.cl @@ -21,7 +21,6 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ - #include "helpers.h" #include "activation_float_helpers.h" @@ -1479,7 +1478,7 @@ __kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost_f16( //3x3 Convolution of elements starting in 0th row pixels0 = convolution_3x3_dilation_stridex2_stridey2_bifrost_f16(src_addr, src.stride_x, src.stride_y, 0, weights_addr, weights_stride_y); //3x3 Convolution of elements starting in 2nd row - pixels1 = convolution_3x3_dilation_stridex2_stridey2_bifrost_f16(src_addr, src.stride_x, src.stride_y, 2, weights_addr, weights_stride_y); + pixels1 = convolution_3x3_dilation_stridex2_stridey2_bifrost_f16(src_addr, src.stride_x, src.stride_y, 2, weights_addr, weights_stride_y); #endif /* DILATION_X==1 && DILATION_Y==1 */ #ifdef HAS_BIAS @@ -1492,6 +1491,153 @@ __kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost_f16( } #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(DEPTH_MULTIPLIER) && defined(DST_CHANNELS) && defined(IS_F16) +#if defined(SRC_DIM1) && defined(SRC_DIM2) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(N0) && defined(DATA_TYPE) && defined(DILATION_X) && defined(DILATION_Y) && defined(CONV_STRIDE_X) && defined(CONV_STRIDE_Y) && defined(CONV_PAD_LEFT) && defined(CONV_PAD_TOP) +/** This function computes the depthwise convolution for NHWC data layout. This kernel assumes that the weights tensor is NOT reshaped + * + * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float + * @note The number of elements processed must be passed at compile time using -DN0 (e.g. -DN0=2) + * @note The depth multiplier must be passed at compile time using -DDEPTH_MULTIPLIER (e.g. -DDEPTH_MULTIPLIER=1) + * @note The first dimension of the input tensor must be passed at compile time using -DSRC_DIM1 (e.g. -DSRC_DIM1=112) + * @note The second dimension of the input tensor must be passed at compile time using -DSRC_DIM2 (e.g. -DSRC_DIM2=80) + * @note The kernel width must be passed at compile time using -DKERNEL_WIDTH (e.g. -DKERNEL_WIDTH=5) + * @note The kernel height must be passed at compile time using -DKERNEL_HEIGHT (e.g. -DKERNEL_HEIGHT=5) + * @note The convolution pad top must be passed at compile time using -DCONV_PAD_TOP (e.g. -DCONV_PAD_TOP=1) + * @note The convolution pad top must be passed at compile time using -DCONV_PAD_LEFT (e.g. -DCONV_PAD_LEFT=1) + * @note The convolution stride along the width must be passed at compile time using -DCONV_STRIDE_X (e.g. -DCONV_STRIDE_Y=X) + * @note The convolution stride along the height must be passed at compile time using -DCONV_STRIDE_Y (e.g. -DCONV_STRIDE_Y=1) + * @note It is possible to select the activation function to apply using -DACTIVATION_TYPE e.g. -DACTIVATION_TYPE=relu + * @note A, B variables required by some activation functions are set using -DA_VAL= and -DB_VAL= respectively + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32 + * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_y * number of elements along Z processed per workitem(in bytes) + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] src_step_w src_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] dst_ptr Pointer to the destination tensor. Supported data types: same as src_ptr + * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Y processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) + * @param[in] dst_step_w dst_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: F16/F32 + * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) + * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) + * @param[in] weights_step_y weights_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes) + * @param[in] weights_step_z weights_stride_z * number of elements along Y processed per workitem(in bytes) + * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor + * @param[in] biases_ptr (Optional) Pointer to the biases vector. Supported data types: same as src_ptr + * @param[in] biases_stride_x (Optional) Stride of the biases vector in X dimension (in bytes) + * @param[in] biases_step_x (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] biases_offset_first_element_in_bytes (Optional) The offset of the first element in the biases vector + */ +__kernel void dwc_MxN_native_fp_nhwc( + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst), + TENSOR3D_DECLARATION(weights), +#if defined(HAS_BIAS) + VECTOR_DECLARATION(biases) +#endif // defined(HAS_BIAS) +) +{ + int x = get_global_id(0); // channels + int y = get_global_id(1); // spatial coordinate x +#if defined(DST_DEPTH) + int z = get_global_id(2) % (int)DST_DEPTH; // spatial coordinate y + int b = get_global_id(2) / (int)DST_DEPTH; // batch +#else // defined(DST_DEPTH) + int z = get_global_id(2); // spatial coordinate y +#endif // defined(DST_DEPTH) + + __global uchar *s_addr = src_ptr + src_offset_first_element_in_bytes + + x * sizeof(DATA_TYPE) * (int)N0; + + __global uchar *d_addr = dst_ptr + dst_offset_first_element_in_bytes + + x * sizeof(DATA_TYPE) * (int)DEPTH_MULTIPLIER * (int)N0 + + y * dst_stride_y + + z * dst_stride_z; + + __global uchar *w_addr = weights_ptr + weights_offset_first_element_in_bytes + + x * sizeof(DATA_TYPE) * (int)DEPTH_MULTIPLIER * (int)N0; + +#if defined(HAS_BIAS) + __global uchar *b_addr = biases_ptr + biases_offset_first_element_in_bytes + + x * sizeof(DATA_TYPE) * (int)DEPTH_MULTIPLIER * (int)N0; +#endif // defined(HAS_BIAS) + +#if defined(DST_DEPTH) + s_addr += b * src_stride_w; + d_addr += b * dst_stride_w; +#endif // defined(DST_DEPTH) + + for(int d = 0; d < (int)DEPTH_MULTIPLIER; ++d) + { + // Each work-item computes N0x1x1 elements + VEC_DATA_TYPE(DATA_TYPE, N0) + res = 0; + + int x_coord = y * CONV_STRIDE_X - (int)CONV_PAD_LEFT; + int y_coord = z * CONV_STRIDE_Y - (int)CONV_PAD_TOP; + + for(int yk = 0; yk < KERNEL_HEIGHT; ++yk) + { + if(y_coord >= 0 && y_coord < SRC_DIM2) + { + int x_coord_tmp = x_coord; + + for(int xk = 0; xk < KERNEL_WIDTH; ++xk) + { + if(x_coord_tmp >= 0 && x_coord_tmp < SRC_DIM1) + { + int s_offset = x_coord_tmp * (int)src_stride_y + y_coord * (int)src_stride_z; + int w_offset = xk * weights_stride_y + yk * weights_stride_z; + + // Load input and weights values + VEC_DATA_TYPE(DATA_TYPE, N0) + i = VLOAD(N0)(0, (__global DATA_TYPE *)(s_addr + s_offset)); + VEC_DATA_TYPE(DATA_TYPE, N0) + w = VLOAD(N0)(0, (__global DATA_TYPE *)(w_addr + w_offset)); + +#if GPU_ARCH == GPU_ARCH_MIDGARD + res += i * w; +#else // GPU_ARCH == GPU_ARCH_MIDGARD + res = fma(i, w, res); +#endif // GPU_ARCH == GPU_ARCH_MIDGARD + } + x_coord_tmp += DILATION_X; + } + } + y_coord += DILATION_Y; + } + +#if defined(HAS_BIAS) + res += VLOAD(N0)(0, (__global DATA_TYPE *)(b_addr)); +#endif // defined(HAS_BIAS) + + res = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, res, A_VAL, B_VAL); + + VSTORE(N0) + (res, 0, (__global DATA_TYPE *)(d_addr)); + + w_addr += sizeof(DATA_TYPE); + d_addr += sizeof(DATA_TYPE); +#if defined(HAS_BIAS) + b_addr += sizeof(DATA_TYPE); +#endif // defined(HAS_BIAS) + } +} +#endif // defined(SRC_DIM1) && defined(SRC_DIM2) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defiend(N0) && defined(DATA_TYPE) && defined(DILATION_X) && defined(DILATION_Y) && defined(CONV_STRIDE_X) && defined(CONV_STRIDE_Y) && defined(CONV_PAD_LEFT) && defined(CONV_PAD_TOP) + #if defined(VEC_SIZE) && defined(SRC_DIM_2) && defined(CONV_PAD_TOP) && defined(CONV_PAD_LEFT) && defined(DATA_TYPE) #if DATA_TYPE != float || DATA_TYPE != half @@ -1501,6 +1647,7 @@ __kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost_f16( #define VEC_FLOAT VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) #if defined(CONV_STRIDE_X) && defined(CONV_STRIDE_Y) + /** This function computes the depthwise convolution for NHWC data layout when the stride along the width or height is not 1. * * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float |