diff options
| author | Michalis Spyrou <michalis.spyrou@arm.com> | 2018-12-07 11:24:36 +0000 |
|---|---|---|
| committer | Michalis Spyrou <michalis.spyrou@arm.com> | 2018-12-07 16:38:23 +0000 |
| commit | 7bb6fb8a81111fcb4804bd30657e5ef73d1449a9 (patch) | |
| tree | 68d456885d22ad015845e333e25036fc50a95e1a /src/core/CL | |
| parent | a84faffd290139be54e4a52ab11da0369262e889 (diff) | |
| download | ComputeLibrary-7bb6fb8a81111fcb4804bd30657e5ef73d1449a9.tar.gz | |
COMPMID-1710 CLElementwiseUnary: Fix accesses smaller than vector size
Change-Id: I26c00ed1941a20fb02fac7d6c03f06adf87fee4d
Reviewed-on: https://review.mlplatform.org/363
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Diffstat (limited to 'src/core/CL')
| -rw-r--r-- | src/core/CL/cl_kernels/elementwise_unary.cl | 34 |
1 files changed, 28 insertions, 6 deletions
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) |