IntelGPU backend: MaxPoolBackprop operation (#1427)

* IntelGPU backend: MaxPoolBackprop operation

* PR1427: Initialization loop added

src/ngraph/runtime/intelgpu/intelgpu_backend.cpp

@@ -144,7 +144,6 @@ static void do_pooling_operation(cldnn::topology& topology,
                                  const Shape& pool_shape,
                                  const Strides& pool_strides,
                                  const Shape& pad_below,
-                                 const Shape& pad_above,
                                  const cldnn::pooling_mode mode)
 {
     arguments_check(op, 1, 1);
@@ -469,32 +468,46 @@ bool runtime::intelgpu::IntelGPUBackend::compile(shared_ptr<Function> func)
         else if ("MaxPool" == op->description())
         {
             const shared_ptr<op::MaxPool> max_pool = static_pointer_cast<op::MaxPool>(op);
-            const Shape& pool_shape = max_pool->get_window_shape();
-            const Strides& pool_strides = max_pool->get_window_movement_strides();
-            const Shape& pad_below = max_pool->get_padding_below();
-            const Shape& pad_above = max_pool->get_padding_above();
 
             do_pooling_operation(topology,
                                  op,
-                                 pool_shape,
-                                 pool_strides,
-                                 pad_below,
-                                 pad_above,
+                                 max_pool->get_window_shape(),
+                                 max_pool->get_window_movement_strides(),
+                                 max_pool->get_padding_below(),
                                  cldnn::pooling_mode::max);
         }
+        else if ("MaxPoolBackprop" == op->description())
+        {
+            arguments_check(op, 2, 1);
+
+            const shared_ptr<op::MaxPoolBackprop> max_pool_b =
+                static_pointer_cast<op::MaxPoolBackprop>(op);
+
+            do_max_pool_backprop_operation(topology,
+                                           get_input_name(op, 0),
+                                           get_input_shape(op, 0),
+                                           get_input_name(op, 1),
+                                           get_input_shape(op, 1),
+                                           get_output_name(op),
+                                           get_output_shape(op),
+                                           get_output_type(op),
+                                           max_pool_b->get_window_shape(),
+                                           max_pool_b->get_window_movement_strides(),
+                                           max_pool_b->get_padding_below());
+        }
         else if ("AvgPool" == op->description())
         {
             const shared_ptr<op::AvgPool> avg_pool = static_pointer_cast<op::AvgPool>(op);
-            const Shape& pool_shape = avg_pool->get_window_shape();
-            const Strides& pool_strides = avg_pool->get_window_movement_strides();
-            const Shape& pad_below = avg_pool->get_padding_below();
-            const Shape& pad_above = avg_pool->get_padding_above();
             const cldnn::pooling_mode mode = avg_pool->get_include_padding_in_avg_computation()
                                                  ? cldnn::pooling_mode::average
                                                  : cldnn::pooling_mode::average_no_padding;
 
-            do_pooling_operation(
-                topology, op, pool_shape, pool_strides, pad_below, pad_above, mode);
+            do_pooling_operation(topology,
+                                 op,
+                                 avg_pool->get_window_shape(),
+                                 avg_pool->get_window_movement_strides(),
+                                 avg_pool->get_padding_below(),
+                                 mode);
         }
         else if ("Broadcast" == op->description())
         {

src/ngraph/runtime/intelgpu/intelgpu_op_custom_kernels.cpp

@@ -228,6 +228,182 @@ void runtime::intelgpu::do_pad_operation(cldnn::topology& topology,
     topology.add(op_pad);
 }
 
+void runtime::intelgpu::do_max_pool_backprop_operation(cldnn::topology& topology,
+                                                       const string& input_name,
+                                                       const Shape& input_shape,
+                                                       const string& delta_name,
+                                                       const Shape& delta_shape,
+                                                       const string& output_name,
+                                                       const Shape& output_shape,
+                                                       const element::Type& output_type,
+                                                       const Shape& win_shape,
+                                                       const Shape& win_stride,
+                                                       const Shape& pad_below)
+{
+    const string entry_point_name = "op_max_pool_backprop_" + output_name;
+    const Shape delta_data(delta_shape.cbegin() + 2, delta_shape.cend());
+    const Shape output_data(output_shape.cbegin() + 2, output_shape.cend());
+    codegen::CodeWriter writer;
+    vector<size_t> gws;
+
+    // The kernel name and parameters
+    writer << "__kernel void " << entry_point_name << "(const __global float input"
+           << array_dims(input_shape) << ", const __global float delta" << array_dims(delta_shape)
+           << ", __global float output" << array_dims(output_shape) << ")\n";
+
+    writer.block_begin();
+    {
+        // Main loop over delta input array.
+        writer << "const uint i0 = get_global_id(0);";
+        gws.push_back(delta_shape.at(0));
+        writer << "// for (uint i0 = 0; i0 < " << delta_shape.at(0) << "; ++i0)\n";
+        writer.block_begin();
+        {
+            writer << "const uint i1 = get_global_id(1);";
+            gws.push_back(delta_shape.at(1));
+            writer << "// for (uint i1 = 0; i1 < " << delta_shape.at(1) << "; ++i1)\n";
+            writer.block_begin();
+            {
+                // Initialization output
+                size_t var_idx = 0;
+                for (auto const& i : output_data)
+                {
+                    writer << "for (uint j" << var_idx << " = 0; j" << var_idx << " < " << i
+                           << "; ++j" << var_idx << ")\n";
+                    writer.block_begin();
+                    ++var_idx;
+                }
+
+                writer << "output[i0][i1]";
+                // Additional dimentions for output
+                for (size_t i = 0; i < output_data.size(); ++i)
+                {
+                    writer << "[j" << i << "]";
+                }
+                writer << " = 0.0f;\n";
+
+                // Closing brackets for Initialization loop
+                for (auto const& i : output_data)
+                {
+                    writer.block_end();
+                }
+                // End of output initialization
+
+                // Loops over other output dimensions
+                var_idx = 2;
+                for (auto const& i : delta_data)
+                {
+                    writer << "for (uint i" << var_idx << " = 0; i" << var_idx << " < " << i
+                           << "; ++i" << var_idx << ")\n";
+                    writer.block_begin();
+                    ++var_idx;
+                }
+
+                // Create variables to save coordinates
+                for (size_t i = 0; i < delta_data.size(); ++i)
+                {
+                    writer << "uint save_i" << i + 2 << " = 0;\n";
+                }
+                writer << "float max_elem = FLT_MIN;\n"
+                       << "uint elem_exists = 0;\n";
+
+                // Loop over window shape
+                var_idx = 0;
+                for (auto const& i : win_shape)
+                {
+                    writer << "for (uint w" << var_idx << " = 0; w" << var_idx << " < " << i
+                           << "; ++w" << var_idx << ")\n";
+                    writer.block_begin();
+                    writer << "const uint win_idx" << var_idx << " = (i" << var_idx + 2 << " * "
+                           << win_stride.at(var_idx) << " /*win_stride*/)"
+                           << " + w" << var_idx << " - " << pad_below.at(var_idx)
+                           << " /*pad_below*/;\n";
+                    ++var_idx;
+                }
+
+                // input coordinate condition
+                writer << "if (";
+                // Generate input coordinate condition
+                for (size_t i = 0; i < win_shape.size(); ++i)
+                {
+                    if (i)
+                    {
+                        writer << " && ";
+                    }
+                    writer << "(win_idx" << i << " < " << input_shape.at(i + 2) << ")";
+                }
+                writer << ")\n";
+                writer.block_begin();
+                {
+                    writer << "const float max_local = input[i0][i1]";
+                    // additional dimensions for input
+                    for (size_t i = 0; i < win_shape.size(); ++i)
+                    {
+                        writer << "[win_idx" << i << "]";
+                    }
+                    writer << ";\n";
+
+                    // find maximum condition
+                    writer << "if (max_local > max_elem)\n";
+                    writer.block_begin();
+                    {
+                        writer << "max_elem = max_local;\n"
+                               << "elem_exists = 1;\n";
+
+                        // Save coordinates
+                        for (size_t i = 0; i < delta_data.size(); ++i)
+                        {
+                            writer << "save_i" << i + 2 << " = win_idx" << i << ";\n";
+                        }
+                    } // End of find maximum condition
+                    writer.block_end();
+
+                } // End of input coordinate condition
+                writer.block_end();
+                // Closing brackets for window shape loop
+                for (auto const& i : win_shape)
+                {
+                    writer.block_end();
+                }
+
+                // Elem_exists condition
+                writer << "if (elem_exists)\n";
+                writer.block_begin();
+                {
+                    writer << "output[i0][i1]";
+                    // Additional dimentions for output
+                    for (size_t i = 0; i < delta_data.size(); ++i)
+                    {
+                        writer << "[save_i" << i + 2 << "]";
+                    }
+                    writer << " += delta" << access_dims(delta_shape) << ";\n";
+                } // End of elem_exists condition
+                writer.block_end();
+                // Closing brackets for delta loop
+                for (auto const& i : delta_data)
+                {
+                    writer.block_end();
+                }
+            } // End of loop over i1
+            writer.block_end();
+        } // End of loop over i0
+        writer.block_end();
+
+    } // End of function bracket
+    writer.block_end();
+
+    const cldnn::layout layout = IntelGPULayout::create_cldnn_layout(output_type, output_shape);
+    const cldnn::custom_gpu_primitive op_max_pool_backprop(output_name,
+                                                           {input_name, delta_name},
+                                                           {writer.get_code()},
+                                                           entry_point_name,
+                                                           get_kernel_args(2, 1),
+                                                           "",
+                                                           layout,
+                                                           gws);
+    topology.add(op_max_pool_backprop);
+}
+
 static void do_1d_scalar_mul(codegen::CodeWriter& writer,
                              string& kernel_name,
                              const Shape& shapeA,

src/ngraph/runtime/intelgpu/intelgpu_op_custom_kernels.hpp

@@ -42,6 +42,18 @@ namespace ngraph
                                   const Shape& pad_below,
                                   const Shape& pad_interior);
 
+            void do_max_pool_backprop_operation(cldnn::topology& topology,
+                                                const std::string& input_name,
+                                                const Shape& input_shape,
+                                                const std::string& delta_name,
+                                                const Shape& delta_shape,
+                                                const std::string& output_name,
+                                                const Shape& output_shape,
+                                                const element::Type& output_type,
+                                                const Shape& win_shape,
+                                                const Shape& win_stride,
+                                                const Shape& pad_below);
+
             void do_dot_operation(cldnn::topology& topology,
                                   const std::string& inputA_name,
                                   const Shape& inputA_shape,

src/ngraph/runtime/intelgpu/unit_test.manifest

@@ -18,9 +18,7 @@ backwards_dot_tensor3_tensor3
 backwards_dot_tensor_scalar
 backwards_dot_tensor_vector
 backwards_floor
-backwards_maxpool_n2c1h5w5_kh3kw3_sh2sw2
 backwards_maxpool_n2_c1_hw5_3x3_str2_max
-backwards_maxpool_n4c1h4w4_kh2kw2_sh1sw1
 backwards_maxpool_n4_c1_hw4_2x2_max
 backwards_minimum
 backwards_replace_slice