diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2018-10-23 15:23:23 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:55:45 +0000 |
| commit | c55beee7ef70fa08a5d217619083b288a74fcb27 (patch) | |
| tree | d4186463153286c0d76f84ea538005e0d5f81b3d | |
| parent | f5ec981e27633cbc34b36a537d41a6e629bc2fc4 (diff) | |
| download | ComputeLibrary-c55beee7ef70fa08a5d217619083b288a74fcb27.tar.gz | |
COMPMID-1029: Collapse CLWinogradInputTransform/CLWinogradOutputTransform
Change-Id: I051748502ca24b9952e7313524bbfd708162efb4
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/155166
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/batchnormalization_layer.cl | 4 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/winograd_input_transform.cl | 601 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/winograd_output_transform.cl | 263 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLWinogradInputTransformKernel.cpp | 38 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp | 14 |
5 files changed, 578 insertions, 342 deletions
diff --git a/src/core/CL/cl_kernels/batchnormalization_layer.cl b/src/core/CL/cl_kernels/batchnormalization_layer.cl index df141269bc..dfd16e0da3 100644 --- a/src/core/CL/cl_kernels/batchnormalization_layer.cl +++ b/src/core/CL/cl_kernels/batchnormalization_layer.cl @@ -273,8 +273,8 @@ __kernel void batchnormalization_layer_nhwc(TENSOR3D_DECLARATION(input), * @param[in] conv_w_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] conv_w_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] conv_w_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] conv_w__stride_w Stride of the source tensor in W dimension (in bytes) - * @param[in] conv_w__step_w input_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] conv_w_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] conv_w_step_w input_stride_w * number of elements along W processed per workitem(in bytes) * @param[in] conv_w_offset_first_element_in_bytes The offset of the first element in the source tensor * @param[in] bn_mean_ptr Pointer to the mean source tensor. Supported data types: same as @p input_ptr * @param[in] bn_mean_stride_x Stride of the mean source tensor in X dimension (in bytes) diff --git a/src/core/CL/cl_kernels/winograd_input_transform.cl b/src/core/CL/cl_kernels/winograd_input_transform.cl index 205e416f5d..9289cb0026 100644 --- a/src/core/CL/cl_kernels/winograd_input_transform.cl +++ b/src/core/CL/cl_kernels/winograd_input_transform.cl @@ -44,6 +44,7 @@ }) #if defined(NUM_TILES_X) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H) +#if defined(SRC_DEPTH) /** This OpenCL kernel computes the input transform when the kernel size is 3x3/3x1 or 1x3 and the output tile is 2x2/2x1 or 1x2 * * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). @@ -70,17 +71,22 @@ * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_2x2_3x3_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2) % SRC_DEPTH; + const int b = get_global_id(2) / SRC_DEPTH; // Compute input address - __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z; + __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z + b * src_stride_w; src_addr = src_addr - ((int)PAD_LEFT * sizeof(DATA_TYPE)) - ((int)PAD_TOP * src_stride_y); @@ -140,7 +146,7 @@ __kernel void winograd_input_transform_2x2_3x3_stepz1_nchw( DATA_TYPE out33 = tmp3.s1 - tmp3.s3; #endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y; + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y + b * dst_stride_w; *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z)) = out00; // in_row0.s0; out00; *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z)) = out01; // in_row0.s1; out01; @@ -189,17 +195,22 @@ __kernel void winograd_input_transform_2x2_3x3_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_2x2_3x3_stepz2_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2) * 2; + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = (get_global_id(2) * 2) % SRC_DEPTH; + const int b = (get_global_id(2) * 2) / SRC_DEPTH; // Compute input address - __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z; + __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z + b * src_stride_w; src_addr = src_addr - ((int)PAD_LEFT * sizeof(DATA_TYPE)) - ((int)PAD_TOP * src_stride_y); @@ -306,7 +317,7 @@ __kernel void winograd_input_transform_2x2_3x3_stepz2_nchw( out33 = (VEC_DATA_TYPE(DATA_TYPE, 2))(tmp3.s1 - tmp3.s3, tmp7.s1 - tmp7.s3); #endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y; + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y + b * dst_stride_w; vstore2(out00, 0, (__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z)); vstore2(out01, 0, (__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z)); @@ -355,17 +366,22 @@ __kernel void winograd_input_transform_2x2_3x3_stepz2_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_4x4_3x3_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2) % SRC_DEPTH; + const int b = get_global_id(2) / SRC_DEPTH; // Compute input address - __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z; + __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z + b * src_stride_w; src_addr = src_addr - ((int)PAD_LEFT * sizeof(DATA_TYPE)) - ((int)PAD_TOP * src_stride_y); @@ -446,7 +462,7 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nchw( #endif // #if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) // Compute destination address - __global DATA_TYPE *dst_addr = (__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y); + __global DATA_TYPE *dst_addr = (__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y + b * dst_stride_w); uint dst_plane_stride = dst_stride_z / sizeof(DATA_TYPE); @@ -637,7 +653,199 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nchw( #endif // #if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) } -#if defined(SRC_DIM_1) && defined(SRC_DIM_2) +/** This OpenCL kernel computes the input transform when the kernel size is 5x5/5x1 or 1x5 and the output tile is 4x4/4x1 or 1x4 when the data layout is NCHW + * + * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). + * @note The pad left and pad top must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (i.e.-DPAD_LEFT=1 and -DPAD_TOP=0). + * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=2 + * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=2 + * @note If this kernel is used to perform Winograd input transform 5x1, -DWINOGRAD_INPUT_TRANSFORM_HORIZONTAL has to be passed at compile time + * @note If this kernel is used to perform Winograd input transform 1x5, -DWINOGRAD_INPUT_TRANSFORM_VERTICAL has to be passed at compile time + * @note The data type must be passed at compile time using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types: float/half. + * + * @param[in] src_ptr Pointer to the source image. Supported data types: F32/F16 + * @param[in] src_stride_x Stride of the source image 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 image 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_offset_first_element_in_bytes The offset of the first element in the source image + * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Y processed per workitem(in bytes) + * @param[in] dst_ptr Pointer to the destination tensor. Supported data types: as @p 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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) + */ +__kernel void winograd_input_transform_4x4_5x5_stepz1_nchw( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2) % SRC_DEPTH; + const int b = get_global_id(2) / SRC_DEPTH; + + // Compute input address + __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z + b * src_stride_w; + + src_addr = src_addr - ((int)PAD_LEFT * sizeof(DATA_TYPE)) - ((int)PAD_TOP * src_stride_y); + + // Load input tile +#if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row0 = vload8(0, (__global DATA_TYPE *)(src_addr)); +#elif defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row0 = (VEC_DATA_TYPE(DATA_TYPE, 8))(*((__global DATA_TYPE *)(src_addr + 0 * src_stride_y)), + *((__global DATA_TYPE *)(src_addr + 1 * src_stride_y)), + *((__global DATA_TYPE *)(src_addr + 2 * src_stride_y)), + *((__global DATA_TYPE *)(src_addr + 3 * src_stride_y)), + *((__global DATA_TYPE *)(src_addr + 4 * src_stride_y)), + *((__global DATA_TYPE *)(src_addr + 5 * src_stride_y)), + *((__global DATA_TYPE *)(src_addr + 6 * src_stride_y)), + *((__global DATA_TYPE *)(src_addr + 7 * src_stride_y))); +#else // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row0 = vload8(0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y)); + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row1 = vload8(0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y)); + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row2 = vload8(0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y)); + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row3 = vload8(0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y)); + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row4 = vload8(0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y)); + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row5 = vload8(0, (__global DATA_TYPE *)(src_addr + 5 * src_stride_y)); + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row6 = vload8(0, (__global DATA_TYPE *)(src_addr + 6 * src_stride_y)); + const VEC_DATA_TYPE(DATA_TYPE, 8) in_row7 = vload8(0, (__global DATA_TYPE *)(src_addr + 7 * src_stride_y)); +#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) + + // Calculate common factors for intermediate tensor + VEC_DATA_TYPE(DATA_TYPE, 8) + tmp0 = in_row0; + VEC_DATA_TYPE(DATA_TYPE, 8) + comm_fact0 = 0.0f; + +#if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) + comm_fact0 += in_row2 + in_row6 - (DATA_TYPE)4.25 * in_row4; + tmp0 += -in_row6 + (DATA_TYPE)5.25 * in_row4 - (DATA_TYPE)5.25 * in_row2; + + VEC_DATA_TYPE(DATA_TYPE, 8) + comm_fact1 = in_row1 + in_row5 - (DATA_TYPE)4.25 * in_row3; + VEC_DATA_TYPE(DATA_TYPE, 8) + comm_fact2 = (DATA_TYPE)0.25 * in_row2 - (DATA_TYPE)1.25 * in_row4 + in_row6; + + const VEC_DATA_TYPE(DATA_TYPE, 8) tmp1 = comm_fact0 + comm_fact1; + const VEC_DATA_TYPE(DATA_TYPE, 8) tmp2 = comm_fact0 - comm_fact1; + + comm_fact0 = (DATA_TYPE)2.5 * in_row3; + comm_fact1 = (DATA_TYPE)0.5 * in_row1 - comm_fact0 + (DATA_TYPE)2.0 * in_row5; + + const VEC_DATA_TYPE(DATA_TYPE, 8) tmp3 = comm_fact1 + comm_fact2; + const VEC_DATA_TYPE(DATA_TYPE, 8) tmp4 = comm_fact2 - comm_fact1; + + comm_fact1 = (DATA_TYPE)2.0 * in_row1 - comm_fact0 + (DATA_TYPE)0.5 * in_row5; + comm_fact2 = (DATA_TYPE)4.0 * in_row2 - (DATA_TYPE)5.0 * in_row4 + in_row6; + + const VEC_DATA_TYPE(DATA_TYPE, 8) tmp5 = comm_fact1 + comm_fact2; + const VEC_DATA_TYPE(DATA_TYPE, 8) tmp6 = comm_fact2 - comm_fact1; + const VEC_DATA_TYPE(DATA_TYPE, 8) tmp7 = in_row7 - in_row1 + (DATA_TYPE)5.25 * in_row3 - (DATA_TYPE)5.25 * in_row5; +#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) + + // Calculate output rows (reuse comm_fact0 vector) + VEC_DATA_TYPE(DATA_TYPE, 8) + out0; + + OUTPUT_ROW_4x4_5x5(out0, tmp0, comm_fact0); + +#if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) + VEC_DATA_TYPE(DATA_TYPE, 8) + out1, out2, out3, out4, out5, out6, out7; + + OUTPUT_ROW_4x4_5x5(out1, tmp1, comm_fact0); + OUTPUT_ROW_4x4_5x5(out2, tmp2, comm_fact0); + OUTPUT_ROW_4x4_5x5(out3, tmp3, comm_fact0); + OUTPUT_ROW_4x4_5x5(out4, tmp4, comm_fact0); + OUTPUT_ROW_4x4_5x5(out5, tmp5, comm_fact0); + OUTPUT_ROW_4x4_5x5(out6, tmp6, comm_fact0); + OUTPUT_ROW_4x4_5x5(out7, tmp7, comm_fact0); +#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) + + // Store values across the channels + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y + b * dst_stride_w; + + *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z)) = out0.s0; + *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z)) = out0.s1; + *((__global DATA_TYPE *)(dst_addr + 2 * dst_stride_z)) = out0.s2; + *((__global DATA_TYPE *)(dst_addr + 3 * dst_stride_z)) = out0.s3; + *((__global DATA_TYPE *)(dst_addr + 4 * dst_stride_z)) = out0.s4; + *((__global DATA_TYPE *)(dst_addr + 5 * dst_stride_z)) = out0.s5; + *((__global DATA_TYPE *)(dst_addr + 6 * dst_stride_z)) = out0.s6; + *((__global DATA_TYPE *)(dst_addr + 7 * dst_stride_z)) = out0.s7; + +#if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) + *((__global DATA_TYPE *)(dst_addr + 8 * dst_stride_z)) = out1.s0; + *((__global DATA_TYPE *)(dst_addr + 9 * dst_stride_z)) = out1.s1; + *((__global DATA_TYPE *)(dst_addr + 10 * dst_stride_z)) = out1.s2; + *((__global DATA_TYPE *)(dst_addr + 11 * dst_stride_z)) = out1.s3; + *((__global DATA_TYPE *)(dst_addr + 12 * dst_stride_z)) = out1.s4; + *((__global DATA_TYPE *)(dst_addr + 13 * dst_stride_z)) = out1.s5; + *((__global DATA_TYPE *)(dst_addr + 14 * dst_stride_z)) = out1.s6; + *((__global DATA_TYPE *)(dst_addr + 15 * dst_stride_z)) = out1.s7; + *((__global DATA_TYPE *)(dst_addr + 16 * dst_stride_z)) = out2.s0; + *((__global DATA_TYPE *)(dst_addr + 17 * dst_stride_z)) = out2.s1; + *((__global DATA_TYPE *)(dst_addr + 18 * dst_stride_z)) = out2.s2; + *((__global DATA_TYPE *)(dst_addr + 19 * dst_stride_z)) = out2.s3; + *((__global DATA_TYPE *)(dst_addr + 20 * dst_stride_z)) = out2.s4; + *((__global DATA_TYPE *)(dst_addr + 21 * dst_stride_z)) = out2.s5; + *((__global DATA_TYPE *)(dst_addr + 22 * dst_stride_z)) = out2.s6; + *((__global DATA_TYPE *)(dst_addr + 23 * dst_stride_z)) = out2.s7; + *((__global DATA_TYPE *)(dst_addr + 24 * dst_stride_z)) = out3.s0; + *((__global DATA_TYPE *)(dst_addr + 25 * dst_stride_z)) = out3.s1; + *((__global DATA_TYPE *)(dst_addr + 26 * dst_stride_z)) = out3.s2; + *((__global DATA_TYPE *)(dst_addr + 27 * dst_stride_z)) = out3.s3; + *((__global DATA_TYPE *)(dst_addr + 28 * dst_stride_z)) = out3.s4; + *((__global DATA_TYPE *)(dst_addr + 29 * dst_stride_z)) = out3.s5; + *((__global DATA_TYPE *)(dst_addr + 30 * dst_stride_z)) = out3.s6; + *((__global DATA_TYPE *)(dst_addr + 31 * dst_stride_z)) = out3.s7; + *((__global DATA_TYPE *)(dst_addr + 32 * dst_stride_z)) = out4.s0; + *((__global DATA_TYPE *)(dst_addr + 33 * dst_stride_z)) = out4.s1; + *((__global DATA_TYPE *)(dst_addr + 34 * dst_stride_z)) = out4.s2; + *((__global DATA_TYPE *)(dst_addr + 35 * dst_stride_z)) = out4.s3; + *((__global DATA_TYPE *)(dst_addr + 36 * dst_stride_z)) = out4.s4; + *((__global DATA_TYPE *)(dst_addr + 37 * dst_stride_z)) = out4.s5; + *((__global DATA_TYPE *)(dst_addr + 38 * dst_stride_z)) = out4.s6; + *((__global DATA_TYPE *)(dst_addr + 39 * dst_stride_z)) = out4.s7; + *((__global DATA_TYPE *)(dst_addr + 40 * dst_stride_z)) = out5.s0; + *((__global DATA_TYPE *)(dst_addr + 41 * dst_stride_z)) = out5.s1; + *((__global DATA_TYPE *)(dst_addr + 42 * dst_stride_z)) = out5.s2; + *((__global DATA_TYPE *)(dst_addr + 43 * dst_stride_z)) = out5.s3; + *((__global DATA_TYPE *)(dst_addr + 44 * dst_stride_z)) = out5.s4; + *((__global DATA_TYPE *)(dst_addr + 45 * dst_stride_z)) = out5.s5; + *((__global DATA_TYPE *)(dst_addr + 46 * dst_stride_z)) = out5.s6; + *((__global DATA_TYPE *)(dst_addr + 47 * dst_stride_z)) = out5.s7; + *((__global DATA_TYPE *)(dst_addr + 48 * dst_stride_z)) = out6.s0; + *((__global DATA_TYPE *)(dst_addr + 49 * dst_stride_z)) = out6.s1; + *((__global DATA_TYPE *)(dst_addr + 50 * dst_stride_z)) = out6.s2; + *((__global DATA_TYPE *)(dst_addr + 51 * dst_stride_z)) = out6.s3; + *((__global DATA_TYPE *)(dst_addr + 52 * dst_stride_z)) = out6.s4; + *((__global DATA_TYPE *)(dst_addr + 53 * dst_stride_z)) = out6.s5; + *((__global DATA_TYPE *)(dst_addr + 54 * dst_stride_z)) = out6.s6; + *((__global DATA_TYPE *)(dst_addr + 55 * dst_stride_z)) = out6.s7; + *((__global DATA_TYPE *)(dst_addr + 56 * dst_stride_z)) = out7.s0; + *((__global DATA_TYPE *)(dst_addr + 57 * dst_stride_z)) = out7.s1; + *((__global DATA_TYPE *)(dst_addr + 58 * dst_stride_z)) = out7.s2; + *((__global DATA_TYPE *)(dst_addr + 59 * dst_stride_z)) = out7.s3; + *((__global DATA_TYPE *)(dst_addr + 60 * dst_stride_z)) = out7.s4; + *((__global DATA_TYPE *)(dst_addr + 61 * dst_stride_z)) = out7.s5; + *((__global DATA_TYPE *)(dst_addr + 62 * dst_stride_z)) = out7.s6; + *((__global DATA_TYPE *)(dst_addr + 63 * dst_stride_z)) = out7.s7; +#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) +} +#endif // defined(SRC_DEPTH) + +#if defined(NUM_TILES_Y) && defined(SRC_DIM_1) && defined(SRC_DIM_2) /** This OpenCL kernel computes the input transform when the output tile is 4x4, 4x1 or 1x4, the filter size 3x3, 3x1 or 1x3 and the data layout is NHWC * * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). @@ -666,16 +874,21 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2) % NUM_TILES_Y; + const int b = get_global_id(2) / NUM_TILES_Y; - __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * sizeof(DATA_TYPE); + __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * sizeof(DATA_TYPE) + b * src_stride_w; // Clamp coordinates. This clamp is valid for all rows int4 y_coord0 = (int4)(y * OUTPUT_TILE_W) + (int4)(0, 1, 2, 3) - (int4)PAD_LEFT; @@ -828,7 +1041,8 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc( #endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) // Compute destination address - __global DATA_TYPE *dst_addr = (__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + x * sizeof(DATA_TYPE) + (y + z * (int)NUM_TILES_X) * dst_stride_y); + __global DATA_TYPE *dst_addr = (__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + x * sizeof(DATA_TYPE) + (y + z * + (int)NUM_TILES_X) * dst_stride_y + b * dst_stride_w); uint dst_plane_stride = dst_stride_z / sizeof(DATA_TYPE); *((__global DATA_TYPE *)dst_addr) = out0; @@ -1048,17 +1262,22 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); + const int x = get_global_id(0); + const int y = get_global_id(1); + const int z = get_global_id(2) % NUM_TILES_Y; + const int b = get_global_id(2) / NUM_TILES_Y; // Compute input address - __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * sizeof(DATA_TYPE); + __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * sizeof(DATA_TYPE) + b * src_stride_w; #if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) // Clamp coordinates. This clamp is valid for all rows @@ -1293,194 +1512,7 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( #endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) // Store values across the channels - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x * sizeof(DATA_TYPE) + (y + z * (int)NUM_TILES_X) * dst_stride_y; - - *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z)) = out0.s0; - *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z)) = out0.s1; - *((__global DATA_TYPE *)(dst_addr + 2 * dst_stride_z)) = out0.s2; - *((__global DATA_TYPE *)(dst_addr + 3 * dst_stride_z)) = out0.s3; - *((__global DATA_TYPE *)(dst_addr + 4 * dst_stride_z)) = out0.s4; - *((__global DATA_TYPE *)(dst_addr + 5 * dst_stride_z)) = out0.s5; - *((__global DATA_TYPE *)(dst_addr + 6 * dst_stride_z)) = out0.s6; - *((__global DATA_TYPE *)(dst_addr + 7 * dst_stride_z)) = out0.s7; - -#if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - *((__global DATA_TYPE *)(dst_addr + 8 * dst_stride_z)) = out1.s0; - *((__global DATA_TYPE *)(dst_addr + 9 * dst_stride_z)) = out1.s1; - *((__global DATA_TYPE *)(dst_addr + 10 * dst_stride_z)) = out1.s2; - *((__global DATA_TYPE *)(dst_addr + 11 * dst_stride_z)) = out1.s3; - *((__global DATA_TYPE *)(dst_addr + 12 * dst_stride_z)) = out1.s4; - *((__global DATA_TYPE *)(dst_addr + 13 * dst_stride_z)) = out1.s5; - *((__global DATA_TYPE *)(dst_addr + 14 * dst_stride_z)) = out1.s6; - *((__global DATA_TYPE *)(dst_addr + 15 * dst_stride_z)) = out1.s7; - *((__global DATA_TYPE *)(dst_addr + 16 * dst_stride_z)) = out2.s0; - *((__global DATA_TYPE *)(dst_addr + 17 * dst_stride_z)) = out2.s1; - *((__global DATA_TYPE *)(dst_addr + 18 * dst_stride_z)) = out2.s2; - *((__global DATA_TYPE *)(dst_addr + 19 * dst_stride_z)) = out2.s3; - *((__global DATA_TYPE *)(dst_addr + 20 * dst_stride_z)) = out2.s4; - *((__global DATA_TYPE *)(dst_addr + 21 * dst_stride_z)) = out2.s5; - *((__global DATA_TYPE *)(dst_addr + 22 * dst_stride_z)) = out2.s6; - *((__global DATA_TYPE *)(dst_addr + 23 * dst_stride_z)) = out2.s7; - *((__global DATA_TYPE *)(dst_addr + 24 * dst_stride_z)) = out3.s0; - *((__global DATA_TYPE *)(dst_addr + 25 * dst_stride_z)) = out3.s1; - *((__global DATA_TYPE *)(dst_addr + 26 * dst_stride_z)) = out3.s2; - *((__global DATA_TYPE *)(dst_addr + 27 * dst_stride_z)) = out3.s3; - *((__global DATA_TYPE *)(dst_addr + 28 * dst_stride_z)) = out3.s4; - *((__global DATA_TYPE *)(dst_addr + 29 * dst_stride_z)) = out3.s5; - *((__global DATA_TYPE *)(dst_addr + 30 * dst_stride_z)) = out3.s6; - *((__global DATA_TYPE *)(dst_addr + 31 * dst_stride_z)) = out3.s7; - *((__global DATA_TYPE *)(dst_addr + 32 * dst_stride_z)) = out4.s0; - *((__global DATA_TYPE *)(dst_addr + 33 * dst_stride_z)) = out4.s1; - *((__global DATA_TYPE *)(dst_addr + 34 * dst_stride_z)) = out4.s2; - *((__global DATA_TYPE *)(dst_addr + 35 * dst_stride_z)) = out4.s3; - *((__global DATA_TYPE *)(dst_addr + 36 * dst_stride_z)) = out4.s4; - *((__global DATA_TYPE *)(dst_addr + 37 * dst_stride_z)) = out4.s5; - *((__global DATA_TYPE *)(dst_addr + 38 * dst_stride_z)) = out4.s6; - *((__global DATA_TYPE *)(dst_addr + 39 * dst_stride_z)) = out4.s7; - *((__global DATA_TYPE *)(dst_addr + 40 * dst_stride_z)) = out5.s0; - *((__global DATA_TYPE *)(dst_addr + 41 * dst_stride_z)) = out5.s1; - *((__global DATA_TYPE *)(dst_addr + 42 * dst_stride_z)) = out5.s2; - *((__global DATA_TYPE *)(dst_addr + 43 * dst_stride_z)) = out5.s3; - *((__global DATA_TYPE *)(dst_addr + 44 * dst_stride_z)) = out5.s4; - *((__global DATA_TYPE *)(dst_addr + 45 * dst_stride_z)) = out5.s5; - *((__global DATA_TYPE *)(dst_addr + 46 * dst_stride_z)) = out5.s6; - *((__global DATA_TYPE *)(dst_addr + 47 * dst_stride_z)) = out5.s7; - *((__global DATA_TYPE *)(dst_addr + 48 * dst_stride_z)) = out6.s0; - *((__global DATA_TYPE *)(dst_addr + 49 * dst_stride_z)) = out6.s1; - *((__global DATA_TYPE *)(dst_addr + 50 * dst_stride_z)) = out6.s2; - *((__global DATA_TYPE *)(dst_addr + 51 * dst_stride_z)) = out6.s3; - *((__global DATA_TYPE *)(dst_addr + 52 * dst_stride_z)) = out6.s4; - *((__global DATA_TYPE *)(dst_addr + 53 * dst_stride_z)) = out6.s5; - *((__global DATA_TYPE *)(dst_addr + 54 * dst_stride_z)) = out6.s6; - *((__global DATA_TYPE *)(dst_addr + 55 * dst_stride_z)) = out6.s7; - *((__global DATA_TYPE *)(dst_addr + 56 * dst_stride_z)) = out7.s0; - *((__global DATA_TYPE *)(dst_addr + 57 * dst_stride_z)) = out7.s1; - *((__global DATA_TYPE *)(dst_addr + 58 * dst_stride_z)) = out7.s2; - *((__global DATA_TYPE *)(dst_addr + 59 * dst_stride_z)) = out7.s3; - *((__global DATA_TYPE *)(dst_addr + 60 * dst_stride_z)) = out7.s4; - *((__global DATA_TYPE *)(dst_addr + 61 * dst_stride_z)) = out7.s5; - *((__global DATA_TYPE *)(dst_addr + 62 * dst_stride_z)) = out7.s6; - *((__global DATA_TYPE *)(dst_addr + 63 * dst_stride_z)) = out7.s7; -#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) -} -#endif // defined(SRC_DIM_1) && defined(SRC_DIM_2) - -/** This OpenCL kernel computes the input transform when the kernel size is 5x5/5x1 or 1x5 and the output tile is 4x4/4x1 or 1x4 when the data layout is NCHW - * - * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). - * @note The pad left and pad top must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (i.e.-DPAD_LEFT=1 and -DPAD_TOP=0). - * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=2 - * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=2 - * @note If this kernel is used to perform Winograd input transform 5x1, -DWINOGRAD_INPUT_TRANSFORM_HORIZONTAL has to be passed at compile time - * @note If this kernel is used to perform Winograd input transform 1x5, -DWINOGRAD_INPUT_TRANSFORM_VERTICAL has to be passed at compile time - * @note The data type must be passed at compile time using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types: float/half. - * - * @param[in] src_ptr Pointer to the source image. Supported data types: F32/F16 - * @param[in] src_stride_x Stride of the source image 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 image 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_offset_first_element_in_bytes The offset of the first element in the source image - * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Y processed per workitem(in bytes) - * @param[in] dst_ptr Pointer to the destination tensor. Supported data types: as @p 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_offset_first_element_in_bytes The offset of the first element in the destination tensor - */ -__kernel void winograd_input_transform_4x4_5x5_stepz1_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) -{ - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); - - // Compute input address - __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x * OUTPUT_TILE_W * sizeof(DATA_TYPE) + y * OUTPUT_TILE_H * src_stride_y + z * src_stride_z; - - src_addr = src_addr - ((int)PAD_LEFT * sizeof(DATA_TYPE)) - ((int)PAD_TOP * src_stride_y); - - // Load input tile -#if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row0 = vload8(0, (__global DATA_TYPE *)(src_addr)); -#elif defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) // !defined(WINOGRAD_FILTER_TRANSFORM_HORIZONTAL) - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row0 = (VEC_DATA_TYPE(DATA_TYPE, 8))(*((__global DATA_TYPE *)(src_addr + 0 * src_stride_y)), - *((__global DATA_TYPE *)(src_addr + 1 * src_stride_y)), - *((__global DATA_TYPE *)(src_addr + 2 * src_stride_y)), - *((__global DATA_TYPE *)(src_addr + 3 * src_stride_y)), - *((__global DATA_TYPE *)(src_addr + 4 * src_stride_y)), - *((__global DATA_TYPE *)(src_addr + 5 * src_stride_y)), - *((__global DATA_TYPE *)(src_addr + 6 * src_stride_y)), - *((__global DATA_TYPE *)(src_addr + 7 * src_stride_y))); -#else // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row0 = vload8(0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y)); - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row1 = vload8(0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y)); - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row2 = vload8(0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y)); - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row3 = vload8(0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y)); - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row4 = vload8(0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y)); - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row5 = vload8(0, (__global DATA_TYPE *)(src_addr + 5 * src_stride_y)); - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row6 = vload8(0, (__global DATA_TYPE *)(src_addr + 6 * src_stride_y)); - const VEC_DATA_TYPE(DATA_TYPE, 8) in_row7 = vload8(0, (__global DATA_TYPE *)(src_addr + 7 * src_stride_y)); -#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - - // Calculate common factors for intermediate tensor - VEC_DATA_TYPE(DATA_TYPE, 8) - tmp0 = in_row0; - VEC_DATA_TYPE(DATA_TYPE, 8) - comm_fact0 = 0.0f; - -#if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - comm_fact0 += in_row2 + in_row6 - (DATA_TYPE)4.25 * in_row4; - tmp0 += -in_row6 + (DATA_TYPE)5.25 * in_row4 - (DATA_TYPE)5.25 * in_row2; - - VEC_DATA_TYPE(DATA_TYPE, 8) - comm_fact1 = in_row1 + in_row5 - (DATA_TYPE)4.25 * in_row3; - VEC_DATA_TYPE(DATA_TYPE, 8) - comm_fact2 = (DATA_TYPE)0.25 * in_row2 - (DATA_TYPE)1.25 * in_row4 + in_row6; - - const VEC_DATA_TYPE(DATA_TYPE, 8) tmp1 = comm_fact0 + comm_fact1; - const VEC_DATA_TYPE(DATA_TYPE, 8) tmp2 = comm_fact0 - comm_fact1; - - comm_fact0 = (DATA_TYPE)2.5 * in_row3; - comm_fact1 = (DATA_TYPE)0.5 * in_row1 - comm_fact0 + (DATA_TYPE)2.0 * in_row5; - - const VEC_DATA_TYPE(DATA_TYPE, 8) tmp3 = comm_fact1 + comm_fact2; - const VEC_DATA_TYPE(DATA_TYPE, 8) tmp4 = comm_fact2 - comm_fact1; - - comm_fact1 = (DATA_TYPE)2.0 * in_row1 - comm_fact0 + (DATA_TYPE)0.5 * in_row5; - comm_fact2 = (DATA_TYPE)4.0 * in_row2 - (DATA_TYPE)5.0 * in_row4 + in_row6; - - const VEC_DATA_TYPE(DATA_TYPE, 8) tmp5 = comm_fact1 + comm_fact2; - const VEC_DATA_TYPE(DATA_TYPE, 8) tmp6 = comm_fact2 - comm_fact1; - const VEC_DATA_TYPE(DATA_TYPE, 8) tmp7 = in_row7 - in_row1 + (DATA_TYPE)5.25 * in_row3 - (DATA_TYPE)5.25 * in_row5; -#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - - // Calculate output rows (reuse comm_fact0 vector) - VEC_DATA_TYPE(DATA_TYPE, 8) - out0; - - OUTPUT_ROW_4x4_5x5(out0, tmp0, comm_fact0); - -#if !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - VEC_DATA_TYPE(DATA_TYPE, 8) - out1, out2, out3, out4, out5, out6, out7; - - OUTPUT_ROW_4x4_5x5(out1, tmp1, comm_fact0); - OUTPUT_ROW_4x4_5x5(out2, tmp2, comm_fact0); - OUTPUT_ROW_4x4_5x5(out3, tmp3, comm_fact0); - OUTPUT_ROW_4x4_5x5(out4, tmp4, comm_fact0); - OUTPUT_ROW_4x4_5x5(out5, tmp5, comm_fact0); - OUTPUT_ROW_4x4_5x5(out6, tmp6, comm_fact0); - OUTPUT_ROW_4x4_5x5(out7, tmp7, comm_fact0); -#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - - // Store values across the channels - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + z * sizeof(DATA_TYPE) + (x + y * (int)NUM_TILES_X) * dst_stride_y; + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x * sizeof(DATA_TYPE) + (y + z * (int)NUM_TILES_X) * dst_stride_y + b * dst_stride_w; *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_z)) = out0.s0; *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_z)) = out0.s1; @@ -1550,8 +1582,10 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nchw( *((__global DATA_TYPE *)(dst_addr + 63 * dst_stride_z)) = out7.s7; #endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) } +#endif // defined(NUM_TILES_Y) && defined(SRC_DIM_1) && defined(SRC_DIM_2) #if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) +#if defined(SRC_DEPTH) /** This OpenCL kernel computes the input transform when the kernel size is 3x1 and the output tile is 2x1 * * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). @@ -1577,10 +1611,14 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_2x1_3x1_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_2x2_3x3_stepz1_nchw(src_ptr, src_stride_x, @@ -1597,7 +1635,9 @@ __kernel void winograd_input_transform_2x1_3x1_stepz1_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 3x1, the output tile is 2x1 and the number of channels is multiple of 2 @@ -1625,10 +1665,14 @@ __kernel void winograd_input_transform_2x1_3x1_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_2x1_3x1_stepz2_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_2x2_3x3_stepz2_nchw(src_ptr, src_stride_x, @@ -1645,7 +1689,9 @@ __kernel void winograd_input_transform_2x1_3x1_stepz2_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 3x1 and the output tile is 4x1 @@ -1673,10 +1719,14 @@ __kernel void winograd_input_transform_2x1_3x1_stepz2_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_4x1_3x1_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_3x3_stepz1_nchw(src_ptr, src_stride_x, @@ -1693,7 +1743,9 @@ __kernel void winograd_input_transform_4x1_3x1_stepz1_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 5x1 and the output tile is 4x1 when the data layout is NCHW @@ -1721,10 +1773,14 @@ __kernel void winograd_input_transform_4x1_3x1_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_4x1_5x1_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_5x5_stepz1_nchw(src_ptr, src_stride_x, @@ -1741,10 +1797,13 @@ __kernel void winograd_input_transform_4x1_5x1_stepz1_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } +#endif // defined(SRC_DEPTH) -#if defined(SRC_DIM_1) && defined(SRC_DIM_2) +#if defined(NUM_TILES_Y) && defined(SRC_DIM_1) && defined(SRC_DIM_2) /** This OpenCL kernel computes the input transform when the kernel size is 3x1 and the output tile is 4x1 for data layout NHWC * * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). @@ -1772,10 +1831,14 @@ __kernel void winograd_input_transform_4x1_5x1_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_4x1_3x1_stepz1_nhwc( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_3x3_stepz1_nhwc(src_ptr, src_stride_x, @@ -1792,7 +1855,9 @@ __kernel void winograd_input_transform_4x1_3x1_stepz1_nhwc( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 5x1 and the output tile is 4x1 for data layout NHWC @@ -1822,10 +1887,14 @@ __kernel void winograd_input_transform_4x1_3x1_stepz1_nhwc( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_4x1_5x1_stepz1_nhwc( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_5x5_stepz1_nhwc(src_ptr, src_stride_x, @@ -1842,12 +1911,15 @@ __kernel void winograd_input_transform_4x1_5x1_stepz1_nhwc( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } -#endif // defined(SRC_DIM_1) && defined(SRC_DIM_2) +#endif // defined(NUM_TILES_Y) && defined(SRC_DIM_1) && defined(SRC_DIM_2) #endif // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) #if defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) +#if defined(SRC_DEPTH) /** This OpenCL kernel computes the input transform when the kernel size is 1x3 and the output tile is 1x2 * * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). @@ -1873,10 +1945,14 @@ __kernel void winograd_input_transform_4x1_5x1_stepz1_nhwc( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_1x2_1x3_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_2x2_3x3_stepz1_nchw(src_ptr, src_stride_x, @@ -1893,7 +1969,9 @@ __kernel void winograd_input_transform_1x2_1x3_stepz1_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 1x3, the output tile is 1x2 and the number of channels is multiple of 2 @@ -1921,10 +1999,14 @@ __kernel void winograd_input_transform_1x2_1x3_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_1x2_1x3_stepz2_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_2x2_3x3_stepz2_nchw(src_ptr, src_stride_x, @@ -1941,7 +2023,9 @@ __kernel void winograd_input_transform_1x2_1x3_stepz2_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 1x3 and the output tile is 1x4 @@ -1969,10 +2053,14 @@ __kernel void winograd_input_transform_1x2_1x3_stepz2_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_1x4_1x3_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_3x3_stepz1_nchw(src_ptr, src_stride_x, @@ -1989,7 +2077,9 @@ __kernel void winograd_input_transform_1x4_1x3_stepz1_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 1x5 and the output tile is 1x4 @@ -2017,10 +2107,14 @@ __kernel void winograd_input_transform_1x4_1x3_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_1x4_1x5_stepz1_nchw( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_5x5_stepz1_nchw(src_ptr, src_stride_x, @@ -2037,10 +2131,13 @@ __kernel void winograd_input_transform_1x4_1x5_stepz1_nchw( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } +#endif // defined(SRC_DEPTH) -#if defined(SRC_DIM_1) && defined(SRC_DIM_2) +#if defined(NUM_TILES_Y) && defined(SRC_DIM_1) && defined(SRC_DIM_2) /** This OpenCL kernel computes the input transform when the kernel size is 1x3 and the output tile is 1x4 for data layout NHWC * * @note The number of tiles in the x axis must be passed at compile time using -DNUM_TILES_X (i.e.-DNUM_TILES_X=5). @@ -2068,10 +2165,14 @@ __kernel void winograd_input_transform_1x4_1x5_stepz1_nchw( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_1x4_1x3_stepz1_nhwc( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_3x3_stepz1_nhwc(src_ptr, src_stride_x, @@ -2088,7 +2189,9 @@ __kernel void winograd_input_transform_1x4_1x3_stepz1_nhwc( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } /** This OpenCL kernel computes the input transform when the kernel size is 1x5 and the output tile is 1x4 for data layout NHWC @@ -2118,10 +2221,14 @@ __kernel void winograd_input_transform_1x4_1x3_stepz1_nhwc( * @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_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void winograd_input_transform_1x4_1x5_stepz1_nhwc( TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst)) + TENSOR3D_DECLARATION(dst), + uint src_stride_w, + uint dst_stride_w) { winograd_input_transform_4x4_5x5_stepz1_nhwc(src_ptr, src_stride_x, @@ -2138,8 +2245,10 @@ __kernel void winograd_input_transform_1x4_1x5_stepz1_nhwc( dst_step_y, dst_stride_z, dst_step_z, - dst_offset_first_element_in_bytes); + dst_offset_first_element_in_bytes, + src_stride_w, + dst_stride_w); } -#endif // defined(SRC_DIM_1) && defined(SRC_DIM_2) +#endif // defined(NUM_TILES_Y) && defined(SRC_DIM_1) && defined(SRC_DIM_2) #endif // defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) #endif // defined(NUM_TILES_X) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
\ No newline at end of file diff --git a/src/core/CL/cl_kernels/winograd_output_transform.cl b/src/core/CL/cl_kernels/winograd_output_transform.cl index 9be51f27ec..bae40f3762 100644 --- a/src/core/CL/cl_kernels/winograd_output_transform.cl +++ b/src/core/CL/cl_kernels/winograd_output_transform.cl @@ -23,7 +23,7 @@ */ #include "helpers.h" -#if defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H) +#if defined(SRC_DEPTH) && defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H) /** This OpenCL kernel performs Winograd output transform when the output tile is 2x2/2x1 or 1x2, the filter size 3x3/3x1 or 1x3 and the data layout is NCHW * * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16 @@ -40,19 +40,23 @@ * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_2x2_3x3_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -60,9 +64,9 @@ __kernel void winograd_output_transform_2x2_3x3_nchw( ) { // Each thread stores a 2x2/2x1 or 1x2 tile accordingly with the filter size - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); - const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); + const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); // Load the values across the 16 or 4 channels to compose the 4x4 or 4x1 tile DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); @@ -119,6 +123,7 @@ __kernel void winograd_output_transform_2x2_3x3_nchw( int x_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; int z_out = get_global_id(0); + int batch = get_global_id(2) / SRC_DEPTH; #if defined(HAS_BIAS) // Add bias @@ -131,7 +136,7 @@ __kernel void winograd_output_transform_2x2_3x3_nchw( #endif // defined(HAS_BIAS) // Get output address - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z; + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w; // Store the output tile #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) @@ -168,19 +173,23 @@ __kernel void winograd_output_transform_2x2_3x3_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_4x4_3x3_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -188,9 +197,9 @@ __kernel void winograd_output_transform_4x4_3x3_nchw( ) { // Each thread stores a 4x4/4x1 or 1x4 tile - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); - const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); + const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); // Load the values across the channels to compose the 6x6 or 6x1 tile DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); @@ -303,6 +312,7 @@ __kernel void winograd_output_transform_4x4_3x3_nchw( int x_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; int z_out = get_global_id(0); + int batch = get_global_id(2) / SRC_DEPTH; #if defined(HAS_BIAS) // Add bias @@ -317,7 +327,7 @@ __kernel void winograd_output_transform_4x4_3x3_nchw( #endif // defined(HAS_BIAS) // Get output address - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z; + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w; // Store the output tile #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) @@ -369,29 +379,33 @@ __kernel void winograd_output_transform_4x4_3x3_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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] dst_size Size of the destination tensor, minus the last padding */ __kernel void winograd_output_transform_4x4_3x3_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size) { // Each thread stores a 4x4/4x1 or 1x4 tile - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); - const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); + const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); // Load the values across the 36 channels to compose the 6x6 or 6x1 tile DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); @@ -505,6 +519,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( int x_out = get_global_id(0); int y_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; int z_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; + int batch = get_global_id(2) / SRC_DEPTH; #if defined(HAS_BIAS) // Add bias @@ -536,7 +551,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( #endif // defined(HAS_BIAS) #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) - int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z); + int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w); offset = min(offset + (int4)(0, 1, 2, 3) * (int4)dst_stride_z, (int4)dst_size); // If address is beyond the last plane, clamp it to dst_size (which points to the last padding). // Store the 1x4 output tile @@ -555,7 +570,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( *((__global DATA_TYPE *)(dst_ptr + mult_y * 3 * dst_stride_y + offset)) = out03; #else // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) // Get output address - int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z); + int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w); offset = min(offset + (int4)(0, 1, 2, 3) * (int4)dst_stride_z, (int4)dst_size); // If address is beyond the last plane, clamp it to dst_size (which points to the last padding). int4 mult_y = min((int4)dst_size - offset, (int4)1); // If out of bound, we don't want to increase dst_stride_y, so we set the multiplier to 0. It will be 1 otherwise. @@ -613,19 +628,23 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_4x4_5x5_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -633,17 +652,18 @@ __kernel void winograd_output_transform_4x4_5x5_nchw( ) { // Each thread stores a 4x4/4x1 or 1x4 tile - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); - const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); + const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); // Compute output address int y_in = get_global_id(1); int x_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; int z_out = get_global_id(0); + int batch = get_global_id(2) / SRC_DEPTH; - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z; + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w; // Load the values across the channels to compose the input tile DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); @@ -818,33 +838,38 @@ __kernel void winograd_output_transform_4x4_5x5_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_4x4_5x5_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size) { // Each thread stores a 4x4/4x1 or 1x4 tile - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); - const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); + const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); int y_in = get_global_id(1); int x_out = get_global_id(0); int y_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; int z_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; + int batch = get_global_id(2) / SRC_DEPTH; // Load the values across the channels to compose the input tile DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); @@ -878,7 +903,7 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( // Store the output tile #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) // Get output address - int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z); + int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w); offset = min(offset + (int4)(0, 1, 2, 3) * (int4)dst_stride_z, (int4)dst_size); // If address is beyond the last plane, clamp it to dst_size (which points to the last padding). *(__global DATA_TYPE *)(dst_ptr + offset.s0) = out00; @@ -1006,7 +1031,7 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( out_col3 += (VEC_DATA_TYPE(DATA_TYPE, 4))b; #endif // defined(HAS_BIAS) // Get output address - int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z); + int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w); offset = min(offset + (int4)(0, 1, 2, 3) * (int4)dst_stride_z, (int4)dst_size); // If address is beyond the last plane, clamp it to dst_size (which points to the last padding). int4 mult_y = min((int4)dst_size - offset, (int4)1); // If out of bound, we don't want to increase dst_stride_y, so we set the multiplier to 0. It will be 1 otherwise. @@ -1046,19 +1071,23 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_2x1_3x1_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -1072,6 +1101,8 @@ __kernel void winograd_output_transform_2x1_3x1_nchw( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1080,6 +1111,8 @@ __kernel void winograd_output_transform_2x1_3x1_nchw( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes #if defined(HAS_BIAS) , @@ -1106,19 +1139,23 @@ __kernel void winograd_output_transform_2x1_3x1_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_4x1_3x1_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -1132,6 +1169,8 @@ __kernel void winograd_output_transform_4x1_3x1_nchw( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1140,6 +1179,8 @@ __kernel void winograd_output_transform_4x1_3x1_nchw( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes #if defined(HAS_BIAS) , @@ -1166,19 +1207,23 @@ __kernel void winograd_output_transform_4x1_3x1_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_4x1_5x1_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -1192,6 +1237,8 @@ __kernel void winograd_output_transform_4x1_5x1_nchw( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1200,6 +1247,8 @@ __kernel void winograd_output_transform_4x1_5x1_nchw( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes #if defined(HAS_BIAS) , @@ -1226,19 +1275,23 @@ __kernel void winograd_output_transform_4x1_5x1_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_4x1_3x1_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) @@ -1251,6 +1304,8 @@ __kernel void winograd_output_transform_4x1_3x1_nhwc( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1259,6 +1314,8 @@ __kernel void winograd_output_transform_4x1_3x1_nhwc( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes, #if defined(HAS_BIAS) bias_ptr, @@ -1284,19 +1341,23 @@ __kernel void winograd_output_transform_4x1_3x1_nhwc( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_4x1_5x1_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) @@ -1309,6 +1370,8 @@ __kernel void winograd_output_transform_4x1_5x1_nhwc( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1317,6 +1380,8 @@ __kernel void winograd_output_transform_4x1_5x1_nhwc( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes, #if defined(HAS_BIAS) bias_ptr, @@ -1344,19 +1409,23 @@ __kernel void winograd_output_transform_4x1_5x1_nhwc( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_1x2_1x3_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -1370,6 +1439,8 @@ __kernel void winograd_output_transform_1x2_1x3_nchw( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1378,6 +1449,8 @@ __kernel void winograd_output_transform_1x2_1x3_nchw( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes #if defined(HAS_BIAS) , @@ -1404,19 +1477,23 @@ __kernel void winograd_output_transform_1x2_1x3_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_1x4_1x3_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -1430,6 +1507,8 @@ __kernel void winograd_output_transform_1x4_1x3_nchw( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1438,6 +1517,8 @@ __kernel void winograd_output_transform_1x4_1x3_nchw( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes #if defined(HAS_BIAS) , @@ -1464,19 +1545,23 @@ __kernel void winograd_output_transform_1x4_1x3_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_1x4_1x5_nchw( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst) + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bias) @@ -1490,6 +1575,8 @@ __kernel void winograd_output_transform_1x4_1x5_nchw( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1498,6 +1585,8 @@ __kernel void winograd_output_transform_1x4_1x5_nchw( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes #if defined(HAS_BIAS) , @@ -1524,19 +1613,23 @@ __kernel void winograd_output_transform_1x4_1x5_nchw( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_1x4_1x3_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) @@ -1549,6 +1642,8 @@ __kernel void winograd_output_transform_1x4_1x3_nhwc( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1557,6 +1652,8 @@ __kernel void winograd_output_transform_1x4_1x3_nhwc( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes, #if defined(HAS_BIAS) bias_ptr, @@ -1582,19 +1679,23 @@ __kernel void winograd_output_transform_1x4_1x3_nhwc( * @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_z * 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[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p 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] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_w Stride of the source 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 */ __kernel void winograd_output_transform_1x4_1x5_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), + TENSOR4D_DECLARATION(src), + TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) @@ -1607,6 +1708,8 @@ __kernel void winograd_output_transform_1x4_1x5_nhwc( src_step_y, src_stride_z, src_step_z, + src_stride_w, + src_step_w, src_offset_first_element_in_bytes, dst_ptr, dst_stride_x, @@ -1615,6 +1718,8 @@ __kernel void winograd_output_transform_1x4_1x5_nhwc( dst_step_y, dst_stride_z, dst_step_z, + dst_stride_w, + dst_step_w, dst_offset_first_element_in_bytes, #if defined(HAS_BIAS) bias_ptr, @@ -1625,4 +1730,4 @@ __kernel void winograd_output_transform_1x4_1x5_nhwc( dst_size); } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) -#endif // defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H) +#endif // defined(SRC_DEPTH) && defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H) diff --git a/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp b/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp index 04067319b0..f76ade1d32 100644 --- a/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp +++ b/src/core/CL/kernels/CLWinogradInputTransformKernel.cpp @@ -116,6 +116,7 @@ void CLWinogradInputTransformKernel::configure(const ICLTensor *input, ICLTensor const PadStrideInfo conv_info = winograd_info.convolution_info; const Size2D output_tile_size = winograd_info.output_tile_size; const Size2D kernel_size = winograd_info.kernel_size; + const DataLayout data_layout = input->info()->data_layout(); const size_t idx_w = get_data_layout_dimension_index(input->info()->data_layout(), DataLayoutDimension::WIDTH); const size_t idx_h = get_data_layout_dimension_index(input->info()->data_layout(), DataLayoutDimension::HEIGHT); @@ -124,7 +125,7 @@ void CLWinogradInputTransformKernel::configure(const ICLTensor *input, ICLTensor const int num_elements_x = input->info()->dimension(idx_w) - (kernel_size.width - 1) + conv_info.pad_left() + conv_info.pad_right(); const int num_elements_y = input->info()->dimension(idx_h) - (kernel_size.height - 1) + conv_info.pad_top() + conv_info.pad_bottom(); - if(input->info()->data_layout() == DataLayout::NCHW) + if(data_layout == DataLayout::NCHW) { // Check if we need to extend the right or bottom border const unsigned int extra_border_right = ((num_elements_x % output_tile_size.width) == 0) ? 0u : static_cast<unsigned int>(output_tile_size.width - 1); @@ -164,12 +165,16 @@ void CLWinogradInputTransformKernel::configure(const ICLTensor *input, ICLTensor build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); build_opts.add_option_if(winograd_info.kernel_size.height == 1, "-DWINOGRAD_INPUT_TRANSFORM_HORIZONTAL"); build_opts.add_option_if(winograd_info.kernel_size.width == 1, "-DWINOGRAD_INPUT_TRANSFORM_VERTICAL"); - - if(input->info()->data_layout() == DataLayout::NHWC) + if(data_layout == DataLayout::NHWC) { + build_opts.add_option("-DNUM_TILES_Y=" + support::cpp11::to_string(_num_tiles_y)); build_opts.add_option("-DSRC_DIM_1=" + support::cpp11::to_string(_input->info()->dimension(1))); build_opts.add_option("-DSRC_DIM_2=" + support::cpp11::to_string(_input->info()->dimension(2))); } + else + { + build_opts.add_option("-DSRC_DEPTH=" + support::cpp11::to_string(_input->info()->dimension(2))); + } // Create kernel std::string kernel_name = "winograd_input_transform_" + output_tile_size.to_string() + "_" + kernel_size.to_string(); @@ -178,7 +183,7 @@ void CLWinogradInputTransformKernel::configure(const ICLTensor *input, ICLTensor const unsigned int tile_max_dim = std::max(output_tile_size.width, output_tile_size.height); // Check optimized kernel if output_dims == 2x2 - if((tile_max_dim == 2) && (input->info()->data_layout() == DataLayout::NCHW)) + if((tile_max_dim == 2) && (data_layout == DataLayout::NCHW)) { _step_z = (_input->info()->dimension(2) % 2) != 0 ? 1 : 2; } @@ -186,7 +191,7 @@ void CLWinogradInputTransformKernel::configure(const ICLTensor *input, ICLTensor // Append stepz and data layout kernel_name += "_stepz"; kernel_name += support::cpp11::to_string(_step_z); - kernel_name += "_" + lower_string(string_from_data_layout(input->info()->data_layout())); + kernel_name += "_" + lower_string(string_from_data_layout(data_layout)); _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); @@ -223,17 +228,30 @@ void CLWinogradInputTransformKernel::run(const Window &window, cl::CommandQueue ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); - const size_t idx_w = get_data_layout_dimension_index(_input->info()->data_layout(), DataLayoutDimension::WIDTH); - const size_t idx_h = get_data_layout_dimension_index(_input->info()->data_layout(), DataLayoutDimension::HEIGHT); - const size_t idx_c = get_data_layout_dimension_index(_input->info()->data_layout(), DataLayoutDimension::CHANNEL); + const DataLayout data_layout = _input->info()->data_layout(); + const size_t idx_w = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); + const size_t idx_h = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); + const size_t idx_c = get_data_layout_dimension_index(data_layout, DataLayoutDimension::CHANNEL); + const size_t total_batches = window.shape().total_size_upper(3); - Window slice = window.first_slice_window_3D(); + // Collapse window + Window window_collapsed = window.collapse_if_possible(ICLKernel::window(), Window::DimZ); + + Window slice = window_collapsed.first_slice_window_3D(); slice.set(idx_w, Window::Dimension(0, _num_tiles_x, 1)); slice.set(idx_h, Window::Dimension(0, _num_tiles_y, 1)); + if(data_layout == DataLayout::NHWC) + { + slice.set(idx_h, Window::Dimension(0, _num_tiles_y * total_batches, 1)); + } ARM_COMPUTE_ERROR_ON(((slice[idx_c].end() - slice[idx_c].start()) % _step_z) != 0); slice.set(idx_c, Window::Dimension(slice[idx_c].start(), slice[idx_c].end(), _step_z)); + unsigned int idx = 2 * num_arguments_per_3D_tensor(); + _kernel.setArg<cl_uint>(idx++, static_cast<unsigned int>(_input->info()->strides_in_bytes()[3])); + _kernel.setArg<cl_uint>(idx++, static_cast<unsigned int>(_output->info()->strides_in_bytes()[3])); + do { unsigned int idx = 0; @@ -242,5 +260,5 @@ void CLWinogradInputTransformKernel::run(const Window &window, cl::CommandQueue enqueue(queue, *this, slice, lws_hint()); } - while(window.slide_window_slice_3D(slice)); + while(window_collapsed.slide_window_slice_3D(slice)); } diff --git a/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp b/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp index 75988c6ca1..dc0a0e7f8f 100644 --- a/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp +++ b/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp @@ -165,6 +165,7 @@ void CLWinogradOutputTransformKernel::configure(const ICLTensor *input, const IC build_opts.add_option("-DOUTPUT_TILE_W=" + support::cpp11::to_string(output_tile_size.width)); build_opts.add_option("-DOUTPUT_TILE_H=" + support::cpp11::to_string(output_tile_size.height)); build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); + build_opts.add_option("-DSRC_DEPTH=" + support::cpp11::to_string(_input->info()->dimension(2))); build_opts.add_option_if(winograd_info.kernel_size.height == 1, "-DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL"); build_opts.add_option_if(winograd_info.kernel_size.width == 1, "-DWINOGRAD_OUTPUT_TRANSFORM_VERTICAL"); @@ -206,8 +207,11 @@ void CLWinogradOutputTransformKernel::run(const Window &window, cl::CommandQueue ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); + // Collapse window + Window window_collapsed = window.collapse_if_possible(ICLKernel::window(), Window::DimZ); + // Get initial windows - Window slice = window.first_slice_window_3D(); + Window slice = window_collapsed.first_slice_window_4D(); slice.set(Window::DimZ, Window::Dimension(0, 1, 1)); // Setup output slice @@ -217,7 +221,7 @@ void CLWinogradOutputTransformKernel::run(const Window &window, cl::CommandQueue if(_bias != nullptr) { - unsigned int idx1 = 2 * num_arguments_per_3D_tensor(); + unsigned int idx1 = 2 * num_arguments_per_4D_tensor(); Window slice_biases; slice_biases.use_tensor_dimensions(_bias->info()->tensor_shape()); add_1D_tensor_argument(idx1, _bias, slice_biases); @@ -225,15 +229,15 @@ void CLWinogradOutputTransformKernel::run(const Window &window, cl::CommandQueue if(_output->info()->data_layout() == DataLayout::NHWC) { - unsigned int idx2 = 2 * num_arguments_per_3D_tensor() + ((_bias != nullptr) ? num_arguments_per_1D_tensor() : 0); + unsigned int idx2 = 2 * num_arguments_per_4D_tensor() + ((_bias != nullptr) ? num_arguments_per_1D_tensor() : 0); _kernel.setArg(idx2, static_cast<int>(_output->info()->total_size() - _output->info()->strides_in_bytes().y())); } do { unsigned int idx = 0; - add_3D_tensor_argument(idx, _input, slice); - add_3D_tensor_argument(idx, _output, slice_out); + add_4D_tensor_argument(idx, _input, slice); + add_4D_tensor_argument(idx, _output, slice_out); enqueue(queue, *this, slice, lws_hint()); } while(window.slide_window_slice_3D(slice) && window.slide_window_slice_3D(slice_out)); |