diff options
Diffstat (limited to 'src/core/CL/cl_kernels/nhwc')
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/direct_convolution.cl | 42 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl | 24 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl | 21 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/indirect_convolution.cl | 40 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/scale.cl | 64 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/transposed_convolution.cl | 14 |
6 files changed, 104 insertions, 101 deletions
diff --git a/src/core/CL/cl_kernels/nhwc/direct_convolution.cl b/src/core/CL/cl_kernels/nhwc/direct_convolution.cl index 4693a1fbcd..81ceeb8846 100644 --- a/src/core/CL/cl_kernels/nhwc/direct_convolution.cl +++ b/src/core/CL/cl_kernels/nhwc/direct_convolution.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022 Arm Limited. + * Copyright (c) 2021-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -66,36 +66,36 @@ * - The weights offset e.g. -DWEI_OFFSET=4 * - The quantized zero value e.g. -DZERO_VALUE=4 * + * @param[in] src_img (Not supported) Read only cl_image object for the source tensor. Included when SRC_TENSOR_TYPE=IMAGE * @param[in] src_ptr Pointer to the source tensor. Supported data type: F16/F32/QASYMM8 - * @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_c The size of the channels dimension of the source tensor + * @param[in] src_w The size of the width dimension of the source tensor + * @param[in] src_h The size of the height dimension of the source tensor + * @param[in] src_n The size of the batches dimension of the source tensor * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] dst_img (Not supported) Write only cl_image object for the destination tensor. Included when DST_TENSOR_TYPE=IMAGE * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: 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 destination 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 destination tensor in W dimension (in bytes) - * @param[in] dst_step_w dst_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] dst_c The size of the channels dimension of the destination tensor + * @param[in] dst_w The size of the width dimension of the destination tensor + * @param[in] dst_h The size of the height dimension of the destination tensor + * @param[in] dst_n The size of the batches dimension of the destination tensor * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] wei_img (Optional) Read only cl_image object for the weights tensor. Included when WEI_TENSOR_TYPE=IMAGE * @param[in] wei_ptr Pointer to the weights tensor. Supported data type: same as @p src_ptr - * @param[in] wei_stride_x Stride of the weights tensor in X dimension (in bytes) - * @param[in] wei_step_x wei_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] wei_stride_y Stride of the weights tensor in Y dimension (in bytes) - * @param[in] wei_step_y wei_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] wei_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] wei_step_z wei_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] wei_stride_w Stride of the weights tensor in W dimension (in bytes) - * @param[in] wei_step_w wei_stride_w * number of elements along W processed per workitem(in bytes) - * @param[in] wei_offset_first_element_in_bytes The offset of the first element in the bias matrix + * @param[in] wei_c The size of the channels dimension of the weights tensor + * @param[in] wei_w The size of the width dimension of the weights tensor + * @param[in] wei_h The size of the height dimension of the weights tensor + * @param[in] wei_n The size of the batches dimension of the weights tensor + * @param[in] wei_offset_first_element_in_bytes The offset of the first element in the weights matrix * @param[in] bia_ptr (Optional) Pointer to the bias tensor Supported data type: same as @p src_ptr (if F32/F16) or S32 (if QASYMM8/QASYMM8_SIGNED) * @param[in] bia_stride_x (Optional) Stride of the bias tensor in X dimension (in bytes) * @param[in] bia_step_x (Optional) bia_stride_x * number of elements along X processed per workitem(in bytes) @@ -103,9 +103,9 @@ */ //! @endcond __kernel void direct_convolution_nhwc( - TENSOR4D_T(src, SRC_TENSOR_TYPE), - TENSOR4D_T(dst, DST_TENSOR_TYPE), - TENSOR4D_T(wei, WEI_TENSOR_TYPE) + TENSOR4D_RO_T(src, SRC_TENSOR_TYPE), + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE), + TENSOR4D_RO_T(wei, WEI_TENSOR_TYPE) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bia) @@ -292,4 +292,4 @@ __kernel void direct_convolution_nhwc( #undef _IDST_HEIGHT #undef _IDST_CHANNELS #undef _IY_MULTIPLIER -}
\ No newline at end of file +} diff --git a/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl b/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl index 345469063a..dcbae220b6 100644 --- a/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl +++ b/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022 Arm Limited. + * Copyright (c) 2021-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -54,6 +54,7 @@ * @note The number of rows to read from the src tensor must be passed at compile time using -DM0_A (e.g., -DM0_A=3). M0_A must be equal to WEI_WIDTH + (M0 - 1) * @note The number of columns to read from the src tensor must be passed at compile time using -DN0_A. It can either be 1 (for DEPTH_MULTIPLIER > 1) or N0 (for DEPTH_MULTIPLIER == 1) * + * @param[in] src_img (Not supported) Read only cl_image object for the source tensor. Included when SRC_TENSOR_TYPE=IMAGE * @param[in] src_ptr Pointer to the source tensor. Supported data type: F16/F32 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) @@ -63,6 +64,7 @@ * @param[in] src_h The size of the height dimension of the source tensor * @param[in] src_n The size of the batches dimension of the source tensor * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] dst_img (Not supported) Write only cl_image object for the destination tensor. Included when DST_TENSOR_TYPE=IMAGE * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: same as @p src_ptr * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) @@ -72,16 +74,16 @@ * @param[in] dst_h The size of the height dimension of the destination tensor * @param[in] dst_n The size of the batches dimension of the destination tensor * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] wei_img (Optional) Read only cl_image object for the weights tensor. Included when WEI_TENSOR_TYPE=IMAGE * @param[in] wei_ptr Pointer to the weights tensor. Supported data type: same as @p src_ptr - * @param[in] wei_stride_x Stride of the weights tensor in X dimension (in bytes) - * @param[in] wei_step_x wei_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] wei_stride_y Stride of the weights tensor in Y dimension (in bytes) - * @param[in] wei_step_y wei_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] wei_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] wei_step_z wei_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] wei_stride_w Stride of the weights tensor in W dimension (in bytes) - * @param[in] wei_step_w wei_stride_w * number of elements along W processed per workitem(in bytes) - * @param[in] wei_offset_first_element_in_bytes The offset of the first element in the bias matrix + * @param[in] wei_c The size of the channels dimension of the weights tensor + * @param[in] wei_w The size of the width dimension of the weights tensor + * @param[in] wei_h The size of the height dimension of the weights tensor + * @param[in] wei_n The size of the batches dimension of the weights tensor + * @param[in] wei_offset_first_element_in_bytes The offset of the first element in the weigts matrix * @param[in] bia_ptr (Optional) Pointer to the bias tensor Supported data type: same as @p src_ptr * @param[in] bia_stride_x (Optional) Stride of the bias tensor in X dimension (in bytes) * @param[in] bia_step_x (Optional) bia_stride_x * number of elements along X processed per workitem(in bytes) @@ -89,9 +91,9 @@ */ //! @endcond __kernel void dwc_native_fp_nhwc( - TENSOR4D_T(src, SRC_TENSOR_TYPE), - TENSOR4D_T(dst, DST_TENSOR_TYPE), - TENSOR4D(wei, WEI_TENSOR_TYPE) + TENSOR4D_RO_T(src, SRC_TENSOR_TYPE), + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE), + TENSOR4D_RO_T(wei, WEI_TENSOR_TYPE) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bia) @@ -206,4 +208,4 @@ __kernel void dwc_native_fp_nhwc( } #endif // defined(WEI_WIDTH) && defined(WEI_HEIGHT) && defined(N0) && defined(M0) && defined(DILATION_X) && defined(DILATION_Y) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) // *INDENT-ON* -// clang-format on
\ No newline at end of file +// clang-format on diff --git a/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl b/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl index e2ffd444dd..2d255e5b61 100644 --- a/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl +++ b/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022 Arm Limited. + * Copyright (c) 2021-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -77,6 +77,7 @@ * @note The number of rows to read from the src tensor must be passed at compile time using -DM0_A (e.g., -DM0_A=3). M0_A must be equal to WEI_WIDTH + (M0 - 1) * @note The number of columns to read from the src tensor must be passed at compile time using -DN0_A. It can either be 1 (for DEPTH_MULTIPLIER > 1) or N0 (for DEPTH_MULTIPLIER == 1) * + * @param[in] src_img (Not supported) Read only cl_image object for the source tensor. Included when SRC_TENSOR_TYPE=IMAGE * @param[in] src_ptr Pointer to the source tensor. Supported data type: QSYMM8/QASYMM8/QASYMM8_SIGNED/QSYMM8_PER_CHANNEL * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) @@ -86,6 +87,7 @@ * @param[in] src_h The size of the height dimension of the source tensor * @param[in] src_n The size of the batches dimension of the source tensor * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] dst_img (Not supported) Write only cl_image object for the destination tensor. Included when DST_TENSOR_TYPE=IMAGE * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: same as @p src_ptr * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) @@ -95,14 +97,15 @@ * @param[in] dst_h The size of the height dimension of the destination tensor * @param[in] dst_n The size of the batches dimension of the destination tensor * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] wei_img (Not supported) Read only cl_image object for the weights tensor. Included when WEI_TENSOR_TYPE=IMAGE * @param[in] wei_ptr Pointer to the weights tensor. Supported data type: same as @p src_ptr - * @param[in] wei_stride_x Stride of the weights tensor in X dimension (in bytes) - * @param[in] wei_step_x wei_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] wei_stride_y Stride of the weights tensor in Y dimension (in bytes) - * @param[in] wei_step_y wei_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] wei_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] wei_step_z wei_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] wei_stride_w Stride of the weights tensor in W dimension (in bytes) + * @param[in] wei_c The size of the channels dimension of the weights tensor + * @param[in] wei_w The size of the width dimension of the weights tensor + * @param[in] wei_h The size of the height dimension of the weights tensor + * @param[in] wei_n The size of the batches dimension of the weights tensor * @param[in] wei_step_w wei_stride_w * number of elements along W processed per workitem(in bytes) * @param[in] wei_offset_first_element_in_bytes The offset of the first element in the weights tensor * @param[in] dst_multipliers_ptr Pointer to the destination multipliers tensor for the per-channel quantization. Supported data type: S32 @@ -120,9 +123,9 @@ */ //! @endcond __kernel void dwc_native_quantized_nhwc( - TENSOR4D_T(src, SRC_TENSOR_TYPE), - TENSOR4D_T(dst, DST_TENSOR_TYPE), - TENSOR4D(wei, WEI_TENSOR_TYPE), + TENSOR4D_RO_T(src, SRC_TENSOR_TYPE), + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE), + TENSOR4D_RO_T(wei, WEI_TENSOR_TYPE), VECTOR_DECLARATION(dst_multipliers), VECTOR_DECLARATION(dst_shifts) #if defined(HAS_BIAS) @@ -269,4 +272,4 @@ __kernel void dwc_native_quantized_nhwc( } #endif // defined(WEI_WIDTH) && defined(WEI_HEIGHT) && defined(N0) && defined(M0) && defined(DILATION_X) && defined(DILATION_Y) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) // *INDENT-ON* -// clang-format on
\ No newline at end of file +// clang-format on diff --git a/src/core/CL/cl_kernels/nhwc/indirect_convolution.cl b/src/core/CL/cl_kernels/nhwc/indirect_convolution.cl index c88f0034c5..aa719bfef0 100644 --- a/src/core/CL/cl_kernels/nhwc/indirect_convolution.cl +++ b/src/core/CL/cl_kernels/nhwc/indirect_convolution.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -41,7 +41,7 @@ * @note The number of M0 rows (width*height) to process must be passed at compile time using -DM0 (e.g. -DM0=2) * - M0 = 1, 2, 3, 4, 5, 6, 7, and 8 * - * @param[out] dst_img (Not supported) CLImage object to the destination tensor (DST_TENSOR_TYPE=IMAGE only) + * @param[out] dst_img (Not supported) Write only cl_image object for the destination tensor. Included when DST_TENSOR_TYPE=IMAGE * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: INT32 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) @@ -54,7 +54,7 @@ */ //! @endcond __kernel void indirect_convolution_address_precalculation( - TENSOR4D_T(dst, DST_TENSOR_TYPE)) + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE)) { const int x = get_global_id(0); const int y = get_global_id(1); @@ -123,7 +123,7 @@ __kernel void indirect_convolution_address_precalculation( * - N0 = 2, 3, 4, 8, 16 * - K0 = 2, 3, 4, 8, 16 (only 4, 8 and 16 if WEI_TENSOR_TYPE=IMAGE) * - * @param[in] src_img (Not supported) CLImage object to the source tensor (SRC_TENSOR_TYPE=IMAGE only) + * @param[in] src_img (Not supported) Read only cl_image object for the source tensor. Included when SRC_TENSOR_TYPE=IMAGE * @param[in] src_ptr Pointer to the source tensor. Supported data type: F16/F32 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) @@ -133,7 +133,7 @@ __kernel void indirect_convolution_address_precalculation( * @param[in] src_h The size of the height dimension of the source tensor * @param[in] src_n The size of the batches dimension of the source tensor * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[in] off_img (Not supported) CLImage object to the indirect buffer tensor (OFF_TENSOR_TYPE=IMAGE only) + * @param[in] off_img (Not supported) Read only cl_image object for the indirect buffer tensor. Included when OFF_TENSOR_TYPE=IMAGE * @param[in] off_ptr Pointer to the indirect buffer tensor. Supported data type: INT32 * @param[in] off_stride_y Stride of the indirect buffer tensor in Y dimension (in bytes) * @param[in] off_stride_z Stride of the indirect buffer tensor in Z dimension (in bytes) @@ -143,8 +143,8 @@ __kernel void indirect_convolution_address_precalculation( * @param[in] off_h The size of the height dimension of the indirect buffer tensor * @param[in] off_n The size of the batches dimension of the indirect buffer tensor * @param[in] off_offset_first_element_in_bytes The offset of the first element in the indirect buffer tensor - * @param[out] dst_img (Not supported) CLImage object to the destination tensor (DST_TENSOR_TYPE=IMAGE only) - * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: same as the input tensor + * @param[out] dst_img (Not supported) Write only cl_image object for the destination tensor. Included when DST_TENSOR_TYPE=IMAGE + * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: same as @p src_ptr * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) @@ -153,8 +153,8 @@ __kernel void indirect_convolution_address_precalculation( * @param[in] dst_h The size of the height dimension of the destination tensor * @param[in] dst_n The size of the batches dimension of the destination tensor * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] wei_img (Optional) CLImage object to the destination tensor (WEI_TENSOR_TYPE=IMAGE only) - * @param[out] wei_ptr Pointer to the weights tensor. Supported data type: same as the input tensor + * @param[out] wei_img (Optional) Read only cl_image object for the weights tensor. Included when WEI_TENSOR_TYPE=IMAGE + * @param[out] wei_ptr Pointer to the weights tensor. Supported data type: same as @p src_ptr * @param[in] wei_stride_y Stride of the weights tensor in Y dimension (in bytes) * @param[in] wei_stride_z Stride of the weights tensor in Z dimension (in bytes) * @param[in] wei_stride_w Stride of the weights tensor in W dimension (in bytes) @@ -163,23 +163,17 @@ __kernel void indirect_convolution_address_precalculation( * @param[in] wei_h The size of the height dimension of the weights tensor * @param[in] wei_n The size of the batches dimension of the weights tensor * @param[in] wei_offset_first_element_in_bytes The offset of the first element in the weights tensor - * @param[out] bia_img (Not supported) CLImage object to the destination tensor (BIA_TENSOR_TYPE=IMAGE only) - * @param[out] bia_ptr (Optional) Pointer to the bias tensor. Supported data type: same as the input tensor - * @param[in] bia_stride_y (Optional) Stride of the bias tensor in Y dimension (in bytes) - * @param[in] bia_stride_z (Optional) Stride of the bias tensor in Z dimension (in bytes) - * @param[in] bia_stride_w (Optional) Stride of the bias tensor in W dimension (in bytes) - * @param[in] bia_c (Optional) The size of the channels dimension of the bias tensor - * @param[in] bia_w (Optional) The size of the width dimension of the bias tensor - * @param[in] bia_h (Optional) The size of the height dimension of the bias tensor - * @param[in] bia_n (Optional) The size of the batches dimension of the bias tensor - * @param[in] bia_offset_first_element_in_bytes (Optional) The offset of the first element in the bias tensor + * @param[in] bia_ptr (Optional) Pointer to the bias tensor Supported data type: same as @p src_ptr + * @param[in] bia_stride_x (Optional) Stride of the bias tensor in X dimension (in bytes) + * @param[in] bia_step_x (Optional) bia_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] bia_offset_first_element_in_bytes (Optional) The offset of the first element in the bias matrix */ //! @endcond __kernel void indirect_convolution_nhwc( - TENSOR4D_T(src, SRC_TENSOR_TYPE), - TENSOR4D_T(off, OFF_TENSOR_TYPE), - TENSOR4D_T(dst, DST_TENSOR_TYPE), - TENSOR4D_T(wei, WEI_TENSOR_TYPE) + TENSOR4D_RO_T(src, SRC_TENSOR_TYPE), + TENSOR4D_RO_T(off, OFF_TENSOR_TYPE), + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE), + TENSOR4D_RO_T(wei, WEI_TENSOR_TYPE) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bia) diff --git a/src/core/CL/cl_kernels/nhwc/scale.cl b/src/core/CL/cl_kernels/nhwc/scale.cl index f6a3e0971b..e071b0f192 100644 --- a/src/core/CL/cl_kernels/nhwc/scale.cl +++ b/src/core/CL/cl_kernels/nhwc/scale.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2022 Arm Limited. + * Copyright (c) 2016-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -61,8 +61,8 @@ */ //! @endcond __kernel void scale_nearest_neighbour_nhwc( - TENSOR4D_T(src, SRC_TENSOR_TYPE), - TENSOR4D_T(dst, DST_TENSOR_TYPE), + TENSOR4D_RO_T(src, SRC_TENSOR_TYPE), + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE), const float scale_x, const float scale_y) { @@ -128,31 +128,33 @@ __kernel void scale_nearest_neighbour_nhwc( * - The source offset e.g. -DOFFSET=4 * - The source scale e.g. -DSCALE=4 * - * @param[in] src_ptr Pointer to the source tensor. Supported data types: U8/S16/F16/F32. - * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) - * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) - * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) - * @param[in] src_c The size of the channels dimension of the source tensor - * @param[in] src_w The size of the width dimension of the source tensor - * @param[in] src_h The size of the height dimension of the source tensor - * @param[in] src_n The size of the batches dimension of the source tensor - * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[in] dst_ptr Pointer to the destination tensor. Supported data types: U8/S16/F16/F32. - * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) - * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) - * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) - * @param[in] dst_c The size of the channels dimension of the destination tensor - * @param[in] dst_w The size of the width dimension of the destination tensor - * @param[in] dst_h The size of the height dimension of the destination tensor - * @param[in] dst_n The size of the batches dimension of the destination tensor - * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[in] scale_x The scale value to apply on the source width - * @param[in] scale_y The scale value to apply on the source height + * @param[in] src_img (Not supported) Read only cl_image object for the source tensor. Included when SRC_TENSOR_TYPE=IMAGE + * @param[in] src_ptr Pointer to the source tensor. Supported data types: U8/S16/F16/F32. + * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] src_c The size of the channels dimension of the source tensor + * @param[in] src_w The size of the width dimension of the source tensor + * @param[in] src_h The size of the height dimension of the source tensor + * @param[in] src_n The size of the batches dimension of the source tensor + * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] dst_img (Not supported) Write only cl_image object for the destination tensor. Included when DST_TENSOR_TYPE=IMAGE + * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: U8/S16/F16/F32. + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) + * @param[in] dst_c The size of the channels dimension of the destination tensor + * @param[in] dst_w The size of the width dimension of the destination tensor + * @param[in] dst_h The size of the height dimension of the destination tensor + * @param[in] dst_n The size of the batches dimension of the destination tensor + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] scale_x The scale value to apply on the source width + * @param[in] scale_y The scale value to apply on the source height */ //! @endcond __kernel void scale_bilinear_nhwc( - TENSOR4D_T(src, SRC_TENSOR_TYPE), - TENSOR4D_T(dst, DST_TENSOR_TYPE), + TENSOR4D_RO_T(src, SRC_TENSOR_TYPE), + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE), const float scale_x, const float scale_y) { @@ -224,11 +226,11 @@ __kernel void scale_bilinear_nhwc( const float a1 = (yi_f - (float)yi); const float b1 = (1.f - a1); - out[0].v = CONVERT_SAT((CONVERT(in00[0].v, VEC_DATA_TYPE(float, N0)) * b * b1) + - (CONVERT(in01[0].v, VEC_DATA_TYPE(float, N0)) * a * b1) + - (CONVERT(in10[0].v, VEC_DATA_TYPE(float, N0)) * b * a1) + - (CONVERT(in11[0].v, VEC_DATA_TYPE(float, N0)) * a * a1), - VEC_DATA_TYPE(DST_DATA_TYPE, N0)); + out[0].v = CONVERT_SAT((CONVERT(in00[0].v, VEC_DATA_TYPE(float, N0)) * b * b1) + + (CONVERT(in01[0].v, VEC_DATA_TYPE(float, N0)) * a * b1) + + (CONVERT(in10[0].v, VEC_DATA_TYPE(float, N0)) * b * a1) + + (CONVERT(in11[0].v, VEC_DATA_TYPE(float, N0)) * a * a1), + VEC_DATA_TYPE(DST_DATA_TYPE, N0)); #endif // defined(IS_FLOATING_POINT) TILE(uint, 1, 1, dst_indirect_y); @@ -240,4 +242,4 @@ __kernel void scale_bilinear_nhwc( T_STORE_INDIRECT_WIDTH_SELECT(DST_DATA_TYPE, 1, N0, PARTIAL_N0, DST_TENSOR_TYPE, dst, cout, dst_stride_y, x_cond, out, dst_indirect_y); } -#endif /* SCALE_BILINEAR */
\ No newline at end of file +#endif /* SCALE_BILINEAR */ diff --git a/src/core/CL/cl_kernels/nhwc/transposed_convolution.cl b/src/core/CL/cl_kernels/nhwc/transposed_convolution.cl index fe6182fc95..1393537283 100644 --- a/src/core/CL/cl_kernels/nhwc/transposed_convolution.cl +++ b/src/core/CL/cl_kernels/nhwc/transposed_convolution.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -64,7 +64,7 @@ * - The weights offset e.g. -DWEI_OFFSET=4 * - The quantized zero value e.g. -DZERO_VALUE=4 * - * + * @param[in] src_img (Not supported) Read only cl_image object for the source tensor. Included when SRC_TENSOR_TYPE=IMAGE * @param[in] src_ptr Pointer to the source tensor. Supported data type: F16/F32 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) @@ -73,6 +73,7 @@ * @param[in] src_w The size of the width dimension of the source tensor * @param[in] src_h The size of the height dimension of the source tensor * @param[in] src_n The size of the batches dimension of the source tensor + * @param[out] dst_img (Not supported) Write only cl_image object for the destination tensor. Included when DST_TENSOR_TYPE=IMAGE * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: same as @p src_ptr * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) @@ -82,6 +83,7 @@ * @param[in] dst_h The size of the height dimension of the destination tensor * @param[in] dst_n The size of the batches dimension of the destination tensor * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] wei_img (Not supported) Read only cl_image object for the weights tensor. Included when WEI_TENSOR_TYPE=IMAGE * @param[in] wei_ptr Pointer to the weights tensor. Supported data type: same as @p src_ptr * @param[in] wei_stride_y Stride of the weights tensor in Y dimension (in bytes) * @param[in] wei_stride_z Stride of the weights tensor in Z dimension (in bytes) @@ -98,9 +100,9 @@ */ //! @endcond __kernel void transposed_convolution_nhwc( - TENSOR4D_T(src, SRC_TENSOR_TYPE), - TENSOR4D_T(dst, DST_TENSOR_TYPE), - TENSOR4D_T(wei, WEI_TENSOR_TYPE) + TENSOR4D_RO_T(src, SRC_TENSOR_TYPE), + TENSOR4D_WO_T(dst, DST_TENSOR_TYPE), + TENSOR4D_RO_T(wei, WEI_TENSOR_TYPE) #if defined(HAS_BIAS) , VECTOR_DECLARATION(bia) @@ -292,4 +294,4 @@ __kernel void transposed_convolution_nhwc( #undef _IDST_HEIGHT #undef _IDST_CHANNELS #undef _IY_MULTIPLIER -}
\ No newline at end of file +} |