Perception: turn back to TensorRT for pfe inference (ApolloAuto#12974)

modules/perception/common/perception_gflags.cc

@@ -39,10 +39,10 @@ DEFINE_string(work_root, "", "Project work root direcotry.");
 
 // lidar_point_pillars
 DEFINE_int32(gpu_id, 0, "The id of gpu used for inference.");
-DEFINE_string(pfe_torch_file,
+DEFINE_string(pfe_onnx_file,
               "/apollo/modules/perception/production/data/perception/lidar/"
-              "models/detection/point_pillars/pfe.pt",
-              "The path of pillars feature extractor torch file.");
+              "models/detection/point_pillars/pfe.onnx",
+              "The path of PFE onnx file.");
 DEFINE_string(rpn_onnx_file,
               "/apollo/modules/perception/production/data/perception/lidar/"
               "models/detection/point_pillars/rpn.onnx",

modules/perception/common/perception_gflags.h

@@ -33,7 +33,7 @@ DECLARE_string(work_root);
 
 // lidar_point_pillars
 DECLARE_int32(gpu_id);
-DECLARE_string(pfe_torch_file);
+DECLARE_string(pfe_onnx_file);
 DECLARE_string(rpn_onnx_file);
 DECLARE_double(normalizing_factor);
 DECLARE_int32(num_point_feature);

modules/perception/lidar/lib/detection/lidar_point_pillars/BUILD

@@ -22,12 +22,12 @@ cc_library(
         ":common",
         ":nms_cuda",
         ":params",
+        ":pfe_cuda",
         ":postprocess_cuda",
         ":preprocess_points",
         ":scatter_cuda",
         "//cyber/common",
         "//modules/perception/common:perception_gflags",
-        "//third_party:libtorch",
         "@local_config_cuda//cuda:cudart",
         "@local_config_tensorrt//:tensorrt",
     ],
@@ -93,6 +93,16 @@ cuda_library(
     ],
 )
 
+cuda_library(
+    name = "pfe_cuda",
+    srcs = ["pfe_cuda.cu"],
+    hdrs = ["pfe_cuda.h"],
+    deps = [
+        ":common",
+        "@local_config_cuda//cuda:cudart",
+    ],
+)
+
 cuda_library(
     name = "postprocess_cuda",
     srcs = ["postprocess_cuda.cu"],

modules/perception/lidar/lib/detection/lidar_point_pillars/params.h

@@ -37,6 +37,7 @@ class Params {
   static constexpr int kMaxNumPillars = 30000;
   static constexpr int kMaxNumPointsPerPillar = 60;
   static constexpr int kNumPointFeature = 5;  // x, y, z, i, delta of time
+  static constexpr int kNumGatherPointFeature = 9;
   static constexpr int kNumAnchor = 200 * 140 * 8 + 220 * 140 * 8;
   static constexpr int kNumOutputBoxFeature = 7;
   static constexpr int kBatchSize = 1;

modules/perception/lidar/lib/detection/lidar_point_pillars/pfe_cuda.cu

@@ -0,0 +1,99 @@
+/******************************************************************************
+ * Copyright 2020 The Apollo 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 "modules/perception/lidar/lib/detection/lidar_point_pillars/pfe_cuda.h"
+
+namespace apollo {
+namespace perception {
+namespace lidar {
+
+__global__ void gather_point_feature_kernel(
+    float* dev_pillar_point_feature, float* dev_num_points_per_pillar,
+    float* dev_pillar_coors, float* dev_pfe_gather_feature,
+    int max_num_points_per_pillar, int num_point_feature,
+    int num_gather_point_feature, float pillar_x_size, float pillar_y_size,
+    float min_x_range, float min_y_range) {
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
+  float point_mean[3];
+  for (int point_id = 0; point_id < dev_num_points_per_pillar[tid];
+       ++point_id) {
+    for (int i = 0; i < 3; ++i) {
+      int feature_id = i + point_id * num_point_feature +
+                       tid * max_num_points_per_pillar * num_point_feature;
+      point_mean[i] += dev_pillar_point_feature[feature_id];
+    }
+  }
+  for (int i = 0; i < 3; ++i) {
+    point_mean[i] = point_mean[i] / dev_num_points_per_pillar[tid];
+  }
+
+  float x_offset = pillar_x_size / 2 + min_x_range;
+  float y_offset = pillar_y_size / 2 + min_y_range;
+  for (int point_id = 0; point_id < dev_num_points_per_pillar[tid];
+       ++point_id) {
+    int point_head_id = point_id * num_point_feature +
+                        tid * max_num_points_per_pillar * num_point_feature;
+    int pfe_gather_head_id = point_id * num_gather_point_feature +
+        tid * max_num_points_per_pillar * num_gather_point_feature;
+    float point_x = dev_pillar_point_feature[point_head_id];
+    float point_y = dev_pillar_point_feature[point_head_id + 1];
+    dev_pfe_gather_feature[pfe_gather_head_id] =
+        sqrt(point_x * point_x + point_y * point_y);
+    for (int i = 2; i < num_point_feature; ++i) {
+      dev_pfe_gather_feature[pfe_gather_head_id + i - 1] =
+          dev_pillar_point_feature[point_head_id + i];
+    }
+    for (int i = 4; i < 7; ++i) {
+      dev_pfe_gather_feature[pfe_gather_head_id + i] =
+          dev_pillar_point_feature[point_head_id + i - 4] - point_mean[i - 4];
+    }
+    dev_pfe_gather_feature[pfe_gather_head_id + 7] =
+        dev_pillar_point_feature[point_head_id] -
+        (dev_pillar_coors[tid * 4 + 3] * pillar_x_size + x_offset);
+    dev_pfe_gather_feature[pfe_gather_head_id + 8] =
+        dev_pillar_point_feature[point_head_id + 1] -
+        (dev_pillar_coors[tid * 4 + 2] * pillar_y_size + y_offset);
+  }
+}
+
+PfeCuda::PfeCuda(int max_num_pillars, int max_num_points_per_pillar,
+                 int num_point_feature, int num_gather_point_feature,
+                 float pillar_x_size, float pillar_y_size, float min_x_range,
+                 float min_y_range, int num_threads)
+    : max_num_pillars_(max_num_pillars),
+      max_num_points_per_pillar_(max_num_points_per_pillar),
+      num_point_feature_(num_point_feature),
+      num_gather_point_feature_(num_gather_point_feature),
+      pillar_x_size_(pillar_x_size),
+      pillar_y_size_(pillar_y_size),
+      min_x_range_(min_x_range),
+      min_y_range_(min_y_range),
+      num_threads_(num_threads) {}
+
+void PfeCuda::GatherPointFeature(float* dev_pillar_point_feature,
+                                 float* dev_num_points_per_pillar,
+                                 float* dev_pillar_coors,
+                                 float* dev_pfe_gather_feature) {
+  const int num_blocks = DIVUP(max_num_pillars_, num_threads_);
+  gather_point_feature_kernel<<<num_blocks, num_threads_>>>(
+      dev_pillar_point_feature, dev_num_points_per_pillar, dev_pillar_coors,
+      dev_pfe_gather_feature, max_num_points_per_pillar_, num_point_feature_,
+      num_gather_point_feature_, pillar_x_size_, pillar_y_size_, min_x_range_,
+      min_y_range_);
+}
+}  // namespace lidar
+}  // namespace perception
+}  // namespace apollo

modules/perception/lidar/lib/detection/lidar_point_pillars/pfe_cuda.h

@@ -0,0 +1,51 @@
+/******************************************************************************
+ * Copyright 2020 The Apollo 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.
+ *****************************************************************************/
+
+#pragma once
+
+#include "modules/perception/lidar/lib/detection/lidar_point_pillars/common.h"
+
+namespace apollo {
+namespace perception {
+namespace lidar {
+
+class PfeCuda {
+ public:
+  PfeCuda(int max_num_pillars, int max_num_points_per_pillar,
+          int num_point_feature, int num_gather_point_feature,
+          float pillar_x_size, float pillar_y_size, float min_x_range,
+          float min_y_range, int num_threads);
+  ~PfeCuda() = default;
+
+  void GatherPointFeature(float* dev_pillar_point_feature,
+                          float* dev_num_points_per_pillar,
+                          float* dev_pillar_coors,
+                          float* dev_pfe_gather_feature);
+
+ private:
+  int max_num_pillars_;
+  int max_num_points_per_pillar_;
+  int num_point_feature_;
+  int num_gather_point_feature_;
+  float pillar_x_size_;
+  float pillar_y_size_;
+  float min_x_range_;
+  float min_y_range_;
+  int num_threads_;
+};
+}  // namespace lidar
+}  // namespace perception
+}  // namespace apollo