Merge branch 'develop' into LapackEigh

paddle/fluid/inference/tensorrt/convert/pool2d_op.cc

@@ -87,6 +87,10 @@ class Pool2dOpConverter : public OpConverter {
     bool adaptive = false;
     if (op_desc.HasAttr("adaptive"))
       adaptive = BOOST_GET_CONST(bool, op_desc.GetAttr("adaptive"));
+    std::string padding_algorithm = "EXPLICIT";
+    if (op_desc.HasAttr("padding_algorithm"))
+      padding_algorithm =
+          BOOST_GET_CONST(std::string, op_desc.GetAttr("padding_algorithm"));
 
     nvinfer1::PoolingType nv_pool_type = nvinfer1::PoolingType::kMAX;
     nvinfer1::ReduceOperation reduce_operation =
@@ -124,6 +128,9 @@ class Pool2dOpConverter : public OpConverter {
         pool_layer->setStride(nv_strides);
         pool_layer->setPadding(nv_paddings);
         pool_layer->setAverageCountExcludesPadding(exclusive);
+        if (padding_algorithm == "SAME") {
+          pool_layer->setPaddingMode(nvinfer1::PaddingMode::kSAME_UPPER);
+        }
         layer = pool_layer;
       } else if (global_pooling) {
         auto *reduce_layer = TRT_ENGINE_ADD_LAYER(engine_, Reduce, *input1,
@@ -159,6 +166,9 @@ class Pool2dOpConverter : public OpConverter {
       auto output_name = op_desc.Output("Out")[0];
       pool_layer->setStride(nv_strides);
       pool_layer->setPadding(nv_paddings);
+      if (padding_algorithm == "SAME") {
+        pool_layer->setPaddingMode(nvinfer1::PaddingMode::kSAME_UPPER);
+      }
       pool_layer->setAverageCountExcludesPadding(exclusive);
       pool_layer->setName(("pool2d (Output: " + output_name + ")").c_str());
       pool_layer->getOutput(0)->setName(output_name.c_str());
@@ -198,6 +208,9 @@ class Pool2dOpConverter : public OpConverter {
                           "trt pool layer in converter could not be created."));
       pool_layer->setStride(nv_strides);
       pool_layer->setPadding(nv_paddings);
+      if (padding_algorithm == "SAME") {
+        pool_layer->setPaddingMode(nvinfer1::PaddingMode::kSAME_UPPER);
+      }
       pool_layer->setAverageCountExcludesPadding(exclusive);
       layer = pool_layer;
     } else {

paddle/fluid/inference/tensorrt/op_teller.cc

@@ -172,6 +172,22 @@ bool OpTeller::Tell(const framework::ir::Node* node, bool use_no_calib_int8,
       std::vector<int> paddings =
           BOOST_GET_CONST(std::vector<int>, desc.GetAttr("paddings"));
       if (paddings.size() > 2) return false;
+      if (desc.HasAttr("exclusive")) {
+        if (BOOST_GET_CONST(bool, desc.GetAttr("exclusive"))) {
+          std::vector<int> ksize =
+              BOOST_GET_CONST(std::vector<int>, desc.GetAttr("ksize"));
+          for (size_t i = 0; i < ksize.size(); i++) {
+            if (ksize[i] <= paddings[i]) {
+              VLOG(3) << "the padding size should be less than the filter size "
+                         "for exclusive-counting pooling.";
+              return false;
+            }
+          }
+        }
+      }
+      if (desc.HasAttr("ceil_mode")) {
+        if (BOOST_GET_CONST(bool, desc.GetAttr("ceil_mode"))) return false;
+      }
       if (desc.Input("X").size() != 1) {
         VLOG(3) << "TRT Pool2d expect 1 input, but got "
                 << desc.Input("X").size();
@@ -440,6 +456,10 @@ bool OpTeller::Tell(const framework::ir::Node* node, bool use_no_calib_int8,
       }
     }
 
+    if (op_type == "anchor_generator") {
+      if (!with_dynamic_shape) return false;
+    }
+
     if (op_type == "yolo_box") {
       if (with_dynamic_shape) return false;
       bool has_attrs =

paddle/fluid/operators/determinant_op.cc

@@ -48,6 +48,8 @@ class DeterminantGradOp : public framework::OperatorWithKernel {
     OP_INOUT_CHECK(ctx->HasInput("Input"), "Input", "Input",
                    "DeterminantGradOp");
     OP_INOUT_CHECK(ctx->HasInput("Out"), "Input", "Out", "DeterminantGradOp");
+    OP_INOUT_CHECK(ctx->HasInput(framework::GradVarName("Out")), "Input",
+                   framework::GradVarName("Out"), "DeterminantGradOp");
     OP_INOUT_CHECK(ctx->HasOutput(framework::GradVarName("Input")), "Output",
                    framework::GradVarName("Input"), "DeterminantGradOp");
 
@@ -117,7 +119,8 @@ class SlogDeterminantGradOp : public framework::OperatorWithKernel {
                    "SlogDeterminantGradOp");
     OP_INOUT_CHECK(ctx->HasInput("Out"), "Input", "Out",
                    "SlogDeterminantGradOp");
-
+    OP_INOUT_CHECK(ctx->HasInput(framework::GradVarName("Out")), "Input",
+                   framework::GradVarName("Out"), "SlogDeterminantGradOp");
     OP_INOUT_CHECK(ctx->HasOutput(framework::GradVarName("Input")), "Output",
                    framework::GradVarName("Input"), "SlogDeterminantGradOp");
 
@@ -179,13 +182,13 @@ REGISTER_OPERATOR(slogdeterminant, ops::SlogDeterminantOp,
                   ops::SlogDeterminantGradOpMaker<paddle::imperative::OpBase>);
 
 REGISTER_OPERATOR(slogdeterminant_grad,
-                  ops::DeterminantGradOp)  // reuse det grad op
+                  ops::SlogDeterminantGradOp)  // reuse det grad op
 
 REGISTER_OP_CPU_KERNEL(
     slogdeterminant, ops::SlogDeterminantKernel<plat::CPUDeviceContext, float>,
     ops::SlogDeterminantKernel<plat::CPUDeviceContext, double>);
 
 REGISTER_OP_CPU_KERNEL(
     slogdeterminant_grad,
-    ops::DeterminantGradKernel<plat::CPUDeviceContext, float>,
-    ops::DeterminantGradKernel<plat::CPUDeviceContext, double>);
+    ops::SlogDeterminantGradKernel<plat::CPUDeviceContext, float>,
+    ops::SlogDeterminantGradKernel<plat::CPUDeviceContext, double>);

paddle/fluid/operators/determinant_op.cu

@@ -14,42 +14,6 @@ limitations under the License. */
 
 #include "paddle/fluid/framework/op_registry.h"
 #include "paddle/fluid/operators/determinant_op.h"
-#include "paddle/fluid/platform/cuda_primitives.h"
-
-namespace paddle {
-namespace operators {
-
-using platform::PADDLE_CUDA_NUM_THREADS;
-using Tensor = framework::Tensor;
-
-template <typename T>
-__global__ void DeterminantGrad(const size_t numel, T* out) {
-  int tid = threadIdx.x + blockIdx.x * blockDim.x;
-  if (tid < numel) {
-    out[tid] = static_cast<T>(1);
-  }
-}
-
-template <typename T>
-class DeterminantGradCUDAKernel : public framework::OpKernel<T> {
- public:
-  void Compute(const framework::ExecutionContext& context) const override {
-    const auto* dout = context.Input<Tensor>(framework::GradVarName("Out"));
-    const T* dout_data = dout->data<T>();
-    auto dout_dim = vectorize(dout->dims());
-
-    auto* dx = context.Output<Tensor>(framework::GradVarName("Input"));
-    T* dx_data = dx->mutable_data<T>(context.GetPlace());
-
-    int64_t numel = dx->numel();
-    for (int64_t idx = 0; idx < numel; idx++) {
-      dx_data[idx] = static_cast<T>(1);
-    }
-  }
-};
-
-}  // namespace operators
-}  // namespace paddle
 
 namespace ops = paddle::operators;
 namespace plat = paddle::platform;