diff options
| author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2018-08-10 09:34:11 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:54:54 +0000 |
| commit | 1d1f32ce7ef6acea4afd4cf6a929436640b72ccd (patch) | |
| tree | 16e6d0902b0663464f382bac86b0220f0f7c2a97 /src/core/CL/cl_kernels/concatenate.cl | |
| parent | 9f5a64379ff30a39718ea2bd7e64c6473da866ce (diff) | |
| download | ComputeLibrary-1d1f32ce7ef6acea4afd4cf6a929436640b72ccd.tar.gz | |
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>
Diffstat (limited to 'src/core/CL/cl_kernels/concatenate.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/concatenate.cl | 14 |
1 files changed, 10 insertions, 4 deletions
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 |