/* * Copyright (c) 2018-2023 Arm Limited. * * SPDX-License-Identifier: MIT * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to * deal in the Software without restriction, including without limitation the * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or * sell copies of the Software, and to permit persons to whom the Software is * furnished to do so, subject to the following conditions: * * The above copyright notice and this permission notice shall be included in all * copies or substantial portions of the Software. * * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ #include "activation_float_helpers.h" #include "helpers.h" #include "tile_helpers.h" #if defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H) #if defined(VEC_SIZE) && VEC_SIZE == 2 #if defined(WINOGRAD_OUTPUT_TRANSFORM_2X2_7X7_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 2x2/2x1 or 1x2, the filter size 7x7/7x1 or 1x7 and the data layout is NHWC * * @note must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16 * @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 output transform 7x1, -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL has to be passed at compile time * @note If this kernel is used to perform Winograd output transform 1x7, -DWINOGRAD_OUTPUT_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. * @note The number of output elements processed along the X direction must be passed at compile time using -DN0 e.g. -DN0=1 * * @param[in] src_ptr Pointer to the source tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] _ISRC_HEIGHT The source tensor's height * @param[in] _IDST_WIDTH The destination tensor's width * @param[in] _IDST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_2x2_7x7_nhwc( TENSOR4D(src, BUFFER), TENSOR4D(dst, BUFFER), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int _ISRC_HEIGHT, const int _IDST_WIDTH, const int _IDST_HEIGHT) { const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM const int mout = GET_SPATIAL_IDX(1, 1, 0); // WINOGRAD OUTPUT TILES #if defined(IS_BATCHED) const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX #else // defined(IS_BATCHED) const int bout = 0; // BATCH SIZE IDX #endif // defined(IS_BATCHED) int x_out = (mout % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (mout / NUM_TILES_X) * OUTPUT_TILE_H; #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) TILE(DATA_TYPE, 8, N0, in); TILE(DATA_TYPE, 2, N0, out); TILE(uint, 8, 1, src_indirect_y); // Calculate the indirect Y for the source tensor LOOP_UNROLLING(int, i, 0, 1, 8, { src_indirect_y[i].v = mout + i *_ISRC_HEIGHT; src_indirect_y[i].v += bout * (int)(_ISRC_HEIGHT * 8); }) // Initialize the input tile LOOP_UNROLLING(int, i, 0, 1, 8, { in[i].v = 0; }) // Load the values across the 8 channels to compose the 8x1 tile T_LOAD_INDIRECT(DATA_TYPE, 8, N0, BUFFER, src, cout, src_stride_y, src_indirect_y, in); // Compute out0 and out01 out[0].v = in[0].v + in[1].v + in[2].v + in[3].v + in[4].v + in[5].v + in[6].v; out[1].v = -in[1].v + in[2].v - (DATA_TYPE)2.f * in[3].v + (DATA_TYPE)2.0f * in[4].v - (DATA_TYPE)3.0f * in[5].v + (DATA_TYPE)3.0f * in[6].v + in[7].v; #if defined(HAS_BIAS) // Add bias TILE(DATA_TYPE, 1, N0, b); T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); T_ELTWISE_BROADCAST_ADD_X(DATA_TYPE, 2, N0, out, b, out); #endif // defined(HAS_BIAS) T_ACTIVATION(DATA_TYPE, 2, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); TILE(uint, 2, 1, dst_indirect_y); #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) LOOP_UNROLLING(int, yk, 0, 1, 2, { int y_c = min(y_out + yk, ((int)_IDST_HEIGHT - 1)); dst_indirect_y[yk].v = x_out + y_c * (int)(_IDST_WIDTH); }) #else // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) LOOP_UNROLLING(int, xk, 0, 1, 2, { int x_c = min(x_out + xk, ((int)_IDST_WIDTH - 1)); dst_indirect_y[xk].v = x_c + y_out * (int)(_IDST_WIDTH); }) #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) // Store the tile in reverse order so the invalid values are overwritten with the valid ones T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 2, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y); #else // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) TILE(DATA_TYPE, 64, N0, in); TILE(DATA_TYPE, 4, N0, out); TILE(DATA_TYPE, 16, N0, tmp); TILE(uint, 64, 1, src_indirect_y); // Calculate the indirect Y for the source tensor LOOP_UNROLLING(int, i, 0, 1, 64, { src_indirect_y[i].v = mout + i *_ISRC_HEIGHT; src_indirect_y[i].v += bout * (int)(_ISRC_HEIGHT * 64); }) // Initialize the input tile LOOP_UNROLLING(int, i, 0, 1, 64, { in[i].v = 0; }) // Load the values across the 64 channels to compose the 8x8 tile T_LOAD_INDIRECT(DATA_TYPE, 64, N0, BUFFER, src, cout, src_stride_y, src_indirect_y, in); LOOP_UNROLLING(int, i, 0, 1, 8, { tmp[i * 2].v = in[0 + i].v + in[8 + i].v + in[16 + i].v + in[24 + i].v + in[32 + i].v + in[40 + i].v + in[48 + i].v; tmp[i * 2 + 1].v = -in[8 + i].v + in[16 + i].v - (DATA_TYPE)2 * in[24 + i].v + (DATA_TYPE)2 * in[32 + i].v + (DATA_TYPE) - 3 * in[40 + i].v + (DATA_TYPE)3 * in[48 + i].v + in[56 + i].v; }) // Compute the 2x2 output tile LOOP_UNROLLING(int, i, 0, 1, 2, { out[i * 2].v = tmp[0 + i].v + tmp[2 + i].v + tmp[4 + i].v + tmp[6 + i].v + tmp[8 + i].v + tmp[10 + i].v + tmp[12 + i].v; out[i * 2 + 1].v = -tmp[2 + i].v + tmp[4 + i].v - (DATA_TYPE)2 * tmp[6 + i].v + (DATA_TYPE)2 * tmp[8 + i].v - (DATA_TYPE)3 * tmp[10 + i].v + (DATA_TYPE)3 * tmp[12 + i].v + tmp[14 + i].v; }) #if defined(HAS_BIAS) // Add bias TILE(DATA_TYPE, 1, N0, b); T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); T_ELTWISE_BROADCAST_ADD_X(DATA_TYPE, 4, N0, out, b, out); #endif // defined(HAS_BIAS) T_ACTIVATION(DATA_TYPE, 4, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); TILE(uint, 4, 1, dst_indirect_y); // Calculate the destination indirect Y LOOP_UNROLLING(int, yk, 0, 1, 2, { LOOP_UNROLLING(int, xk, 0, 1, 2, { int x_c = min(x_out + xk, ((int)_IDST_WIDTH - 1)); int y_c = min(y_out + yk, ((int)_IDST_HEIGHT - 1)); dst_indirect_y[xk + yk * 2].v = x_c + y_c *_IDST_WIDTH; dst_indirect_y[xk + yk * 2].v += bout * (int)(_IDST_WIDTH * _IDST_HEIGHT); }) }) // Store the tile in reverse order so the invalid values are overwritten with the valid ones T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 4, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y); #endif // !defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_2X2_7X7_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC) #endif // defined(VEC_SIZE) && VEC_SIZE == 2 #if defined(VEC_SIZE) && VEC_SIZE == 4 #if defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_3X3_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC) /** This OpenCL kernel performs Winograd output 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 width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4 * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4 * @note If this kernel is used to perform Winograd output transform 3x1, -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL has to be passed at compile time * @note If this kernel is used to perform Winograd output transform 1x3, -DWINOGRAD_OUTPUT_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. * @note The number of output elements processed along the X direction must be passed at compile time using -DN0 e.g. -DN0=1 * * @param[in] src_ptr Pointer to the source tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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 * @param[in] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_4x4_3x3_nhwc( TENSOR4D(src, BUFFER), TENSOR4D(dst, BUFFER), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM const int mout = GET_SPATIAL_IDX(1, 1, 0); // WINOGRAD OUTPUT TILES #if defined(IS_BATCHED) const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX #else // defined(IS_BATCHED) const int bout = 0; // BATCH SIZE IDX #endif // defined(IS_BATCHED) #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) TILE(DATA_TYPE, 6, N0, in); TILE(DATA_TYPE, 4, N0, out); TILE(uint, 6, 1, src_indirect_y); LOOP_UNROLLING(int, i, 0, 1, 6, { src_indirect_y[i].v = mout + i *SRC_HEIGHT; src_indirect_y[i].v += bout * (int)(SRC_HEIGHT * 6); }) // Initialize the input tile LOOP_UNROLLING(int, i, 0, 1, 6, { in[i].v = 0; }) // Load the values across the 36 channels to compose the 6x6 or 6x1 tile T_LOAD_INDIRECT(DATA_TYPE, 6, N0, BUFFER, src, cout, src_stride_y, src_indirect_y, in); // Compute out00, out01, out02 and out03 out[0].v = in[0].v + in[1].v + in[2].v + in[3].v + in[4].v; out[1].v = in[1].v - in[2].v + (DATA_TYPE)2.0f * in[3].v - (DATA_TYPE)2.0f * in[4].v; out[2].v = in[1].v + in[2].v + (DATA_TYPE)4.0f * in[3].v + (DATA_TYPE)4.0f * in[4].v; out[3].v = in[1].v - in[2].v + (DATA_TYPE)8.0f * in[3].v - (DATA_TYPE)8.0f * in[4].v + in[5].v; #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ELTWISE_BROADCAST_ADD_X(DATA_TYPE, 4, N0, out, b, out); #endif // HAS_BIAS int x_out = (mout % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (mout / NUM_TILES_X) * OUTPUT_TILE_H; T_ACTIVATION(DATA_TYPE, 4, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); TILE(uint, 4, 1, dst_indirect_y); // Calculate the destination indirect Y #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) LOOP_UNROLLING(int, yk, 0, 1, 4, { int y_c = min(y_out + yk, ((int)DST_HEIGHT - 1)); dst_indirect_y[yk].v = x_out + y_c *DST_WIDTH; dst_indirect_y[yk].v += bout * (int)(DST_WIDTH * DST_HEIGHT); }) #else // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) LOOP_UNROLLING(int, xk, 0, 1, 4, { int x_c = min(x_out + xk, ((int)DST_WIDTH - 1)); dst_indirect_y[xk].v = x_c + y_out *DST_WIDTH; dst_indirect_y[xk].v += bout * (int)(DST_WIDTH * DST_HEIGHT); }) #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) // Store the tile in reverse order so the invalid values are overwritten with the valid ones T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 4, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y); #else // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) // Calculate the indirect Y for the source tensor TILE(DATA_TYPE, 36, N0, in); TILE(DATA_TYPE, 4, N0, tmp); TILE(uint, 36, 1, src_indirect_y); LOOP_UNROLLING(int, i, 0, 1, 36, { src_indirect_y[i].v = mout + i *SRC_HEIGHT; src_indirect_y[i].v += bout * (int)(SRC_HEIGHT * 36); }) // Initialize the input tile LOOP_UNROLLING(int, i, 0, 1, 36, { in[i].v = 0; }) // Load the values across the 36 channels to compose the 6x6 or 6x1 tile T_LOAD_INDIRECT(DATA_TYPE, 36, N0, BUFFER, src, cout, src_stride_y, src_indirect_y, in); LOOP_UNROLLING(int, i, 0, 1, 6, { tmp[0].v = in[6 + i].v + in[12 + i].v; tmp[1].v = in[6 + i].v - in[12 + i].v; tmp[2].v = in[18 + i].v + in[24 + i].v; tmp[3].v = in[18 + i].v - in[24 + i].v; tmp[3].v = tmp[3].v + tmp[3].v; in[i].v = in[i].v + tmp[0].v + tmp[2].v; in[6 + i].v = tmp[3].v + tmp[1].v; in[12 + i].v = fma(tmp[2].v, (VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[0].v); in[18 + i].v = fma(tmp[3].v, (VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[1].v) + in[30 + i].v; }) // Compute the output tile TILE(DATA_TYPE, 16, N0, out); LOOP_UNROLLING(int, i, 0, 1, 4, { tmp[0].v = in[6 * i + 1].v + in[6 * i + 2].v; tmp[1].v = in[6 * i + 1].v - in[6 * i + 2].v; tmp[2].v = in[6 * i + 3].v + in[6 * i + 4].v; tmp[3].v = in[6 * i + 3].v - in[6 * i + 4].v; tmp[3].v = tmp[3].v + tmp[3].v; out[4 * i + 0].v = in[6 * i + 0].v + tmp[0].v + tmp[2].v; out[4 * i + 1].v = tmp[3].v + tmp[1].v; out[4 * i + 2].v = fma(tmp[2].v, (VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[0].v); out[4 * i + 3].v = fma(tmp[3].v, (VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[1].v) + in[6 * i + 5].v; }) #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ELTWISE_BROADCAST_ADD_X(DATA_TYPE, 16, N0, out, b, out); #endif // HAS_BIAS int x_out = (mout % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (mout / NUM_TILES_X) * OUTPUT_TILE_H; T_ACTIVATION(DATA_TYPE, 16, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); TILE(uint, 16, 1, dst_indirect_y); // Calculate the destination indirect Y LOOP_UNROLLING(int, yk, 0, 1, 4, { LOOP_UNROLLING(int, xk, 0, 1, 4, { int x_c = min(x_out + xk, ((int)DST_WIDTH - 1)); int y_c = min(y_out + yk, ((int)DST_HEIGHT - 1)); dst_indirect_y[xk + yk * 4].v = x_c + y_c *DST_WIDTH; dst_indirect_y[xk + yk * 4].v += bout * (int)(DST_WIDTH * DST_HEIGHT); }) }) // Store the tile in reverse order so the invalid values are overwritten with the valid ones T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 16, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y); #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_3X3_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC) #if defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_5X5_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 4x4/4x1 or 1x4, the filter size 5x5/5x1 or 1x5 and the data layout is NHWC * * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4 * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4 * @note If this kernel is used to perform Winograd output transform 5x1, -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL has to be passed at compile time * @note If this kernel is used to perform Winograd output transform 1x5, -DWINOGRAD_OUTPUT_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. * @note The number of output elements processed along the X direction must be passed at compile time using -DN0 e.g. -DN0=1 * * @param[in] src_ptr Pointer to the source tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_4x4_5x5_nhwc( TENSOR4D(src, BUFFER), TENSOR4D(dst, BUFFER), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM const int mout = GET_SPATIAL_IDX(1, 1, 0); // WINOGRAD OUTPUT TILES #if defined(IS_BATCHED) const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX #else // defined(IS_BATCHED) const int bout = 0; // BATCH SIZE IDX #endif // defined(IS_BATCHED) #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) TILE(DATA_TYPE, 8, N0, in); TILE(DATA_TYPE, 4, N0, out); TILE(DATA_TYPE, 4, N0, tmp); TILE(uint, 8, 1, src_indirect_y); LOOP_UNROLLING(int, i, 0, 1, 8, { src_indirect_y[i].v = mout + i *SRC_HEIGHT; src_indirect_y[i].v += bout * (int)(SRC_HEIGHT * 8); }) // Initialize the input tile LOOP_UNROLLING(int, i, 0, 1, 8, { in[i].v = 0; }) // "in" contains 1x8 or 8x1 tile here T_LOAD_INDIRECT(DATA_TYPE, 8, N0, BUFFER, src, cout, src_stride_y, src_indirect_y, in); // A^T * in, and in this degenerate case out consists of 1 column/row tmp[0].v = in[1].v - in[2].v; tmp[1].v = (DATA_TYPE)2.0f * (in[3].v - in[4].v); tmp[2].v = (DATA_TYPE)2.0f * (in[5].v + in[6].v); tmp[3].v = in[3].v + in[4].v; out[0].v = in[0].v + in[1].v + in[2].v + tmp[3].v + (DATA_TYPE)4.0f * tmp[2].v; out[1].v = tmp[0].v + tmp[1].v + (DATA_TYPE)4.0f * (in[5].v - in[6].v); out[2].v = in[1].v + in[2].v + (DATA_TYPE)4.0f * tmp[3].v + tmp[2].v; out[3].v = tmp[0].v + (DATA_TYPE)4.0f * tmp[1].v + in[5].v - in[6].v + in[7].v; #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ELTWISE_BROADCAST_ADD_X(DATA_TYPE, 4, N0, out, b, out); #endif // HAS_BIAS int x_out = (mout % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (mout / NUM_TILES_X) * OUTPUT_TILE_H; T_ACTIVATION(DATA_TYPE, 4, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); TILE(uint, 4, 1, dst_indirect_y); // Calculate the destination indirect Y #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) LOOP_UNROLLING(int, yk, 0, 1, 4, { int y_c = min(y_out + yk, ((int)DST_HEIGHT - 1)); dst_indirect_y[yk].v = x_out + y_c *DST_WIDTH; dst_indirect_y[yk].v += bout * (int)(DST_WIDTH * DST_HEIGHT); }) #else // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) LOOP_UNROLLING(int, xk, 0, 1, 4, { int x_c = min(x_out + xk, ((int)DST_WIDTH - 1)); dst_indirect_y[xk].v = x_c + y_out *DST_WIDTH; dst_indirect_y[xk].v += bout * (int)(DST_WIDTH * DST_HEIGHT); }) #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) // Store the tile in reverse order so the invalid values are overwritten with the valid ones T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 4, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y); #else // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) // Calculate the indirect Y for the source tensor TILE(DATA_TYPE, 64, N0, in); TILE(DATA_TYPE, 6, N0, tmp); TILE(uint, 64, 1, src_indirect_y); LOOP_UNROLLING(int, i, 0, 1, 64, { src_indirect_y[i].v = mout + i *SRC_HEIGHT; src_indirect_y[i].v += bout * (int)(SRC_HEIGHT * 64); }) // Initialize the input tile LOOP_UNROLLING(int, i, 0, 1, 64, { in[i].v = 0; }) // "in" here is 8x8 tile T_LOAD_INDIRECT(DATA_TYPE, 64, N0, BUFFER, src, cout, src_stride_y, src_indirect_y, in); // A^T * in LOOP_UNROLLING(int, i, 0, 1, 8, { tmp[0].v = in[8 + i].v + in[16 + i].v; tmp[1].v = in[8 + i].v - in[16 + i].v; tmp[2].v = in[24 + i].v + in[32 + i].v; tmp[3].v = in[24 + i].v - in[32 + i].v; tmp[3].v = tmp[3].v + tmp[3].v; tmp[4].v = in[40 + i].v + in[48 + i].v; tmp[4].v = tmp[4].v + tmp[4].v; tmp[5].v = in[40 + i].v - in[48 + i].v; // 4x8 matrix as a result in[i].v = in[i].v + tmp[0].v + fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[4].v, tmp[2].v); in[8 + i].v = tmp[1].v + fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[5].v, tmp[3].v); in[16 + i].v = tmp[0].v + fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[2].v, tmp[4].v); in[24 + i].v = tmp[1].v + fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[3].v, tmp[5].v) + in[56 + i].v; }) // Compute the output tile TILE(DATA_TYPE, 16, N0, out); // in * A, with in = A^T * in as above LOOP_UNROLLING(int, i, 0, 1, 4, { tmp[0].v = in[8 * i + 1].v + in[8 * i + 2].v; tmp[1].v = in[8 * i + 1].v - in[8 * i + 2].v; tmp[2].v = in[8 * i + 3].v + in[8 * i + 4].v; tmp[3].v = in[8 * i + 3].v - in[8 * i + 4].v; tmp[3].v = tmp[3].v + tmp[3].v; tmp[4].v = in[8 * i + 5].v + in[8 * i + 6].v; tmp[4].v = tmp[4].v + tmp[4].v; tmp[5].v = in[8 * i + 5].v - in[8 * i + 6].v; // 4x4 tile out[4 * i].v = in[8 * i].v + tmp[0].v + fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[4].v, tmp[2].v); out[4 * i + 1].v = tmp[1].v + fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[5].v, tmp[3].v); out[4 * i + 2].v = fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[2].v, tmp[0].v) + tmp[4].v; out[4 * i + 3].v = fma((VEC_DATA_TYPE(DATA_TYPE, N0))4.0f, tmp[3].v, tmp[1].v) + tmp[5].v + in[8 * i + 7].v; }) #if defined(HAS_BIAS) TILE(DATA_TYPE, 1, N0, b); T_LOAD(DATA_TYPE, 1, N0, BUFFER, bias, cout, 0, 1, 0, b); // c = c + bias[broadcasted] T_ELTWISE_BROADCAST_ADD_X(DATA_TYPE, 16, N0, out, b, out); #endif // HAS_BIAS int x_out = (mout % NUM_TILES_X) * OUTPUT_TILE_W; int y_out = (mout / NUM_TILES_X) * OUTPUT_TILE_H; T_ACTIVATION(DATA_TYPE, 16, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); TILE(uint, 16, 1, dst_indirect_y); // Calculate the destination indirect Y LOOP_UNROLLING(int, yk, 0, 1, 4, { LOOP_UNROLLING(int, xk, 0, 1, 4, { int x_c = min(x_out + xk, ((int)DST_WIDTH - 1)); int y_c = min(y_out + yk, ((int)DST_HEIGHT - 1)); dst_indirect_y[xk + yk * 4].v = x_c + y_c *DST_WIDTH; dst_indirect_y[xk + yk * 4].v += bout * (int)(DST_WIDTH * DST_HEIGHT); }) }) // Store the tile in reverse order so the invalid values are overwritten with the valid ones T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 16, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y); #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_5X5_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC) #endif // defined(VEC_SIZE) && VEC_SIZE == 4 #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) #if defined(VEC_SIZE) && VEC_SIZE == 2 #if defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 2x1, the filter size 7x1 and the data layout is NHWC * * @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=1 * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24 * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32 * @note -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL 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 tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_2x1_7x1_nhwc( TENSOR4D_DECLARATION(src), TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { winograd_output_transform_2x2_7x7_nhwc(src_ptr, src_stride_x, src_step_x, src_stride_y, 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, dst_step_x, dst_stride_y, 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, bias_stride_x, bias_step_x, bias_offset_first_element_in_bytes, #endif // defined(HAS_BIAS) dst_size, SRC_HEIGHT, DST_WIDTH, DST_HEIGHT); } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC) #endif // defined(VEC_SIZE) && VEC_SIZE == 2 #if defined(VEC_SIZE) && VEC_SIZE == 4 #if defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 4x1, the filter size 3x1 and the data layout is NHWC * * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4 * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=1 * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24 * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32 * @note -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL 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 tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_4x1_3x1_nhwc( TENSOR4D_DECLARATION(src), TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { winograd_output_transform_4x4_3x3_nhwc(src_ptr, src_stride_x, src_step_x, src_stride_y, 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, dst_step_x, dst_stride_y, 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, bias_stride_x, bias_step_x, bias_offset_first_element_in_bytes, #endif // defined(HAS_BIAS) dst_size, SRC_HEIGHT, DST_WIDTH, DST_HEIGHT); } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC) #if defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 4x1, the filter size 5x1 and the data layout is NHWC * * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4 * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=1 * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24 * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32 * @note -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL 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 tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_4x1_5x1_nhwc( TENSOR4D_DECLARATION(src), TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { winograd_output_transform_4x4_5x5_nhwc(src_ptr, src_stride_x, src_step_x, src_stride_y, 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, dst_step_x, dst_stride_y, 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, bias_stride_x, bias_step_x, bias_offset_first_element_in_bytes, #endif // defined(HAS_BIAS) dst_size, SRC_HEIGHT, DST_WIDTH, DST_HEIGHT); } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC) #endif // defined(VEC_SIZE) && VEC_SIZE == 4 #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) #if defined(VEC_SIZE) && VEC_SIZE == 2 #if defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 1x2, the filter size 1x7 and the data layout is NHWC * * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=1 * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=2 * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24 * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32 * @note -DWINOGRAD_OUTPUT_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 tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_1x2_1x7_nhwc( TENSOR4D_DECLARATION(src), TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { winograd_output_transform_2x2_7x7_nhwc(src_ptr, src_stride_x, src_step_x, src_stride_y, 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, dst_step_x, dst_stride_y, 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, bias_stride_x, bias_step_x, bias_offset_first_element_in_bytes, #endif // defined(HAS_BIAS) dst_size, SRC_HEIGHT, DST_WIDTH, DST_HEIGHT); } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC) #endif // defined(VEC_SIZE) && VEC_SIZE == 2 #if defined(VEC_SIZE) && VEC_SIZE == 4 #if defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 1x4, the filter size 1x3 and the data layout is NHWC * * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=1 * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4 * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24 * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32 * @note -DWINOGRAD_OUTPUT_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 tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_1x4_1x3_nhwc( TENSOR4D_DECLARATION(src), TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { winograd_output_transform_4x4_3x3_nhwc(src_ptr, src_stride_x, src_step_x, src_stride_y, 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, dst_step_x, dst_stride_y, 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, bias_stride_x, bias_step_x, bias_offset_first_element_in_bytes, #endif // defined(HAS_BIAS) dst_size, SRC_HEIGHT, DST_WIDTH, DST_HEIGHT); } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC) #if defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC) /** This OpenCL kernel performs Winograd output transform when the output tile is 1x4, the filter size 1x5 and the data layout is NHWC * * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=1 * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4 * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24 * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32 * @note -DWINOGRAD_OUTPUT_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 tensor. Supported data types: F32/F16 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_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] 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] SRC_HEIGHT The source tensor's height * @param[in] DST_WIDTH The destination tensor's width * @param[in] DST_HEIGHT The destination tensor's height */ __kernel void winograd_output_transform_1x4_1x5_nhwc( TENSOR4D_DECLARATION(src), TENSOR4D_DECLARATION(dst), #if defined(HAS_BIAS) VECTOR_DECLARATION(bias), #endif // defined(HAS_BIAS) int dst_size, const int SRC_HEIGHT, const int DST_WIDTH, const int DST_HEIGHT) { winograd_output_transform_4x4_5x5_nhwc(src_ptr, src_stride_x, src_step_x, src_stride_y, 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, dst_step_x, dst_stride_y, 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, bias_stride_x, bias_step_x, bias_offset_first_element_in_bytes, #endif // defined(HAS_BIAS) dst_size, SRC_HEIGHT, DST_WIDTH, DST_HEIGHT); } #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC) #endif // defined(VEC_SIZE) && VEC_SIZE == 4 #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) #endif // defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)