From 894066de8cc26d1a3aca62dcaa6b30a2a1116028 Mon Sep 17 00:00:00 2001 From: George Wort Date: Fri, 15 Feb 2019 15:12:52 +0000 Subject: COMPMID-1844: Implement CLCrop Change-Id: I8822c37adc45960705dc3f32a53214795ba3cf39 Signed-off-by: George Wort Reviewed-on: https://review.mlplatform.org/c/789 Reviewed-by: Manuel Bottini Tested-by: Arm Jenkins Reviewed-by: Pablo Marquez --- src/core/CL/cl_kernels/memset.cl | 15 +++++++++------ 1 file changed, 9 insertions(+), 6 deletions(-) (limited to 'src/core/CL/cl_kernels/memset.cl') diff --git a/src/core/CL/cl_kernels/memset.cl b/src/core/CL/cl_kernels/memset.cl index 80b34ebdf4..7d8e0ef53f 100644 --- a/src/core/CL/cl_kernels/memset.cl +++ b/src/core/CL/cl_kernels/memset.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018 ARM Limited. + * Copyright (c) 2018-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -41,24 +41,27 @@ * @param[in] value The value used to fill the pages of the tensor */ __kernel void memset( - IMAGE_DECLARATION(tensor)) + TENSOR3D_DECLARATION(tensor)) { - Image tensor = CONVERT_TO_IMAGE_STRUCT(tensor); + Tensor3D tensor = CONVERT_TO_TENSOR3D_STRUCT(tensor); -#if defined(VEC_SIZE) && defined(LAST_ACCESSED_X) +#if defined(VEC_SIZE) + +#if defined(LAST_ACCESSED_X) // Check if access on width gets out of bounds // If it does shift access vector to access elements within bounds const int xi = (int)(get_global_id(0) * VEC_SIZE); tensor.ptr -= max(xi - (int)LAST_ACCESSED_X, 0) * tensor_stride_x; +#endif // defined(LAST_ACCESSED_X) VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) data = (DATA_TYPE)(CONSTANT_VALUE); VSTORE(VEC_SIZE) (data, 0, (__global DATA_TYPE *)tensor.ptr); -#else // !defined(VEC_SIZE) || !defined(LAST_ACCESSED_X) +#else // !defined(VEC_SIZE) *((__global DATA_TYPE *)(tensor.ptr)) = (DATA_TYPE)(CONSTANT_VALUE); -#endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) +#endif // defined(VEC_SIZE) } #endif // Check for compile time constants -- cgit v1.2.1