diff options
Diffstat (limited to 'src/core/CL/cl_kernels/nhwc/indirect_convolution.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/indirect_convolution.cl | 40 |
1 files changed, 17 insertions, 23 deletions
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) |