Skip to content

Commit

Permalink
Perception: allow customizing rotation angles of an anchor
Browse files Browse the repository at this point in the history
  • Loading branch information
@jeroldchen @jinghaomiao
jeroldchen authored and jinghaomiao committed May 18, 2020
1 parent 48604ec commit a95004b
Show file tree
Hide file tree
Showing 6 changed files with 58 additions and 47 deletions.
7 changes: 3 additions & 4 deletions modules/perception/lidar/lib/detection/lidar_point_pillars/anchor_mask_cuda.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -178,7 +178,7 @@ AnchorMaskCuda::AnchorMaskCuda(const int num_inds_for_scan,
const int num_anchor_x_inds,
const int num_anchor_y_inds,
const int num_class,
const int num_anchor_r_inds,
const int num_anchor_per_loc,
const float min_x_range,
const float min_y_range,
const float pillar_x_size,
Expand All @@ -189,7 +189,7 @@ AnchorMaskCuda::AnchorMaskCuda(const int num_inds_for_scan,
num_anchor_x_inds_(num_anchor_x_inds),
num_anchor_y_inds_(num_anchor_y_inds),
num_class_(num_class),
num_anchor_r_inds_(num_anchor_r_inds),
num_anchor_per_loc_(num_anchor_per_loc),
min_x_range_(min_x_range),
min_y_range_(min_y_range),
pillar_x_size_(pillar_x_size),
Expand All @@ -211,8 +211,7 @@ void AnchorMaskCuda::DoAnchorMaskCuda(
GPU_CHECK(cudaMemcpy(dev_sparse_pillar_map, dev_cumsum_along_y,
num_inds_for_scan_ * num_inds_for_scan_ * sizeof(int),
cudaMemcpyDeviceToDevice));
make_anchor_mask_kernel<<<num_anchor_x_inds_ * num_class_ *
num_anchor_r_inds_,
make_anchor_mask_kernel<<<num_anchor_x_inds_ * num_anchor_per_loc_,
num_anchor_y_inds_>>>(
dev_box_anchors_min_x, dev_box_anchors_min_y, dev_box_anchors_max_x,
dev_box_anchors_max_y, dev_sparse_pillar_map, dev_anchor_mask,
Expand Down
10 changes: 5 additions & 5 deletions modules/perception/lidar/lib/detection/lidar_point_pillars/anchor_mask_cuda.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,7 @@ class AnchorMaskCuda {
const int num_anchor_x_inds_;
const int num_anchor_y_inds_;
const int num_class_;
const int num_anchor_r_inds_;
const int num_anchor_per_loc_;
const float min_x_range_;
const float min_y_range_;
const float pillar_x_size_;
Expand All @@ -64,9 +64,9 @@ class AnchorMaskCuda {
* @param[in] num_anchor_x_inds Number of x-indexes for anchors
* @param[in] num_anchor_y_inds Number of y-indexes for anchors
* @param[in] num_class Number of classes
* @param[in] num_anchor_r_inds Number of rotation-indexes for anchors
* @param[in] min_x_range Minimum x value for pointcloud
* @param[in] min_y_range Minimum y value for pointcloud
* @param[in] num_anchor_per_loc Number of anchors per location
* @param[in] min_x_range Minimum x value for point cloud
* @param[in] min_y_range Minimum y value for point cloud
* @param[in] pillar_x_size Size of x-dimension for a pillar
* @param[in] pillar_y_size Size of y-dimension for a pillar
* @param[in] grid_x_size Number of pillars in x-coordinate
Expand All @@ -75,7 +75,7 @@ class AnchorMaskCuda {
*/
AnchorMaskCuda(const int num_inds_for_scan, const int num_anchor_x_inds,
const int num_anchor_y_inds, const int num_class,
const int num_anchor_r_inds, const float min_x_range,
const int num_anchor_per_loc, const float min_x_range,
const float min_y_range, const float pillar_x_size,
const float pillar_y_size, const int grid_x_size,
const int grid_y_size);
Expand Down
71 changes: 40 additions & 31 deletions modules/perception/lidar/lib/detection/lidar_point_pillars/point_pillars.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -53,9 +53,10 @@ const int PointPillars::kGridZSize = 1;
const int PointPillars::kRpnInputSize = 64 * kGridXSize * kGridYSize;
const int PointPillars::kNumAnchorXInds = kGridXSize * 0.5;
const int PointPillars::kNumAnchorYInds = kGridYSize * 0.5;
const int PointPillars::kNumAnchorRInds = 2;
const int PointPillars::kNumAnchorPerLoc = 6;
// kNumAnchorPerLoc need to be equal to length of kAnchorRo
const int PointPillars::kNumAnchor =
kNumAnchorXInds * kNumAnchorYInds * kNumClass * kNumAnchorRInds;
kNumAnchorXInds * kNumAnchorYInds * kNumAnchorPerLoc;
const int PointPillars::kNumOutputBoxFeature = 7;
const int PointPillars::kRpnBoxOutputSize = kNumAnchor * kNumOutputBoxFeature;
const int PointPillars::kRpnClsOutputSize = kNumAnchor * kNumClass;
Expand All @@ -81,6 +82,11 @@ const float PointPillars::kSensorHeight = 1.73f;
const std::vector<float> PointPillars::kAnchorDxSizes{1.6f, 0.6f, 0.6f};
const std::vector<float> PointPillars::kAnchorDySizes{3.9f, 1.76f, 0.8f};
const std::vector<float> PointPillars::kAnchorDzSizes{1.56f, 1.73f, 1.73f};
const std::vector<int> PointPillars::kNumAnchorRo{2, 2, 2};
// kNumAnchorRo's length equal to kNumClass
const std::vector<float> PointPillars::kAnchorRo{0, M_PI/2,
0, M_PI/2,
0, M_PI/2};

PointPillars::PointPillars(const bool reproduce_result_mode,
const float score_threshold,
Expand All @@ -107,7 +113,7 @@ PointPillars::PointPillars(const bool reproduce_result_mode,

anchor_mask_cuda_ptr_.reset(
new AnchorMaskCuda(kNumIndsForScan, kNumAnchorXInds, kNumAnchorYInds,
kNumClass, kNumAnchorRInds, kMinXRange, kMinYRange,
kNumClass, kNumAnchorPerLoc, kMinXRange, kMinYRange,
kPillarXSize, kPillarYSize, kGridXSize, kGridYSize));

scatter_cuda_ptr_.reset(
Expand All @@ -117,7 +123,7 @@ PointPillars::PointPillars(const bool reproduce_result_mode,
const float float_max = std::numeric_limits<float>::max();
postprocess_cuda_ptr_.reset(new PostprocessCuda(
float_min, float_max, kNumAnchorXInds, kNumAnchorYInds, kNumClass,
kNumAnchorRInds, score_threshold_, kNumThreads,
kNumAnchorPerLoc, score_threshold_, kNumThreads,
nms_overlap_threshold_, kNumBoxCorners, kNumOutputBoxFeature));

DeviceMemoryMalloc();
Expand Down Expand Up @@ -328,23 +334,21 @@ void PointPillars::GenerateAnchors(float* anchors_px_, float* anchors_py_,
anchor_y_count[i] = static_cast<float>(i) * y_stride + y_offset;
}

float anchor_r_count[kNumAnchorRInds];
anchor_r_count[0] = 0;
anchor_r_count[1] = M_PI / 2;

for (int y = 0; y < kNumAnchorYInds; ++y) {
for (int x = 0; x < kNumAnchorXInds; ++x) {
int ro_count = 0;
for (int c = 0; c < kNumClass; ++c) {
for (int r = 0; r < kNumAnchorRInds; ++r) {
int ind = y * kNumAnchorXInds * kNumClass * kNumAnchorRInds +
x * kNumClass * kNumAnchorRInds + c * kNumAnchorRInds + r;
for (int i = 0; i < kNumAnchorRo[c]; ++i) {
int ind = y * kNumAnchorXInds * kNumAnchorPerLoc +
x * kNumAnchorPerLoc + ro_count;
anchors_px_[ind] = anchor_x_count[x];
anchors_py_[ind] = anchor_y_count[y];
anchors_ro_[ind] = anchor_r_count[r];
anchors_ro_[ind] = kAnchorRo[ro_count];
anchors_pz_[ind] = -1 * kSensorHeight;
anchors_dx_[ind] = kAnchorDxSizes[c];
anchors_dy_[ind] = kAnchorDySizes[c];
anchors_dz_[ind] = kAnchorDzSizes[c];
ro_count++;
}
}
}
Expand Down Expand Up @@ -390,31 +394,36 @@ void PointPillars::ConvertAnchors2BoxAnchors(float* anchors_px,
memset(flipped_anchors_dy, 0, kNumAnchor * sizeof(float));
for (int x = 0; x < kNumAnchorXInds; ++x) {
for (int y = 0; y < kNumAnchorYInds; ++y) {
int ro_count = 0;
for (int c = 0; c < kNumClass; ++c) {
int base_ind = x * kNumAnchorYInds * kNumClass * kNumAnchorRInds +
y * kNumClass * kNumAnchorRInds + c * kNumAnchorRInds;
flipped_anchors_dx[base_ind + 0] = kAnchorDxSizes[c];
flipped_anchors_dy[base_ind + 0] = kAnchorDySizes[c];
flipped_anchors_dx[base_ind + 1] = kAnchorDySizes[c];
flipped_anchors_dy[base_ind + 1] = kAnchorDxSizes[c];
for (int i = 0; i < kNumAnchorRo[c]; ++i) {
int base_ind = x * kNumAnchorYInds * kNumAnchorPerLoc +
y * kNumAnchorPerLoc + ro_count;
if (kAnchorRo[ro_count] <= 0.78) {
flipped_anchors_dx[base_ind] = kAnchorDxSizes[c];
flipped_anchors_dy[base_ind] = kAnchorDySizes[c];
} else {
flipped_anchors_dx[base_ind] = kAnchorDySizes[c];
flipped_anchors_dy[base_ind] = kAnchorDxSizes[c];
}
ro_count++;
}
}
}
}
for (int x = 0; x < kNumAnchorXInds; ++x) {
for (int y = 0; y < kNumAnchorYInds; ++y) {
for (int c = 0; c < kNumClass; ++c) {
for (int r = 0; r < kNumAnchorRInds; ++r) {
int ind = x * kNumAnchorYInds * kNumClass * kNumAnchorRInds +
y * kNumClass * kNumAnchorRInds + c * kNumAnchorRInds + r;
box_anchors_min_x_[ind] =
anchors_px[ind] - flipped_anchors_dx[ind] / 2.0f;
box_anchors_min_y_[ind] =
anchors_py[ind] - flipped_anchors_dy[ind] / 2.0f;
box_anchors_max_x_[ind] =
anchors_px[ind] + flipped_anchors_dx[ind] / 2.0f;
box_anchors_max_y_[ind] =
anchors_py[ind] + flipped_anchors_dy[ind] / 2.0f;
}
for (int i = 0; i < kNumAnchorPerLoc; ++i) {
int ind = x * kNumAnchorYInds * kNumAnchorPerLoc +
y * kNumAnchorPerLoc + i;
box_anchors_min_x_[ind] =
anchors_px[ind] - flipped_anchors_dx[ind] / 2.0f;
box_anchors_min_y_[ind] =
anchors_py[ind] - flipped_anchors_dy[ind] / 2.0f;
box_anchors_max_x_[ind] =
anchors_px[ind] + flipped_anchors_dx[ind] / 2.0f;
box_anchors_max_y_[ind] =
anchors_py[ind] + flipped_anchors_dy[ind] / 2.0f;
}
}
}
Expand Down
5 changes: 4 additions & 1 deletion modules/perception/lidar/lib/detection/lidar_point_pillars/point_pillars.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -112,7 +112,8 @@ class PointPillars {
static const int kRpnInputSize;
static const int kNumAnchorXInds;
static const int kNumAnchorYInds;
static const int kNumAnchorRInds;
// static const int kNumAnchorRInds;
static const int kNumAnchorPerLoc;
static const int kNumAnchor;
static const int kNumOutputBoxFeature;
static const int kRpnBoxOutputSize;
Expand All @@ -137,6 +138,8 @@ class PointPillars {
static const std::vector<float> kAnchorDxSizes;
static const std::vector<float> kAnchorDySizes;
static const std::vector<float> kAnchorDzSizes;
static const std::vector<int> kNumAnchorRo;
static const std::vector<float> kAnchorRo;

// initialize in initializer list
const bool reproduce_result_mode_;
Expand Down
6 changes: 3 additions & 3 deletions modules/perception/lidar/lib/detection/lidar_point_pillars/postprocess_cuda.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -187,15 +187,15 @@ __global__ void sort_boxes_by_indexes_kernel(
PostprocessCuda::PostprocessCuda(
const float float_min, const float float_max, const int num_anchor_x_inds,
const int num_anchor_y_inds, const int num_class,
const int num_anchor_r_inds, const float score_threshold,
const int num_anchor_per_loc, const float score_threshold,
const int num_threads, const float nms_overlap_threshold,
const int num_box_corners, const int num_output_box_feature)
: float_min_(float_min),
float_max_(float_max),
num_anchor_x_inds_(num_anchor_x_inds),
num_anchor_y_inds_(num_anchor_y_inds),
num_class_(num_class),
num_anchor_r_inds_(num_anchor_r_inds),
num_anchor_per_loc_(num_anchor_per_loc),
score_threshold_(score_threshold),
num_threads_(num_threads),
nms_overlap_threshold_(nms_overlap_threshold),
Expand All @@ -216,7 +216,7 @@ void PostprocessCuda::DoPostprocessCuda(
int* dev_filtered_dir, float* dev_box_for_nms,
int* dev_filter_count, std::vector<float>* out_detection,
std::vector<int>* out_label) {
filter_kernel<<<num_anchor_x_inds_ * num_class_ * num_anchor_r_inds_,
filter_kernel<<<num_anchor_x_inds_ * num_anchor_per_loc_,
num_anchor_y_inds_>>>(
rpn_box_output, rpn_cls_output, rpn_dir_output, dev_anchor_mask,
dev_anchors_px, dev_anchors_py, dev_anchors_pz, dev_anchors_dx,
Expand Down
6 changes: 3 additions & 3 deletions modules/perception/lidar/lib/detection/lidar_point_pillars/postprocess_cuda.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,7 @@ class PostprocessCuda {
const int num_anchor_x_inds_;
const int num_anchor_y_inds_;
const int num_class_;
const int num_anchor_r_inds_;
const int num_anchor_per_loc_;
const float score_threshold_;
const int num_threads_;
const float nms_overlap_threshold_;
Expand All @@ -75,7 +75,7 @@ class PostprocessCuda {
* @param[in] num_anchor_x_inds Number of x-indexes for anchors
* @param[in] num_anchor_y_inds Number of y-indexes for anchors
* @param[in] num_class Number of object's classes
* @param[in] num_anchor_r_inds Number of rotation-indexes for anchors
* @param[in] num_anchor_per_loc Number of anchors per location
* @param[in] score_threshold Score threshold for filtering output
* @param[in] num_threads Number of threads when launching cuda kernel
* @param[in] nms_overlap_threshold IOU threshold for NMS
Expand All @@ -86,7 +86,7 @@ class PostprocessCuda {
*/
PostprocessCuda(const float float_min, const float float_max,
const int num_anchor_x_inds, const int num_anchor_y_inds,
const int num_class, const int num_anchor_r_inds,
const int num_class, const int num_anchor_per_loc,
const float score_threshold, const int num_threads,
const float nms_overlap_threshold, const int num_box_corners,
const int num_output_box_feature);
Expand Down

0 comments on commit a95004b

Please sign in to comment.