From 7bb6fb8a81111fcb4804bd30657e5ef73d1449a9 Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Fri, 7 Dec 2018 11:24:36 +0000 Subject: COMPMID-1710 CLElementwiseUnary: Fix accesses smaller than vector size Change-Id: I26c00ed1941a20fb02fac7d6c03f06adf87fee4d Reviewed-on: https://review.mlplatform.org/363 Tested-by: Arm Jenkins Reviewed-by: Michele Di Giorgio --- src/core/CL/cl_kernels/elementwise_unary.cl | 34 ++++++++++++++++++++++++----- 1 file changed, 28 insertions(+), 6 deletions(-) (limited to 'src') diff --git a/src/core/CL/cl_kernels/elementwise_unary.cl b/src/core/CL/cl_kernels/elementwise_unary.cl index bccb47ed1f..92db569d97 100644 --- a/src/core/CL/cl_kernels/elementwise_unary.cl +++ b/src/core/CL/cl_kernels/elementwise_unary.cl @@ -24,7 +24,9 @@ #include "helpers.h" #include "warp_helpers.h" -#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(OPERATION) +#if defined(DATA_TYPE) && defined(OPERATION) + +#if defined(VEC_SIZE) && defined(LAST_ACCESSED_X) /** Calculate reverse square root * * @param[in] input Pointer to the first element. @@ -46,6 +48,29 @@ inline VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) exponential(const VEC_DATA_TYPE(DATA_T { return exp(input); } +#else // !defined(VEC_SIZE) || !defined(LAST_ACCESSED_X) +/** Calculate reverse square root + * + * @param[in] input Single element. + * + * @return reverse square root + */ +inline DATA_TYPE inverse_sqrt(const DATA_TYPE input) +{ + return rsqrt(input); +} + +/** Calculate exponential + * + * @param[in] input Single element. + * + * @return exponential + */ +inline DATA_TYPE exponential(const DATA_TYPE input) +{ + return exp(input); +} +#endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) /** Applies element wise unary operator in a tensor. * @@ -78,10 +103,7 @@ __kernel void elementwise_unary( VSTORE(VEC_SIZE) (OPERATION(data), 0, (__global DATA_TYPE *)out.ptr); #else // !defined(VEC_SIZE) || !defined(LAST_ACCESSED_X) - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr); - VSTORE(VEC_SIZE) - (OPERATION(data), 0, (__global DATA_TYPE *)out.ptr); + *((__global DATA_TYPE *)(out.ptr)) = (DATA_TYPE)(OPERATION(*((__global DATA_TYPE *)in.ptr))); #endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) } -#endif // defined(DATA_TYPE) && defined(VEC_SIZE) && defined(OPERATION) +#endif // defined(DATA_TYPE) && defined(OPERATION) -- cgit v1.2.1