From c55beee7ef70fa08a5d217619083b288a74fcb27 Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Tue, 23 Oct 2018 15:23:23 +0100 Subject: COMPMID-1029: Collapse CLWinogradInputTransform/CLWinogradOutputTransform Change-Id: I051748502ca24b9952e7313524bbfd708162efb4 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/155166 Tested-by: bsgcomp Reviewed-by: Gian Marco Iodice --- src/core/CL/cl_kernels/winograd_input_transform.cl | 601 ++++++++++++--------- 1 file changed, 355 insertions(+), 246 deletions(-) (limited to 'src/core/CL/cl_kernels/winograd_input_transform.cl') 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 -- cgit v1.2.1