Changes in cl kernel files related to #if defined conditions (#4106)

1. Add missing #if defined conditions
2. Remove the Vec_size from #if defined check since kernel
   should execute even if it is not passed as an preprocessor
   argument with an assumption VEC_SIZE=1 i.e., processing
   on one element at a time.

Files modified are:
1. embedding_lookup.cl
2. exp.cl
3. hashtable_lookup.cl
4. neg_tensor.cl

Signed-off-by: shubham <shub98.gupta@samsung.com>

diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/embedding_lookup.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/embedding_lookup.cl
index 9105e16..348458f 100644 (file)
--- a/libs/ARMComputeEx/src/core/CL/cl_kernels/embedding_lookup.cl
+++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/embedding_lookup.cl
@@ -16,7 +16,11 @@
  */
 #include "helpers.h"
 
-#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_OUT) && defined(NUM_DIMS)
+#ifndef VEC_SIZE
+#define VEC_SIZE 1
+#endif
+
+#if defined(DATA_TYPE) && defined(DEPTH_OUT) && defined(NUM_DIMS)
 /** Perform embedding_lookup of input tensor
  *
  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
@@ -77,4 +81,4 @@ __kernel void embedding_lookup(TENSOR4D_DECLARATION(input),
     VSTORE(VEC_SIZE)(CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr), VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)),
                      0, (__global DATA_TYPE *)out.ptr);
 }
-#endif // defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_OUT) && defined(NUM_DIMS)
+#endif // defined(DATA_TYPE) && defined(DEPTH_OUT) && defined(NUM_DIMS)

diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/exp.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/exp.cl
index 6ca24c6..69d94f3 100644 (file)
--- a/libs/ARMComputeEx/src/core/CL/cl_kernels/exp.cl
+++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/exp.cl
@@ -16,6 +16,11 @@
  */
 #include "helpers.h"
 
+#ifndef VEC_SIZE
+#define VEC_SIZE 1
+#endif
+
+#if defined(DATA_TYPE)
 /** Perform an exponential operation on an input tensor.
  *
  * @attention Data type can be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=float
@@ -49,3 +54,4 @@ __kernel void exp_layer(
     VSTORE(VEC_SIZE)
     (exp(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input.ptr)), 0, (__global DATA_TYPE *)output.ptr);
 }
+#endif // defined(DATA_TYPE)

diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/hashtable_lookup.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/hashtable_lookup.cl
index b60a6cc..ed74098 100644 (file)
--- a/libs/ARMComputeEx/src/core/CL/cl_kernels/hashtable_lookup.cl
+++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/hashtable_lookup.cl
@@ -16,7 +16,11 @@
  */
 #include "helpers.h"
 
-#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_OUT) && defined(NUM_DIMS)
+#ifndef VEC_SIZE
+#define VEC_SIZE 1
+#endif
+
+#if defined(DATA_TYPE) && defined(DEPTH_OUT) && defined(NUM_DIMS)
 /** Perform hashtable_lookup of input tensor
  *
  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
@@ -81,4 +85,4 @@ __kernel void hashtable_lookup(TENSOR4D_DECLARATION(input),
     VSTORE(VEC_SIZE)(CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr), VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)),
                      0, (__global DATA_TYPE *)out.ptr);
 }
-#endif // defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_OUT) && defined(NUM_DIMS)
+#endif // defined(DATA_TYPE) && defined(DEPTH_OUT) && defined(NUM_DIMS)

diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/neg_tensor.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/neg_tensor.cl
index 43b48f2..e3aa463 100644 (file)
--- a/libs/ARMComputeEx/src/core/CL/cl_kernels/neg_tensor.cl
+++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/neg_tensor.cl
@@ -16,6 +16,11 @@
  */
 #include "helpers.h"
 
+#ifndef VEC_SIZE
+#define VEC_SIZE 1
+#endif
+
+#if defined(DATA_TYPE)
 /** Performs a negation of input tensor.
  *
  * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
@@ -40,3 +45,4 @@ __kernel void neg_tensor(
     VSTORE(VEC_SIZE)
     (-VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input.ptr), 0, (__global DATA_TYPE *)output.ptr);
 }
+#endif // defined(DATA_TYPE)