diff options
| author | Usama Arif <usama.arif@arm.com> | 2019-03-26 13:44:01 +0000 |
|---|---|---|
| committer | Usama Arif <usama.arif@arm.com> | 2019-03-26 15:47:46 +0000 |
| commit | ae0001e4e9af43e03398a8aa36c038cf4fb2600a (patch) | |
| tree | c15543469a3fb30f936d29fbc545de32ecb44bb8 /src/core/CL/cl_kernels/l2_normalize.cl | |
| parent | ec6997563a7cccf58431267cca39435ecd57cd32 (diff) | |
| download | ComputeLibrary-ae0001e4e9af43e03398a8aa36c038cf4fb2600a.tar.gz | |
COMPMID-2077 Optimise CLL2NormalizeLayerKernel
Change-Id: If3866d90f84d62578b7e26f524fa4d7757a5970c
Signed-off-by: Usama Arif <usama.arif@arm.com>
Reviewed-on: https://review.mlplatform.org/c/906
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/l2_normalize.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/l2_normalize.cl | 20 |
1 files changed, 13 insertions, 7 deletions
diff --git a/src/core/CL/cl_kernels/l2_normalize.cl b/src/core/CL/cl_kernels/l2_normalize.cl index 5f66efbcc4..70b8b363fa 100644 --- a/src/core/CL/cl_kernels/l2_normalize.cl +++ b/src/core/CL/cl_kernels/l2_normalize.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -31,26 +31,32 @@ * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/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) + * @param[in] src_step_y src_stride_y * number of elements along X processed per workitem(in bytes) * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor * @param[in] sum_ptr Pointer to the source tensor. Supported data types: F16/F32 * @param[in] sum_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] sum_step_x sum_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] sum_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] sum_step_y sum_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] sum_offset_first_element_in_bytes The offset of the first element in the source tensor * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor * @param[in] epsilon Epsilon value */ __kernel void l2_normalize_x( - VECTOR_DECLARATION(src), - VECTOR_DECLARATION(sum), - VECTOR_DECLARATION(dst), + IMAGE_DECLARATION(src), + IMAGE_DECLARATION(sum), + IMAGE_DECLARATION(dst), DATA_TYPE epsilon) { - Vector src = CONVERT_TO_VECTOR_STRUCT(src); - Vector sum = CONVERT_TO_VECTOR_STRUCT(sum); - Vector dst = CONVERT_TO_VECTOR_STRUCT(dst); + Image src = CONVERT_TO_IMAGE_STRUCT(src); + Image sum = CONVERT_TO_IMAGE_STRUCT(sum); + Image dst = CONVERT_TO_IMAGE_STRUCT(dst); VEC_DATA_TYPE(DATA_TYPE, 16) in = vload16(0, (__global DATA_TYPE *)src.ptr); |