commit 1d1f32ce7ef6acea4afd4cf6a929436640b72ccd
author Gian Marco Iodice <gianmarco.iodice@arm.com> Fri Aug 10 09:34:11 2018 +0100
committer Anthony Barbier <anthony.barbier@arm.com> Fri Nov 02 16:54:54 2018 +0000
tree 16e6d0902b0663464f382bac86b0220f0f7c2a97
parent 9f5a64379ff30a39718ea2bd7e64c6473da866ce

COMPMID-1188 - Passed WIDTH_OFFSET at compile time in CLWidthDepthConcatenateLayerKernel

Change-Id: Icab813cd432174608621ee6a87015aeb10ab822d
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/143570
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>

src/core/CL/cl_kernels/concatenate.cl

diff --git a/src/core/CL/cl_kernels/concatenate.cl b/src/core/CL/cl_kernels/concatenate.cl
index 6ec8383..16c4363 100644
--- a/src/core/CL/cl_kernels/concatenate.cl
+++ b/src/core/CL/cl_kernels/concatenate.cl
@@ -23,9 +23,14 @@
  */
 #include "helpers.h"
 
+#if defined(DATA_TYPE)
+#if defined(WIDTH_OFFSET)
 /** This kernel concatenates the input tensor into the output tensor along the first dimension
  *
- * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QASYMM8, F16, F32
+ * @note The data type has to be passed at compile time using -DDATA_TYPE. i.e. -DDATA_TYPE=float
+ * @note The offset for the first spatial dimension has to be passed at compile time using -DWIDTH_OFFSET. i.e. -DWIDTH_OFFSET=128
+ *
+ * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: U8/S8/QASYMM8/U16/S16/F16/U32/F32
  * @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)
@@ -45,8 +50,7 @@
  */
 __kernel void concatenate_width(
     TENSOR3D_DECLARATION(src),
-    TENSOR3D_DECLARATION(dst),
-    int offset)
+    TENSOR3D_DECLARATION(dst))
 {
     Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
     Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
@@ -55,8 +59,9 @@
     source_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)src.ptr);
 
     VSTORE(VEC_SIZE)
-    (source_values, 0, (__global DATA_TYPE *)(dst.ptr + offset));
+    (source_values, 0, (__global DATA_TYPE *)(dst.ptr) + WIDTH_OFFSET);
 }
+#endif // defined(WIDTH_OFFSET)
 
 /** This kernel concatenates the input tensor into the output tensor along the third dimension
  *
@@ -92,3 +97,4 @@
     VSTORE(VEC_SIZE)
     (source_values, 0, (__global DATA_TYPE *)(dst.ptr + offsets.z));
 }
+#endif // defined(DATA_TYPE)
\ No newline at end of file