add npu pad2d op converter (#1896)

test=develop

lite/npu/bridge/CMakeLists.txt

@@ -20,6 +20,7 @@ lite_cc_library(npu_bridge_transpose_op SRCS transpose_op.cc DEPS ${npu_bridge_d
 lite_cc_library(npu_bridge_split_op SRCS split_op.cc DEPS ${npu_bridge_deps})
 lite_cc_library(npu_bridge_concat_op SRCS concat_op.cc DEPS ${npu_bridge_deps})
 lite_cc_library(npu_bridge_shuffle_channel_op SRCS shuffle_channel_op.cc DEPS ${npu_bridge_deps})
+lite_cc_library(npu_bridge_pad2d_op SRCS pad2d_op.cc DEPS ${npu_bridge_deps})
 
 set(npu_bridges
         npu_bridge_registry
@@ -40,6 +41,7 @@ set(npu_bridges
         npu_bridge_split_op
         npu_bridge_concat_op
         npu_bridge_shuffle_channel_op
+        npu_bridge_pad2d_op
         CACHE INTERNAL "npu_bridges")
 
 lite_cc_library(npu_test_helper SRCS test_helper.cc DEPS npu_helper ${npu_ddk_libs} ${npu_bridges} ${npu_kernels} ${ops})
@@ -60,5 +62,6 @@ lite_cc_test(test_npu_bridge_transpose_op SRCS transpose_op_test.cc DEPS npu_tes
 lite_cc_test(test_npu_bridge_split_op SRCS split_op_test.cc DEPS npu_test_helper)
 lite_cc_test(test_npu_bridge_concat_op SRCS concat_op_test.cc DEPS npu_test_helper)
 lite_cc_test(test_npu_bridge_shuffle_channel_op SRCS shuffle_channel_op_test.cc DEPS npu_test_helper)
+lite_cc_test(test_npu_bridge_pad2d_op SRCS pad2d_op_test.cc DEPS npu_test_helper)
 
 message(STATUS "+++++ npu_bridges: ${npu_bridges}")

lite/npu/bridge/pad2d_op.cc

+// Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "ai_ddk_lib/include/graph/buffer.h"
+#include "ai_ddk_lib/include/graph/graph.h"
+#include "ai_ddk_lib/include/graph/model.h"
+#include "ai_ddk_lib/include/graph/op/all_ops.h"
+#include "ai_ddk_lib/include/graph/operator.h"
+#include "ai_ddk_lib/include/graph/operator_reg.h"
+#include "lite/npu/bridge/registry.h"
+#include "lite/npu/bridge/utils.h"
+
+namespace paddle {
+namespace lite {
+namespace npu {
+namespace bridge {
+
+node_map_type Pad2dConverter(const std::shared_ptr<lite::OpLite> pad2d_op,
+                             const node_map_type& inputs_map) {
+  auto scope = pad2d_op->scope();
+  auto op_info = pad2d_op->op_info();
+  auto op_type = op_info->Type();
+  auto unique_op_type = UniqueName(op_type);
+  LOG(INFO) << "Converting " + op_type + "...";
+
+  std::shared_ptr<ge::op::Pad> pad2d_node =
+      std::make_shared<ge::op::Pad>(unique_op_type);
+  auto x_var_name = op_info->Input("X").front();
+  pad2d_node->set_input_x(*inputs_map.at(x_var_name));
+  OpList::Global().add(inputs_map.at(x_var_name));
+  OpList::Global().add(pad2d_node);
+
+  auto mode = op_info->GetAttr<std::string>("mode");
+  if (mode == "constant") {
+    pad2d_node->set_attr_mode(0);
+  } else if (mode == "reflect") {
+    LOG(FATAL) << "NPU doesn't support this pad mod: " << mode;
+    pad2d_node->set_attr_mode(1);
+  } else {
+    LOG(FATAL) << "NPU doesn't support this pad mod: " << mode;
+  }
+
+  auto x_dims = scope->FindTensor(x_var_name)->dims();
+  auto padding = op_info->GetAttr<std::vector<int>>("paddings");
+  CHECK_EQ(padding.size(), 4);
+  int xds = x_dims.size();
+  padding.insert(padding.begin(), xds * 2 - 4, 0);
+  auto npu_padding =
+      std::make_shared<ge::op::Const>(unique_op_type + "/padding");
+  npu_padding->set_attr_value(CreateTensorAndFillData<int>(padding, {xds, 2}));
+  pad2d_node->set_input_padding(*npu_padding);
+  OpList::Global().add(npu_padding);
+
+  if (mode == "constant") {
+    auto pad_value = op_info->GetAttr<float>("pad_value");
+    auto npu_pad_value =
+        std::make_shared<ge::op::Const>(unique_op_type + "/pad_value");
+    npu_pad_value->set_attr_value(CreateTensorAndFillData<float>({pad_value}));
+    pad2d_node->set_input_constant_values(*npu_pad_value);
+    OpList::Global().add(npu_pad_value);
+
+    pad2d_node->set_attr_T(0);  // type of pad_value:  0:float  3:int32
+  }
+
+  node_map_type outputs_map;
+  outputs_map[op_info->Output("Out").front()] = pad2d_node;
+  return outputs_map;
+}
+
+}  // namespace bridge
+}  // namespace npu
+}  // namespace lite
+}  // namespace paddle
+
+REGISTER_NPU_BRIDGE(pad2d, paddle::lite::npu::bridge::Pad2dConverter);

lite/npu/bridge/pad2d_op_test.cc

+// Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "lite/operators/pad2d_op.h"
+#include <gtest/gtest.h>
+#include "lite/core/op_registry.h"
+#include "lite/npu/bridge/registry.h"
+#include "lite/npu/bridge/test_helper.h"
+
+namespace paddle {
+namespace lite {
+namespace npu {
+namespace bridge {
+
+template <typename dtype>
+void pad2d_ref(const std::shared_ptr<operators::Pad2dOpLite> op) {
+  Scope* scope = op->scope();
+  const OpInfo* op_info = op->op_info();
+  auto x = scope->FindMutableTensor(op_info->Input("X").front());
+  auto out = scope->FindMutableTensor(op_info->Output("Out").front());
+
+  auto paddings = op_info->GetAttr<std::vector<int>>("paddings");
+  int pad_top = paddings[0];
+  int pad_bottom = paddings[1];
+  int pad_left = paddings[2];
+  int pad_right = paddings[3];
+
+  auto mode = op_info->GetAttr<std::string>("mode");
+  int pad_mode;
+  if (mode == "constant") {
+    pad_mode = 0;
+  } else if (mode == "reflect") {
+    pad_mode = 1;
+  } else if (mode == "edge") {
+    pad_mode = 2;
+  } else {
+    LOG(FATAL) << "Unknown mode type";
+  }
+  float pad_value = op_info->GetAttr<float>("pad_value");
+
+  auto out_dims = out->dims();
+  int n = out_dims[0];
+  int c = out_dims[1];
+  int h = out_dims[2];
+  int w = out_dims[3];
+
+  int in_w = w - pad_left - pad_right;
+  int in_h = h - pad_bottom - pad_top;
+  int spatial_size_out = w * h;
+  int spatial_size_in = in_w * in_h;
+
+  auto x_data = x->data<float>();
+  auto out_data = out->mutable_data<float>();
+#pragma omp parallel for
+  for (int i = 0; i < n * c; ++i) {
+    const float* din_batch = x_data + i * spatial_size_in;
+    float* dout_batch = out_data + i * spatial_size_out;
+    int in_y = 0;
+    int in_x = 0;
+    for (int y = 0; y < h; ++y) {
+      for (int x = 0; x < w; ++x) {
+        switch (pad_mode) {
+          case 0:
+            in_y = y - pad_top;
+            in_x = x - pad_left;
+            dout_batch[y * w + x] =
+                (in_x >= 0 && in_x < in_w) && (in_y >= 0 && in_y < in_h)
+                    ? din_batch[in_y * in_w + in_x]
+                    : pad_value;
+            break;
+          case 1:
+            in_x =
+                std::min(std::max(pad_left, x), in_w + pad_left - 1) - pad_left;
+            in_y = std::min(std::max(pad_top, y), in_h + pad_top - 1) - pad_top;
+            dout_batch[y * w + x] = din_batch[in_y * in_w + in_x];
+            break;
+          case 2:
+            in_y = y - pad_top;
+            in_x = x - pad_left;
+            in_y = std::max(in_y, -in_y);
+            in_y = std::min(in_y, 2 * in_h - in_y - 2);
+            in_x = std::max(in_x, -in_x);
+            in_x = std::min(in_x, 2 * in_w - in_x - 2);
+            dout_batch[y * w + x] = din_batch[in_y * in_w + in_x];
+            break;
+          default:
+            LOG(ERROR) << "ERROR: unknown pad mode:" << pad_mode;
+        }
+      }
+    }
+  }
+}
+
+void test_pad2d(int bs,
+                int ic,
+                int ih,
+                int iw,
+                std::vector<int> paddings,
+                float pad_value,
+                std::string mode) {
+  // prepare input&output variables
+  Scope scope;
+  std::string x_var_name = "x";
+  std::string out_var_name = "out";
+  std::string out_ref_var_name = "out_ref";
+  auto* x = scope.NewTensor(x_var_name);
+  auto* out = scope.NewTensor(out_var_name);
+  auto* out_ref = scope.NewTensor(out_ref_var_name);
+  x->Resize({bs, ic, ih, iw});
+
+  // initialize input&output data
+  //  FillTensor<float, int>(x);
+  auto x_data = x->mutable_data<float>();
+
+  // initialize op desc
+  cpp::OpDesc opdesc;
+  opdesc.SetType("pad2d");
+  opdesc.SetInput("X", {x_var_name});
+  opdesc.SetOutput("Out", {out_var_name});
+  opdesc.SetAttr("paddings", paddings);
+  opdesc.SetAttr("pad_value", pad_value);
+  opdesc.SetAttr("mode", mode);
+  opdesc.SetAttr("data_format", std::string("NCHW"));
+
+  auto op = CreateOp<operators::Pad2dOpLite>(opdesc, &scope);
+  pad2d_ref<float>(op);
+  out_ref->CopyDataFrom(*out);
+
+  LauchOp(op, {x_var_name}, {out_var_name});
+
+  // compare results
+  auto* out_data = out->mutable_data<float>();
+  auto* out_ref_data = out_ref->mutable_data<float>();
+  for (int i = 0; i < out->numel(); i++) {
+    EXPECT_NEAR(out_data[i], out_ref_data[i], 1e-2) << "-----" << i;
+  }
+}
+
+TEST(NPUBridges, pad2d) {
+#if 1
+  for (auto bs : {1, 4, 7}) {
+    for (auto ic : {1, 4, 7}) {
+      for (auto ih : {1, 4, 7}) {
+        for (auto iw : {1, 4, 7}) {
+          for (auto paddings : {/*std::vector<int>{0, 0, 0, 0},*/
+                                std::vector<int>{0, 0, 0, 1},
+                                std::vector<int>{0, 1, 0, 2},
+                                std::vector<int>{1, 2, 3, 4}}) {
+            // npu not support pad_value!=0
+            for (auto pad_value : {0.f /*,1.f*/}) {
+              // npu only support constant
+              for (auto mode : {"constant" /*, "reflect", "edge"*/}) {
+                if (mode == "edge") continue;
+                VLOG(3) << "bs: " << bs << "  ic: " << ic << "  ih: " << ih
+                        << "  iw: " << iw << "  paddings: {" << paddings[0]
+                        << "," << paddings[1] << "," << paddings[2] << ","
+                        << paddings[3] << "}"
+                        << "  pad_value: " << pad_value << "  mode: " << mode;
+                test_pad2d(bs, ic, ih, iw, paddings, pad_value, mode);
+              }
+            }
+          }
+        }
+      }
+    }
+  }
+#else
+  test_pad2d(1, 1, 1, 1, {0, 0, 0, 1}, 0, "constant");
+#endif
+}
+
+}  // namespace bridge
+}  // namespace npu
+}  // namespace lite
+}  // namespace paddle
+
+USE_LITE_OP(pad2d);
+USE_NPU_BRIDGE(pad2d);