IntelGPU backend: Convert operation (#1394)

src/ngraph/runtime/intelgpu/intelgpu_backend.cpp

@@ -328,6 +328,25 @@ bool runtime::intelgpu::IntelGPUBackend::compile(shared_ptr<Function> func)
                                      reversed_axes);
             }
         }
+        else if ("Convert" == op->description())
+        {
+            arguments_check(op, 1, 1);
+
+            if (get_input_type(op) == get_output_type(op))
+            {
+                do_equal_propagation(topology, get_input_name(op), get_output_name(op));
+            }
+            else
+            {
+                do_convert_operation(topology,
+                                     get_input_name(op),
+                                     get_input_shape(op),
+                                     get_input_type(op),
+                                     get_output_name(op),
+                                     get_output_shape(op),
+                                     get_output_type(op));
+            }
+        }
         else if ("Concat" == op->description())
         {
             if (op->get_inputs().empty() || op->get_outputs().size() != 1)

src/ngraph/runtime/intelgpu/intelgpu_op_custom_kernels.cpp

@@ -732,3 +732,44 @@ void runtime::intelgpu::do_reverse_operation(cldnn::topology& topology,
                                                  gws);
     topology.add(op_reverse);
 }
+
+void runtime::intelgpu::do_convert_operation(cldnn::topology& topology,
+                                             const string& input_name,
+                                             const Shape& input_shape,
+                                             const element::Type& input_type,
+                                             const string& output_name,
+                                             const Shape& output_shape,
+                                             const element::Type& output_type)
+{
+    const cldnn::layout layout = IntelGPULayout::create_cldnn_layout(output_type, output_shape);
+    const string entry_point_name = "convert_" + output_name;
+    const string& input_type_name = input_type.c_type_string();
+    const string& output_type_name = output_type.c_type_string();
+    codegen::CodeWriter writer;
+    vector<size_t> gws;
+
+    writer << "__kernel void " << entry_point_name << "(const __global " << input_type_name
+           << " input" << array_dims(input_shape) << ", __global " << output_type_name << " output"
+           << array_dims(output_shape) << ")\n";
+
+    writer.block_begin();
+    {
+        gws = generate_loops(writer, output_shape, true);
+
+        writer << "output" << access_dims(output_shape) << " = convert_" << output_type_name
+               << "(input" << access_dims(output_shape) << ");\n";
+
+        generate_loops(writer, output_shape, false);
+    }
+    writer.block_end();
+
+    const cldnn::custom_gpu_primitive op_convert(output_name,
+                                                 {input_name},
+                                                 {writer.get_code()},
+                                                 entry_point_name,
+                                                 get_kernel_args(1, 1),
+                                                 "",
+                                                 layout,
+                                                 gws);
+    topology.add(op_convert);
+}

src/ngraph/runtime/intelgpu/intelgpu_op_custom_kernels.hpp

@@ -92,6 +92,14 @@ namespace ngraph
                                       const element::Type& output_type,
                                       const AxisSet& reversed_axes);
 
+            void do_convert_operation(cldnn::topology& topology,
+                                      const std::string& input_name,
+                                      const Shape& input_shape,
+                                      const element::Type& input_type,
+                                      const std::string& output_name,
+                                      const Shape& output_shape,
+                                      const element::Type& output_type);
+
             // Helper functions used in cldnn::custom_gpu_primitive kernels
             std::vector<cldnn_arg> get_kernel_args(size_t input, size_t output);
             std::string array_dims(const Shape& dimentions);

src/ngraph/runtime/intelgpu/unit_test.manifest

@@ -46,7 +46,6 @@ broadcast_vector_rowwise_int64
 ceiling
 concat_matrix_int64
 constant_multi_use
-convert_float32_bool
 convert_int32_bool
 convert_int32_float32
 convolution_2d_1item