Extend cl image support to input and output tensors

- Add support for texture image to input and output of direct
convolution
- Extend T_LOAD2D_INDIRECT macro to read values from cl image storages

Resolves COMPMID-5715

Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Change-Id: Idb0410f53f6d0763cd9e39895a7cbf9bc826d33a
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/8904
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>

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)