Changes in Comparison op files (#4080)

1. Add missing comments for op_code
2. Add #if definition to verify the preprocessors before executing
3. initialise op_code = 0 and changing the op_code values assigned
   to operation to have runtime error handling
4. Indentation of the kernel code.

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

diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op.cl
index 1664e68..8dbaba9 100644 (file)
--- a/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op.cl
+++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op.cl
@@ -16,12 +16,15 @@
  */
 #include "helpers.h"
 
+#if defined(DATA_TYPE_IN) && defined(DATA_TYPE_OUT) && defined(OP_CODE)
 /** Returns truth value of comparison operators.
  * Comparison operators may be equal, not_equal etc.
  *
  * @attention The input and output data types need to be passed at compile time using -DDATA_TYPE_IN, -DDATA_TYPE_OUT,
  * e.g. -DDATA_TYPE_IN=float, -DDATA_TYPE_OUT = uchar
  * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
+ * @attention Operation type(code) specifying which operation to perform should be passed as preprocessor argument using
+ * -DOP_CODE = number. e.g. -DOP_CODE=1
  *
  * @param[in]  input1_ptr                            Pointer to the source tensor. Supported data types: U8/S8/U16/S16/F16/U32/S32/F32
  * @param[in]  input1_stride_x                       Stride of the source tensor in X dimension (in bytes)
@@ -59,15 +62,21 @@ __kernel void comparison_op(
     Tensor3D input2 = CONVERT_TO_TENSOR3D_STRUCT(input2);
     Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
 
-#if OP_CODE == 0 //EQUAL
-    VSTORE(VEC_SIZE)
-    (CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input1.ptr) == VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input2.ptr),
-             VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)),
-     0, (__global DATA_TYPE_OUT *)output.ptr);
-#elif OP_CODE == 1 //NOT_EQUAL
-    VSTORE(VEC_SIZE)
-    (CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input1.ptr) != VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input2.ptr),
-             VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)),
-     0, (__global DATA_TYPE_OUT *)output.ptr);
-#endif
+    #if OP_CODE == 1 //EQUAL
+     VSTORE(VEC_SIZE)
+     (CONVERT(VLOAD(VEC_SIZE)
+             (0, (__global DATA_TYPE_IN *)input1.ptr) == VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input2.ptr),
+             VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)),0, (__global DATA_TYPE_OUT *)output.ptr);
+
+    #elif OP_CODE == 2 //NOT_EQUAL
+     VSTORE(VEC_SIZE)
+     (CONVERT(VLOAD(VEC_SIZE)
+             (0, (__global DATA_TYPE_IN *)input1.ptr) != VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)input2.ptr),
+             VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)), 0, (__global DATA_TYPE_OUT *)output.ptr);
+
+    #else // OP NOT SUPPORTED
+     return;
+
+    #endif
 }
+#endif // defined(DATA_TYPE_IN) && defined(DATA_TYPE_OUT) && defined(OP_CODE)

diff --git a/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op_quantized.cl b/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op_quantized.cl
index 41c90b7..697400e 100644 (file)
--- a/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op_quantized.cl
+++ b/libs/ARMComputeEx/src/core/CL/cl_kernels/comparison_op_quantized.cl
@@ -17,7 +17,10 @@
 #include "helpers.h"
 #define SUB(x, y) (x) - (y)
 
-#if defined(OFFSET_IN1) && defined(OFFSET_IN2) && defined(SCALE_IN1) && defined(SCALE_IN2) && defined(VEC_SIZE)
+#ifndef VEC_SIZE
+#define VEC_SIZE  1
+
+#if defined(OFFSET_IN1) && defined(OFFSET_IN2) && defined(SCALE_IN1) && defined(SCALE_IN2) && defined(DATA_TYPE_OUT)
 
 #define VEC_FLOAT VEC_DATA_TYPE(float, VEC_SIZE)
 #define VEC_INT VEC_DATA_TYPE(int, VEC_SIZE)
@@ -26,6 +29,8 @@
 /** Returns the truth value of comparison .
  * @attention Offset and Scale of both input should be given as a preprocessor argument using -DOFFSET_IN1=int, -DOFFSET_IN2=int, -DSCALE_IN1=float and -DSCALE_IN2=float. e.g. -DOFFSET_IN1=1, -DOFFSET_IN2=0, -DSCALE_IN1=0.5, -DSCALE_IN2=0.5
  * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
+ * @attention Operation type(code) specifying which operation to perform should be passed as preprocessor argument using
+ * -DOP_CODE = number. e.g. -DOP_CODE=1
  *
  * @param[in]  input1_ptr                            Pointer to the source tensor. Supported data types: QASYMM8
  * @param[in]  input1_stride_x                       Stride of the source tensor in X dimension (in bytes)
@@ -72,10 +77,16 @@ __kernel void comparison_op_qasymm8(
 
     const VEC_FLOAT in1f32  = CONVERT(in_a, VEC_FLOAT) * (VEC_FLOAT)((float)SCALE_IN1);
     const VEC_FLOAT in2f32  = CONVERT(in_b, VEC_FLOAT) * (VEC_FLOAT)((float)SCALE_IN2);
-#if OPCODE == 0 //EQUAL QUANTIZED
-    VSTORE(VEC_SIZE)(CONVERT(in1f32 == in2f32, VEC_OUT), 0, (__global DATA_TYPE_OUT *)out.ptr);
-#elif OPCODE == 1 //NOT EQUAL QUANTIZED
-    VSTORE(VEC_SIZE)(CONVERT(in1f32 != in2f32, VEC_OUT), 0, (__global DATA_TYPE_OUT *)out.ptr);
-#endif
+
+    #if OPCODE == 1 //EQUAL QUANTIZED
+      VSTORE(VEC_SIZE)(CONVERT(in1f32 == in2f32, VEC_OUT), 0, (__global DATA_TYPE_OUT *)out.ptr);
+
+    #elif OPCODE == 2 //NOT EQUAL QUANTIZED
+      VSTORE(VEC_SIZE)(CONVERT(in1f32 != in2f32, VEC_OUT), 0, (__global DATA_TYPE_OUT *)out.ptr);
+
+    #else // OP NOT SUPPORTED
+      return;
+
+    #endif
 }
-#endif // defined(OFFSET_IN1) && defined(OFFSET_IN2) && defined(SCALE_IN1) && defined(SCALE_IN2) && defined(VEC_SIZE)
+#endif // defined(OFFSET_IN1) && defined(OFFSET_IN2) && defined(SCALE_IN1) && defined(SCALE_IN2) && defined(DATA_TYPE_OUT)

diff --git a/libs/ARMComputeEx/src/core/CL/kernels/CLComparisonOpKernel.cpp b/libs/ARMComputeEx/src/core/CL/kernels/CLComparisonOpKernel.cpp
index 136ed3a..1b390ed 100644 (file)
--- a/libs/ARMComputeEx/src/core/CL/kernels/CLComparisonOpKernel.cpp
+++ b/libs/ARMComputeEx/src/core/CL/kernels/CLComparisonOpKernel.cpp
@@ -69,18 +69,18 @@ void CLComparisonOpKernel::configure(const ICLTensor *input1, const ICLTensor *i
 
   // Create kernel
   std::string kernel_name = "comparison_op";
-  int op_code;
+  int op_code = 0;
 
   switch (op)
   {
     case ComparisonOperation::EQUAL:
-      op_code = 0;
+      op_code = 1;
       break;
     case ComparisonOperation::NOT_EQUAL:
-      op_code = 1;
+      op_code = 2;
       break;
     default:
-      throw std::runtime_error("Not supported, yet");
+      throw std::runtime_error(" Operation not supported, yet");
   }
 
   std::set<std::string> build_opts;