Skip to content

Commit

Permalink
Perception: upgrade onnx models and adapt for newer offline training …
Browse files Browse the repository at this point in the history
…program (ApolloAuto#11248)
  • Loading branch information
jeroldchen authored May 15, 2020
1 parent 4289c34 commit a29ae24
Show file tree
Hide file tree
Showing 13 changed files with 660 additions and 639 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -80,7 +80,7 @@ __global__ void nms_kernel(const int n_boxes, const float nms_overlap_thresh,
if (row_start == col_start) {
start = threadIdx.x + 1;
}
for (int i = start; i < col_size; i++) {
for (int i = start; i < col_size; ++i) {
if (devIoU(cur_box, block_boxes + i * num_box_corners) >
nms_overlap_thresh) {
t |= 1ULL << i;
Expand Down Expand Up @@ -122,14 +122,14 @@ void NmsCuda::DoNmsCuda(const int host_filter_count,
std::vector<uint64_t> remv(col_blocks);
memset(&remv[0], 0, sizeof(uint64_t) * col_blocks);

for (int i = 0; i < host_filter_count; i++) {
for (int i = 0; i < host_filter_count; ++i) {
int nblock = i / num_threads_;
int inblock = i % num_threads_;

if (!(remv[nblock] & (1ULL << inblock))) {
out_keep_inds[(*out_num_to_keep)++] = i;
uint64_t *p = &host_mask[0] + i * col_blocks;
for (int j = nblock; j < col_blocks; j++) {
for (int j = nblock; j < col_blocks; ++j) {
remv[j] |= p[j];
}
}
Expand Down

Large diffs are not rendered by default.

Original file line number Diff line number Diff line change
Expand Up @@ -104,6 +104,7 @@ class PointPillars {
static const int kNumClass;
static const int kMaxNumPillars;
static const int kMaxNumPointsPerPillar;
static const int kNumPointFeature;
static const int kPfeOutputSize;
static const int kGridXSize;
static const int kGridYSize;
Expand Down Expand Up @@ -174,22 +175,16 @@ class PointPillars {
int* dev_cumsum_along_x_;
int* dev_cumsum_along_y_;

float* dev_pillar_x_;
float* dev_pillar_y_;
float* dev_pillar_z_;
float* dev_pillar_i_;

float* dev_x_coors_for_sub_shaped_;
float* dev_y_coors_for_sub_shaped_;
float* dev_pillar_feature_mask_;
float* dev_pillar_point_feature_;
float* dev_pillar_coors_;

float* dev_box_anchors_min_x_;
float* dev_box_anchors_min_y_;
float* dev_box_anchors_max_x_;
float* dev_box_anchors_max_y_;
int* dev_anchor_mask_;

void* pfe_buffers_[9];
void* pfe_buffers_[4];
void* rpn_buffers_[4];

float* dev_scattered_feature_;
Expand Down Expand Up @@ -274,15 +269,15 @@ class PointPillars {

/**
* @brief Preproces points
* @param[in] in_points_array pointcloud array
* @param[in] in_points_array Point cloud array
* @param[in] in_num_points Number of points
* @details Call CPU or GPU preprocess
*/
void Preprocess(const float* in_points_array, const int in_num_points);

/**
* @brief Preproces by CPU
* @param[in] in_points_array pointcloud array
* @param[in] in_points_array Point cloud array
* @param[in] in_num_points Number of points
* @details The output from preprocessCPU is reproducible, while preprocessGPU
* is not
Expand All @@ -291,34 +286,29 @@ class PointPillars {

/**
* @brief Preproces by GPU
* @param[in] in_points_array pointcloud array
* @param[in] in_points_array Point cloud array
* @param[in] in_num_points Number of points
* @details Faster preprocess comapared with CPU preprocess
* @details Faster preprocess compared with CPU preprocess
*/
void PreprocessGPU(const float* in_points_array, const int in_num_points);

/**
* @brief Convert anchors to box form like min_x, min_y, max_x, max_y anchors
* @param[in] anchors_px_ Represents x-coordinate value for a corresponding
* anchor
* @param[in] anchors_py_ Represents y-coordinate value for a corresponding
* anchor
* @param[in] anchors_dx_ Represents x-dimension value for a corresponding
* anchor
* @param[in] anchors_dy_ Represents y-dimension value for a corresponding
* anchor
* @param[in] box_anchors_min_x_ Represents minimum x value for a
* correspomding anchor
* @param[in] box_anchors_min_y_ Represents minimum y value for a
* correspomding anchor
* @param[in] box_anchors_max_x_ Represents maximum x value for a
* correspomding anchor
* @param[in] box_anchors_max_y_ Represents maximum y value for a
* correspomding anchor
* @param[in] anchors_px_
* Represents x-coordinate value for a corresponding anchor
* @param[in] anchors_py_
* Represents y-coordinate value for a corresponding anchor
* @param[in] box_anchors_min_x_
* Represents minimum x value for a corresponding anchor
* @param[in] box_anchors_min_y_
* Represents minimum y value for a corresponding anchor
* @param[in] box_anchors_max_x_
* Represents maximum x value for a corresponding anchor
* @param[in] box_anchors_max_y_
* Represents maximum y value for a corresponding anchor
* @details Make box anchors for nms
*/
void ConvertAnchors2BoxAnchors(float* anchors_px_, float* anchors_py_,
float* anchors_dx_, float* anchors_dy_,
float* box_anchors_min_x_,
float* box_anchors_min_y_,
float* box_anchors_max_x_,
Expand All @@ -339,25 +329,23 @@ class PointPillars {
* @param[in] nms_overlap_threshold IOU threshold for NMS
* @param[in] pfe_onnx_file Pillar Feature Extractor ONNX file path
* @param[in] rpn_onnx_file Region Proposal Network ONNX file path
* @details Variables could be chaned through rosparam
* @details Variables could be changed through point_pillars_detection
*/
PointPillars(const bool reproduce_result_mode,
const float score_threshold,
PointPillars(const bool reproduce_result_mode, const float score_threshold,
const float nms_overlap_threshold,
const std::string pfe_onnx_file,
const std::string rpn_onnx_file);
~PointPillars();

/**
* @brief Call PointPillars for the inference
* @param[in] in_points_array Pointcloud array
* @param[in] in_points_array Point cloud array
* @param[in] in_num_points Number of points
* @param[out] out_detections Network output bounding box
* @param[out] out_labels Network output object's label
* @details This is an interface for the algorithm
*/
void DoInference(const float* in_points_array,
const int in_num_points,
void DoInference(const float* in_points_array, const int in_num_points,
std::vector<float>* out_detections,
std::vector<int>* out_labels);
};
Expand Down
Loading

0 comments on commit a29ae24

Please sign in to comment.