COMPMID-1844: Implement CLCrop
Change-Id: I8822c37adc45960705dc3f32a53214795ba3cf39
Signed-off-by: George Wort <george.wort@arm.com>
Reviewed-on: https://review.mlplatform.org/c/789
Reviewed-by: Manuel Bottini <manuel.bottini@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Pablo Marquez <pablo.tello@arm.com>
diff --git a/src/core/CL/cl_kernels/memset.cl b/src/core/CL/cl_kernels/memset.cl
index 80b34eb..7d8e0ef 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