diff --git a/paddle/fluid/inference/tensorrt/convert/CMakeLists.txt b/paddle/fluid/inference/tensorrt/convert/CMakeLists.txt index 3864f337bdadc61e7531304e2cf2ee52a25253f2..7f5e5a622b9e8578f9b39aa9024d653d7b17138f 100644 --- a/paddle/fluid/inference/tensorrt/convert/CMakeLists.txt +++ b/paddle/fluid/inference/tensorrt/convert/CMakeLists.txt @@ -13,6 +13,7 @@ nv_test(test_trt_fc_op SRCS test_fc_op.cc fc_op.cc DEPS ${FLUID_CORE_MODULES} tensorrt_engine mul_op SERIAL) nv_test(test_trt_activation_op SRCS test_activation_op.cc activation_op.cc DEPS ${FLUID_CORE_MODULES} tensorrt_engine activation_op SERIAL) - +nv_test(test_trt_conv_op SRCS test_conv2d_op.cc conv2d_op.cc + DEPS ${FLUID_CORE_MODULES} tensorrt_engine conv_op SERIAL) nv_test(test_trt_pool2d_op SRCS test_pool2d_op.cc pool2d_op.cc DEPS ${FLUID_CORE_MODULES} tensorrt_engine pool_op SERIAL) diff --git a/paddle/fluid/inference/tensorrt/convert/conv2d_op.cc b/paddle/fluid/inference/tensorrt/convert/conv2d_op.cc index 8e7e23377d4b2fe7afd51f1f58048fc4ed3c6d99..dba1d50b2d1c487ced8e6ca51f2d257641ad5fc7 100644 --- a/paddle/fluid/inference/tensorrt/convert/conv2d_op.cc +++ b/paddle/fluid/inference/tensorrt/convert/conv2d_op.cc @@ -20,11 +20,60 @@ namespace tensorrt { class Conv2dOpConverter : public OpConverter { public: - Conv2dOpConverter() {} void operator()(const framework::proto::OpDesc& op, const framework::Scope& scope, bool test_mode) override { LOG(INFO) << "convert a fluid conv2d op to tensorrt conv layer without bias"; + + framework::OpDesc op_desc(op, nullptr); + PADDLE_ENFORCE_EQ(op_desc.Input("Input").size(), 1); + PADDLE_ENFORCE_EQ(op_desc.Input("Filter").size(), 1); // Y is a weight + PADDLE_ENFORCE_EQ(op_desc.Output("Output").size(), 1); + + auto* X = engine_->GetITensor(op_desc.Input("Input").front()); + // Declare weights + auto* Y_v = scope.FindVar(op_desc.Input("Filter").front()); + PADDLE_ENFORCE_NOT_NULL(Y_v); + auto* Y_t = Y_v->GetMutable(); + auto* weight_data = Y_t->mutable_data(platform::CPUPlace()); + + PADDLE_ENFORCE_EQ(Y_t->dims().size(), 4UL); + const int n_output = Y_t->dims()[0]; + const int filter_h = Y_t->dims()[2]; + const int filter_w = Y_t->dims()[3]; + + const int groups = boost::get(op_desc.GetAttr("groups")); + const std::vector dilations = + boost::get>(op_desc.GetAttr("dilations")); + const std::vector strides = + boost::get>(op_desc.GetAttr("strides")); + const std::vector paddings = + boost::get>(op_desc.GetAttr("paddings")); + + nvinfer1::DimsHW nv_ksize(filter_h, filter_w); + nvinfer1::DimsHW nv_dilations(dilations[0], dilations[1]); + nvinfer1::DimsHW nv_strides(strides[0], strides[1]); + nvinfer1::DimsHW nv_paddings(paddings[0], paddings[1]); + + TensorRTEngine::Weight weight{nvinfer1::DataType::kFLOAT, + static_cast(weight_data), + Y_t->memory_size() / sizeof(float)}; + + TensorRTEngine::Weight bias{nvinfer1::DataType::kFLOAT, nullptr, 0}; + auto* layer = TRT_ENGINE_ADD_LAYER( + engine_, Convolution, *const_cast(X), n_output, + nv_ksize, weight.get(), bias.get()); + PADDLE_ENFORCE(layer != nullptr); + layer->setStride(nv_strides); + layer->setPadding(nv_paddings); + layer->setDilation(nv_dilations); + layer->setNbGroups(groups); + + auto output_name = op_desc.Output("Output").front(); + engine_->SetITensor(output_name, layer->getOutput(0)); + if (test_mode) { + engine_->DeclareOutput(output_name); + } } }; diff --git a/paddle/fluid/inference/tensorrt/convert/test_conv2d_op.cc b/paddle/fluid/inference/tensorrt/convert/test_conv2d_op.cc new file mode 100644 index 0000000000000000000000000000000000000000..f8711c6b60d74639529624c25429bc245de46479 --- /dev/null +++ b/paddle/fluid/inference/tensorrt/convert/test_conv2d_op.cc @@ -0,0 +1,57 @@ +/* Copyright (c) 2018 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 +#include "paddle/fluid/inference/tensorrt/convert/op_converter.h" +#include "paddle/fluid/inference/tensorrt/convert/ut_helper.h" + +namespace paddle { +namespace inference { +namespace tensorrt { + +TEST(conv2d_op, test) { + std::unordered_set parameters({"conv2d-Y"}); + framework::Scope scope; + TRTConvertValidation validator(5, parameters, scope, 1 << 15); + + validator.DeclInputVar("conv2d-X", nvinfer1::Dims3(2, 5, 5)); + validator.DeclParamVar("conv2d-Y", nvinfer1::Dims4(3, 2, 3, 3)); + validator.DeclOutputVar("conv2d-Out", nvinfer1::Dims3(3, 5, 5)); + + // Prepare Op description + framework::OpDesc desc; + desc.SetType("conv2d"); + desc.SetInput("Input", {"conv2d-X"}); + desc.SetInput("Filter", {"conv2d-Y"}); + desc.SetOutput("Output", {"conv2d-Out"}); + + const std::vector strides({1, 1}); + const std::vector paddings({1, 1}); + const std::vector dilations({1, 1}); + const int groups = 1; + + desc.SetAttr("strides", strides); + desc.SetAttr("paddings", paddings); + desc.SetAttr("dilations", dilations); + desc.SetAttr("groups", groups); + + validator.SetOp(*desc.Proto()); + + validator.Execute(3); +} + +} // namespace tensorrt +} // namespace inference +} // namespace paddle +USE_OP(conv2d); diff --git a/paddle/fluid/inference/tensorrt/convert/test_op_converter.cc b/paddle/fluid/inference/tensorrt/convert/test_op_converter.cc index 9b79f86b0edba983019bd932f52b08711ff36d41..d6651a5b244ba31a01220e6299cb2016ae61fe64 100644 --- a/paddle/fluid/inference/tensorrt/convert/test_op_converter.cc +++ b/paddle/fluid/inference/tensorrt/convert/test_op_converter.cc @@ -25,12 +25,42 @@ TEST(OpConverter, ConvertBlock) { framework::ProgramDesc prog; auto* block = prog.MutableBlock(0); auto* conv2d_op = block->AppendOp(); + + // init trt engine + cudaStream_t stream_; + std::unique_ptr engine_; + engine_.reset(new TensorRTEngine(5, 1 << 15, &stream_)); + engine_->InitNetwork(); + PADDLE_ENFORCE_EQ(cudaStreamCreate(&stream_), 0); + + engine_->DeclareInput("conv2d-X", nvinfer1::DataType::kFLOAT, + nvinfer1::Dims3(2, 5, 5)); + conv2d_op->SetType("conv2d"); + conv2d_op->SetInput("Input", {"conv2d-X"}); + conv2d_op->SetInput("Filter", {"conv2d-Y"}); + conv2d_op->SetOutput("Output", {"conv2d-Out"}); - OpConverter converter; + const std::vector strides({1, 1}); + const std::vector paddings({1, 1}); + const std::vector dilations({1, 1}); + const int groups = 1; + + conv2d_op->SetAttr("strides", strides); + conv2d_op->SetAttr("paddings", paddings); + conv2d_op->SetAttr("dilations", dilations); + conv2d_op->SetAttr("groups", groups); + + // init scope framework::Scope scope; - converter.ConvertBlock(*block->Proto(), {}, scope, - nullptr /*TensorRTEngine*/); + std::vector dim_vec = {3, 2, 3, 3}; + auto* x = scope.Var("conv2d-Y"); + auto* x_tensor = x->GetMutable(); + x_tensor->Resize(framework::make_ddim(dim_vec)); + + OpConverter converter; + converter.ConvertBlock(*block->Proto(), {"conv2d-Y"}, scope, + engine_.get() /*TensorRTEngine*/); } } // namespace tensorrt