diff --git a/lite/npu/bridge/CMakeLists.txt b/lite/npu/bridge/CMakeLists.txt index e4c1cca5a8e82a3de2fb70676a68f707ae1eba97..cf3ad9905588d2501952f3eba0f39336e199b54b 100644 --- a/lite/npu/bridge/CMakeLists.txt +++ b/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}") diff --git a/lite/npu/bridge/pad2d_op.cc b/lite/npu/bridge/pad2d_op.cc new file mode 100644 index 0000000000000000000000000000000000000000..e99b7da3a6f7302a3fba9be4da61d72b1086dfde --- /dev/null +++ b/lite/npu/bridge/pad2d_op.cc @@ -0,0 +1,86 @@ +// 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 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 pad2d_node = + std::make_shared(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("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>("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(unique_op_type + "/padding"); + npu_padding->set_attr_value(CreateTensorAndFillData(padding, {xds, 2})); + pad2d_node->set_input_padding(*npu_padding); + OpList::Global().add(npu_padding); + + if (mode == "constant") { + auto pad_value = op_info->GetAttr("pad_value"); + auto npu_pad_value = + std::make_shared(unique_op_type + "/pad_value"); + npu_pad_value->set_attr_value(CreateTensorAndFillData({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); diff --git a/lite/npu/bridge/pad2d_op_test.cc b/lite/npu/bridge/pad2d_op_test.cc new file mode 100644 index 0000000000000000000000000000000000000000..404bb4a62f08639c8e9003e6d03547e5d8d569cd --- /dev/null +++ b/lite/npu/bridge/pad2d_op_test.cc @@ -0,0 +1,189 @@ +// 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 +#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 +void pad2d_ref(const std::shared_ptr 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>("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("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("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(); + auto out_data = out->mutable_data(); +#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 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(x); + auto x_data = x->mutable_data(); + + // 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(opdesc, &scope); + pad2d_ref(op); + out_ref->CopyDataFrom(*out); + + LauchOp(op, {x_var_name}, {out_var_name}); + + // compare results + auto* out_data = out->mutable_data(); + auto* out_ref_data = out_ref->mutable_data(); + 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{0, 0, 0, 0},*/ + std::vector{0, 0, 0, 1}, + std::vector{0, 1, 0, 2}, + std::vector{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);