IntelGPU backend: And, Or operations (#1337)

* IntelGPU backend: And, Or operations

* Code format update: intelgpu_backend.cpp and intelgpu_op_custom_kernels.cpp

* Update logical operations

src/ngraph/runtime/intelgpu/intelgpu_backend.cpp

@@ -71,7 +71,7 @@ static void arguments_check(const shared_ptr<Node>& op, size_t input, size_t out
 
 static void argument_type_check(const element::Type& type)
 {
-    if (type != element::f32)
+    if (type != element::f32 && type != element::boolean)
     {
         ostringstream os;
         os << "Kernel data type " << type << " is not supported";
@@ -147,9 +147,13 @@ static void do_logical_operation(cldnn::topology& topology,
 
     const string& inputA_name = op->get_inputs().at(0).get_tensor().get_name();
     const Shape& inputA_shape = op->get_inputs().at(0).get_shape();
+    const string& inputA_type =
+        op->get_inputs().at(0).get_tensor().get_element_type().c_type_string();
     argument_type_check(op->get_inputs().at(0).get_tensor().get_element_type());
     const string& inputB_name = op->get_inputs().at(1).get_tensor().get_name();
     const Shape& inputB_shape = op->get_inputs().at(1).get_shape();
+    const string& inputB_type =
+        op->get_inputs().at(1).get_tensor().get_element_type().c_type_string();
     argument_type_check(op->get_inputs().at(1).get_tensor().get_element_type());
     const string& output_name = op->get_outputs().begin()->get_tensor().get_name();
     const Shape& output_shape = op->get_outputs().begin()->get_shape();
@@ -158,8 +162,10 @@ static void do_logical_operation(cldnn::topology& topology,
     runtime::intelgpu::do_logic_kernel(topology,
                                        inputA_name,
                                        inputA_shape,
+                                       inputA_type,
                                        inputB_name,
                                        inputB_shape,
+                                       inputB_type,
                                        output_name,
                                        output_shape,
                                        output_type,
@@ -603,6 +609,14 @@ bool runtime::intelgpu::IntelGPUBackend::compile(shared_ptr<Function> func)
         {
             do_logical_operation(topology, op, " <= ");
         }
+        else if ("And" == op->description())
+        {
+            do_logical_operation(topology, op, " && ");
+        }
+        else if ("Or" == op->description())
+        {
+            do_logical_operation(topology, op, " || ");
+        }
         else if ("Subtract" == op->description())
         {
             do_eltwise_operation(topology, op, cldnn::eltwise_mode::sub);

src/ngraph/runtime/intelgpu/intelgpu_layout.cpp

@@ -82,10 +82,10 @@ cldnn::tensor runtime::intelgpu::IntelGPULayout::create_cldnn_tensor(const Shape
 {
     vector<size_t> idx(4, 1);
     size_t index = 0;
-    const size_t total_zise = shape_size<Shape>(element_shape);
+    const size_t total_size = shape_size<Shape>(element_shape);
 
     // clDNN requires at least scalar tensor size. We can't create zero sized tensors
-    if (total_zise != 0)
+    if (total_size != 0)
     {
         for (auto i = element_shape.crbegin(); i != element_shape.crend() && index < 3;
              ++i, ++index)

src/ngraph/runtime/intelgpu/intelgpu_op_custom_kernels.cpp

@@ -606,8 +606,10 @@ void runtime::intelgpu::do_select_operation(cldnn::topology& topology,
 void runtime::intelgpu::do_logic_kernel(cldnn::topology& topology,
                                         const string& inputA_name,
                                         const Shape& inputA_shape,
+                                        const string& inputA_type,
                                         const string& inputB_name,
                                         const Shape& inputB_shape,
+                                        const string& inputB_type,
                                         const string& output_name,
                                         const Shape& output_shape,
                                         const element::Type& output_type,
@@ -617,8 +619,8 @@ void runtime::intelgpu::do_logic_kernel(cldnn::topology& topology,
     const string entry_point_name = "logic_" + output_name;
     codegen::CodeWriter writer;
 
-    writer << "__kernel void " << entry_point_name << "(const __global float inputA"
-           << array_dims(inputA_shape) << ", const __global float inputB"
+    writer << "__kernel void " << entry_point_name << "(const __global " << inputA_type << " inputA"
+           << array_dims(inputA_shape) << ", const __global " << inputB_type << " inputB"
            << array_dims(inputB_shape) << ", __global char output" << array_dims(output_shape)
            << ")\n";
 
@@ -636,6 +638,7 @@ void runtime::intelgpu::do_logic_kernel(cldnn::topology& topology,
 
         writer << "if (inputA" << access_dims(inputA_shape) << operation << "inputB"
                << access_dims(inputB_shape) << ")\n";
+
         writer.block_begin();
         {
             writer << "output" << access_dims(output_shape) << " = 1;\n";

src/ngraph/runtime/intelgpu/intelgpu_op_custom_kernels.hpp

@@ -73,8 +73,10 @@ namespace ngraph
             void do_logic_kernel(cldnn::topology& topology,
                                  const std::string& inputA_name,
                                  const Shape& inputA_shape,
+                                 const std::string& inputA_type,
                                  const std::string& inputB_name,
                                  const Shape& inputB_shape,
+                                 const std::string& inputB_type,
                                  const std::string& output_name,
                                  const Shape& output_shape,
                                  const element::Type& output_type,