From 1d1f32ce7ef6acea4afd4cf6a929436640b72ccd Mon Sep 17 00:00:00 2001 From: Gian Marco Iodice Date: Fri, 10 Aug 2018 09:34:11 +0100 Subject: 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 Reviewed-by: Anthony Barbier --- src/core/CL/cl_kernels/concatenate.cl | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) (limited to 'src/core/CL/cl_kernels') diff --git a/src/core/CL/cl_kernels/concatenate.cl b/src/core/CL/cl_kernels/concatenate.cl index 6ec8383c52..16c4363899 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 @@ __kernel void concatenate_width( 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 @@ __kernel void concatenate_depth( VSTORE(VEC_SIZE) (source_values, 0, (__global DATA_TYPE *)(dst.ptr + offsets.z)); } +#endif // defined(DATA_TYPE) \ No newline at end of file -- cgit v1.2.1