OpenPCDet复现pointpillars，win10

已于 2023-05-11 16:47:51 修改
阅读量1.7k 收藏 15
点赞数
文章标签： conda
于 2023-05-09 21:45:18 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/sinounuo/article/details/130537461
版权

一、安装OpenPCDet

安装虚拟环境：conda create -n OpenPCDet python == 3.9

打开虚拟环境：conda activate OpenPCDet 

单独安装torch：（用轮子装）

pip install torch==1.9.1+cu111 torchvision==0.10.1+cu111 torchaudio==0.9.1 -f https://download.pytorch.org/whl/torch_stable.html

OpenPCDet源码网站上可以看出，目前OpenPCDet支持的最高torch是1.10

 Git地址：

mirrors / open-mmlab / OpenPCDet · GitCode

但是torchvision==0.11.0+cu111只有Linux的版本，win10下没有，所以把torch再降一点：借鉴以下内容：

安装pytorch1.10.0/cu111时报错：no matching distribution found for torchvision==0.11.0+cu111_error: no matching distribution found for torchvis-CSDN博客

Clone OpenPCDet项目 ：在你自己指定的文件夹下，打开git bash here，输入

git clone https://github.com/open-mmlab/OpenPCDet.git
torch安装完成后，cd 到clone好的OpenPCDet目录下：pip install -r requirements.txt


出错：Could not build wheels for SharedArray，我查看了conda网站上的内容，sharedarray是只支持Linux系统，不支持Windows的。

  1. 首先在依赖文件（requirement.txt）中，SharedArray pip安装是会失败的，因为其默认版本不是windows版本：GitHub - imaginary-friend94/Shared-Array-for-Windows: Share numpy arrays between processes

 按照以下内容在win10上安装sharedarray：先把git上的代码clone下来，然后里面文件张这样：

 然后修改setup.py和shared_memory_python.cpp：参考：

在Windows上安装openPCDet遇到的一些问题总结_iou3d_nms_cuda_努力学pc的博客-CSDN博客

修改后，编译：python setup.py develop

在安装pcdet库之前需要对pcdet\ops\iou3d_nms\src里面的文件进行修改。

iou3d_nms.cpp修改之后：

/*
3D IoU Calculation and Rotated NMS(modified from 2D NMS written by others)
Written by Shaoshuai Shi
All Rights Reserved 2019-2020.
*/

#include <torch/serialize/tensor.h>
#include <torch/extension.h>
#include <vector>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <stdint.h>
#include "iou3d_nms.h"

#define CHECK_CUDA(x) do { \
  if (!x.type().is_cuda()) { \
    fprintf(stderr, "%s must be CUDA tensor at %s:%d\n", #x, __FILE__, __LINE__); \
    exit(-1); \
  } \
} while (0)
#define CHECK_CONTIGUOUS(x) do { \
  if (!x.is_contiguous()) { \
    fprintf(stderr, "%s must be contiguous tensor at %s:%d\n", #x, __FILE__, __LINE__); \
    exit(-1); \
  } \
} while (0)
#define CHECK_INPUT(x) CHECK_CUDA(x);CHECK_CONTIGUOUS(x)

#define DIVUP(m,n) ((m) / (n) + ((m) % (n) > 0))

#define CHECK_ERROR(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

const int THREADS_PER_BLOCK_NMS = sizeof(uint64_t) * 8;

void boxesoverlapLauncher(const int num_a, const float *boxes_a, const int num_b, const float *boxes_b, float *ans_overlap);
void boxesioubevLauncher(const int num_a, const float *boxes_a, const int num_b, const float *boxes_b, float *ans_iou);
void nmsLauncher(const float *boxes, uint64_t * mask, int boxes_num, float nms_overlap_thresh);
void nmsNormalLauncher(const float *boxes, uint64_t * mask, int boxes_num, float nms_overlap_thresh);


int boxes_overlap_bev_gpu(at::Tensor boxes_a, at::Tensor boxes_b, at::Tensor ans_overlap){
    // params boxes_a: (N, 7) [x, y, z, dx, dy, dz, heading]
    // params boxes_b: (M, 7) [x, y, z, dx, dy, dz, heading]
    // params ans_overlap: (N, M)

    CHECK_INPUT(boxes_a);
    CHECK_INPUT(boxes_b);
    CHECK_INPUT(ans_overlap);

    int num_a = boxes_a.size(0);
    int num_b = boxes_b.size(0);

    const float * boxes_a_data = boxes_a.data<float>();
    const float * boxes_b_data = boxes_b.data<float>();
    float * ans_overlap_data = ans_overlap.data<float>();

    boxesoverlapLauncher(num_a, boxes_a_data, num_b, boxes_b_data, ans_overlap_data);

    return 1;
}

int boxes_iou_bev_gpu(at::Tensor boxes_a, at::Tensor boxes_b, at::Tensor ans_iou){
    // params boxes_a: (N, 7) [x, y, z, dx, dy, dz, heading]
    // params boxes_b: (M, 7) [x, y, z, dx, dy, dz, heading]
    // params ans_overlap: (N, M)
    CHECK_INPUT(boxes_a);
    CHECK_INPUT(boxes_b);
    CHECK_INPUT(ans_iou);

    int num_a = boxes_a.size(0);
    int num_b = boxes_b.size(0);

    const float * boxes_a_data = boxes_a.data<float>();
    const float * boxes_b_data = boxes_b.data<float>();
    float * ans_iou_data = ans_iou.data<float>();

    boxesioubevLauncher(num_a, boxes_a_data, num_b, boxes_b_data, ans_iou_data);

    return 1;
}

int nms_gpu(at::Tensor boxes, at::Tensor keep, float nms_overlap_thresh){
    // params boxes: (N, 7) [x, y, z, dx, dy, dz, heading]
    // params keep: (N)
    CHECK_INPUT(boxes);
    CHECK_CONTIGUOUS(keep);

    int boxes_num = boxes.size(0);
    const float * boxes_data = boxes.data<float>();
    long long * keep_data = keep.data<long long>();

    const int col_blocks = DIVUP(boxes_num, THREADS_PER_BLOCK_NMS);

    uint64_t *mask_data = NULL;
    CHECK_ERROR(cudaMalloc((void**)&mask_data, boxes_num * col_blocks * sizeof(uint64_t)));
    nmsLauncher(boxes_data, mask_data, boxes_num, nms_overlap_thresh);

    // uint64_t mask_cpu[boxes_num * col_blocks];
    // uint64_t *mask_cpu = new uint64_t [boxes_num * col_blocks];
    std::vector<uint64_t> mask_cpu(boxes_num * col_blocks);

//    printf("boxes_num=%d, col_blocks=%d\n", boxes_num, col_blocks);
    CHECK_ERROR(cudaMemcpy(&mask_cpu[0], mask_data, boxes_num * col_blocks * sizeof(uint64_t),
                           cudaMemcpyDeviceToHost));

    cudaFree(mask_data);

    //uint64_t remv_cpu(col_blocks);
    std::vector<uint64_t> remv_cpu(col_blocks);
    memset(&remv_cpu[0], 0, col_blocks * sizeof(uint64_t));

    int num_to_keep = 0;

    for (int i = 0; i < boxes_num; i++){
        int nblock = i / THREADS_PER_BLOCK_NMS;
        int inblock = i % THREADS_PER_BLOCK_NMS;

        if (!(remv_cpu[nblock] & (1ULL << inblock))){
            keep_data[num_to_keep++] = i;
            uint64_t *p = &mask_cpu[0] + i * col_blocks;
            for (int j = nblock; j < col_blocks; j++){
                remv_cpu[j] |= p[j];
            }
        }
    }
    if ( cudaSuccess != cudaGetLastError() ) printf( "Error!\n" );

    return num_to_keep;
}


int nms_normal_gpu(at::Tensor boxes, at::Tensor keep, float nms_overlap_thresh){
    // params boxes: (N, 7) [x, y, z, dx, dy, dz, heading]
    // params keep: (N)

    CHECK_INPUT(boxes);
    CHECK_CONTIGUOUS(keep);

    int boxes_num = boxes.size(0);
    const float * boxes_data = boxes.data<float>();
    long long * keep_data = keep.data<long long>();

    const int col_blocks = DIVUP(boxes_num, THREADS_PER_BLOCK_NMS);

    uint64_t *mask_data = NULL;
    CHECK_ERROR(cudaMalloc((void**)&mask_data, boxes_num * col_blocks * sizeof(uint64_t)));
    nmsNormalLauncher(boxes_data, mask_data, boxes_num, nms_overlap_thresh);

    // uint64_t mask_cpu[boxes_num * col_blocks];
    // uint64_t *mask_cpu = new uint64_t [boxes_num * col_blocks];
    std::vector<uint64_t> mask_cpu(boxes_num * col_blocks);

//    printf("boxes_num=%d, col_blocks=%d\n", boxes_num, col_blocks);
    CHECK_ERROR(cudaMemcpy(&mask_cpu[0], mask_data, boxes_num * col_blocks * sizeof(uint64_t),
                           cudaMemcpyDeviceToHost));

    cudaFree(mask_data);

    std::vector<uint64_t> remv_cpu(col_blocks);
    memset(&remv_cpu[0], 0, col_blocks * sizeof(uint64_t));

    int num_to_keep = 0;

    for (int i = 0; i < boxes_num; i++){
        int nblock = i / THREADS_PER_BLOCK_NMS;
        int inblock = i % THREADS_PER_BLOCK_NMS;

        if (!(remv_cpu[nblock] & (1ULL << inblock))){
            keep_data[num_to_keep++] = i;
            uint64_t *p = &mask_cpu[0] + i * col_blocks;
            for (int j = nblock; j < col_blocks; j++){
                remv_cpu[j] |= p[j];
            }
        }
    }
    if ( cudaSuccess != cudaGetLastError() ) printf( "Error!\n" );

    return num_to_keep;
}
 

iou3d_nms_kernel.cu修改之后：

/*
3D IoU Calculation and Rotated NMS(modified from 2D NMS written by others)
Written by Shaoshuai Shi
All Rights Reserved 2019-2020.
*/


#include <stdio.h>
#include <stdint.h>
#define THREADS_PER_BLOCK 16
#define DIVUP(m, n) ((m) / (n) + ((m) % (n) > 0))

// #define DEBUG
const int THREADS_PER_BLOCK_NMS = sizeof(unsigned long long) * 8;
//const float EPS = 1e-8;
struct Point {
    float x, y;
    __device__ Point() {}
    __device__ Point(double _x, double _y){
        x = _x, y = _y;
    }

    __device__ void set(float _x, float _y){
        x = _x; y = _y;
    }

    __device__ Point operator +(const Point &b)const{
        return Point(x + b.x, y + b.y);
    }

    __device__ Point operator -(const Point &b)const{
        return Point(x - b.x, y - b.y);
    }
};

__device__ inline float cross(const Point &a, const Point &b){
    return a.x * b.y - a.y * b.x;
}

__device__ inline float cross(const Point &p1, const Point &p2, const Point &p0){
    return (p1.x - p0.x) * (p2.y - p0.y) - (p2.x - p0.x) * (p1.y - p0.y);
}

__device__ int check_rect_cross_cuda(const Point &p1, const Point &p2, const Point &q1, const Point &q2){
    int ret = min(p1.x,p2.x) <= max(q1.x,q2.x)  &&
              min(q1.x,q2.x) <= max(p1.x,p2.x) &&
              min(p1.y,p2.y) <= max(q1.y,q2.y) &&
              min(q1.y,q2.y) <= max(p1.y,p2.y);
    return ret;
}

__device__ inline int check_in_box2d(const float *box, const Point &p){
    //params: (7) [x, y, z, dx, dy, dz, heading]
    const float MARGIN = 1e-2;

    float center_x = box[0], center_y = box[1];
    float angle_cos = cos(-box[6]), angle_sin = sin(-box[6]);  // rotate the point in the opposite direction of box
    float rot_x = (p.x - center_x) * angle_cos + (p.y - center_y) * (-angle_sin);
    float rot_y = (p.x - center_x) * angle_sin + (p.y - center_y) * angle_cos;

    return (fabs(rot_x) < box[3] / 2 + MARGIN && fabs(rot_y) < box[4] / 2 + MARGIN);
}

__device__ inline int intersection(const Point &p1, const Point &p0, const Point &q1, const Point &q0, Point &ans){
    // fast exclusion
    if (check_rect_cross_cuda(p0, p1, q0, q1) == 0) return 0;

    // check cross standing
    float s1 = cross(q0, p1, p0);
    float s2 = cross(p1, q1, p0);
    float s3 = cross(p0, q1, q0);
    float s4 = cross(q1, p1, q0);

    if (!(s1 * s2 > 0 && s3 * s4 > 0)) return 0;

    // calculate intersection of two lines
    float s5 = cross(q1, p1, p0);
    if(fabs(s5 - s1) > 1e-8){
        ans.x = (s5 * q0.x - s1 * q1.x) / (s5 - s1);
        ans.y = (s5 * q0.y - s1 * q1.y) / (s5 - s1);

    }
    else{
        float a0 = p0.y - p1.y, b0 = p1.x - p0.x, c0 = p0.x * p1.y - p1.x * p0.y;
        float a1 = q0.y - q1.y, b1 = q1.x - q0.x, c1 = q0.x * q1.y - q1.x * q0.y;
        float D = a0 * b1 - a1 * b0;

        ans.x = (b0 * c1 - b1 * c0) / D;
        ans.y = (a1 * c0 - a0 * c1) / D;
    }

    return 1;
}

__device__ inline void rotate_around_center(const Point &center, const float angle_cos, const float angle_sin, Point &p){
    float new_x = (p.x - center.x) * angle_cos + (p.y - center.y) * (-angle_sin) + center.x;
    float new_y = (p.x - center.x) * angle_sin + (p.y - center.y) * angle_cos + center.y;
    p.set(new_x, new_y);
}

__device__ inline int point_cmp(const Point &a, const Point &b, const Point &center){
    return atan2(a.y - center.y, a.x - center.x) > atan2(b.y - center.y, b.x - center.x);
}

__device__ inline float box_overlap(const float *box_a, const float *box_b){
    // params box_a: [x, y, z, dx, dy, dz, heading]
    // params box_b: [x, y, z, dx, dy, dz, heading]

    float a_angle = box_a[6], b_angle = box_b[6];
    float a_dx_half = box_a[3] / 2, b_dx_half = box_b[3] / 2, a_dy_half = box_a[4] / 2, b_dy_half = box_b[4] / 2;
    float a_x1 = box_a[0] - a_dx_half, a_y1 = box_a[1] - a_dy_half;
    float a_x2 = box_a[0] + a_dx_half, a_y2 = box_a[1] + a_dy_half;
    float b_x1 = box_b[0] - b_dx_half, b_y1 = box_b[1] - b_dy_half;
    float b_x2 = box_b[0] + b_dx_half, b_y2 = box_b[1] + b_dy_half;

    Point center_a(box_a[0], box_a[1]);
    Point center_b(box_b[0], box_b[1]);

#ifdef DEBUG
    printf("a: (%.3f, %.3f, %.3f, %.3f, %.3f), b: (%.3f, %.3f, %.3f, %.3f, %.3f)\n", a_x1, a_y1, a_x2, a_y2, a_angle,
           b_x1, b_y1, b_x2, b_y2, b_angle);
    printf("center a: (%.3f, %.3f), b: (%.3f, %.3f)\n", center_a.x, center_a.y, center_b.x, center_b.y);
#endif

    Point box_a_corners[5];
    box_a_corners[0].set(a_x1, a_y1);
    box_a_corners[1].set(a_x2, a_y1);
    box_a_corners[2].set(a_x2, a_y2);
    box_a_corners[3].set(a_x1, a_y2);

    Point box_b_corners[5];
    box_b_corners[0].set(b_x1, b_y1);
    box_b_corners[1].set(b_x2, b_y1);
    box_b_corners[2].set(b_x2, b_y2);
    box_b_corners[3].set(b_x1, b_y2);

    // get oriented corners
    float a_angle_cos = cos(a_angle), a_angle_sin = sin(a_angle);
    float b_angle_cos = cos(b_angle), b_angle_sin = sin(b_angle);

    for (int k = 0; k < 4; k++){
#ifdef DEBUG
        printf("before corner %d: a(%.3f, %.3f), b(%.3f, %.3f) \n", k, box_a_corners[k].x, box_a_corners[k].y, box_b_corners[k].x, box_b_corners[k].y);
#endif
        rotate_around_center(center_a, a_angle_cos, a_angle_sin, box_a_corners[k]);
        rotate_around_center(center_b, b_angle_cos, b_angle_sin, box_b_corners[k]);
#ifdef DEBUG
        printf("corner %d: a(%.3f, %.3f), b(%.3f, %.3f) \n", k, box_a_corners[k].x, box_a_corners[k].y, box_b_corners[k].x, box_b_corners[k].y);
#endif
    }

    box_a_corners[4] = box_a_corners[0];
    box_b_corners[4] = box_b_corners[0];

    // get intersection of lines
    Point cross_points[16];
    Point poly_center;
    int cnt = 0, flag = 0;

    poly_center.set(0, 0);
    for (int i = 0; i < 4; i++){
        for (int j = 0; j < 4; j++){
            flag = intersection(box_a_corners[i + 1], box_a_corners[i], box_b_corners[j + 1], box_b_corners[j], cross_points[cnt]);
            if (flag){
                poly_center = poly_center + cross_points[cnt];
                cnt++;
#ifdef DEBUG
                printf("Cross points (%.3f, %.3f): a(%.3f, %.3f)->(%.3f, %.3f), b(%.3f, %.3f)->(%.3f, %.3f) \n",
                    cross_points[cnt - 1].x, cross_points[cnt - 1].y,
                    box_a_corners[i].x, box_a_corners[i].y, box_a_corners[i + 1].x, box_a_corners[i + 1].y,
                    box_b_corners[i].x, box_b_corners[i].y, box_b_corners[i + 1].x, box_b_corners[i + 1].y);
#endif
            }
        }
    }

    // check corners
    for (int k = 0; k < 4; k++){
        if (check_in_box2d(box_a, box_b_corners[k])){
            poly_center = poly_center + box_b_corners[k];
            cross_points[cnt] = box_b_corners[k];
            cnt++;
#ifdef DEBUG
                printf("b corners in a: corner_b(%.3f, %.3f)", cross_points[cnt - 1].x, cross_points[cnt - 1].y);
#endif
        }
        if (check_in_box2d(box_b, box_a_corners[k])){
            poly_center = poly_center + box_a_corners[k];
            cross_points[cnt] = box_a_corners[k];
            cnt++;
#ifdef DEBUG
                printf("a corners in b: corner_a(%.3f, %.3f)", cross_points[cnt - 1].x, cross_points[cnt - 1].y);
#endif
        }
    }

    poly_center.x /= cnt;
    poly_center.y /= cnt;

    // sort the points of polygon
    Point temp;
    for (int j = 0; j < cnt - 1; j++){
        for (int i = 0; i < cnt - j - 1; i++){
            if (point_cmp(cross_points[i], cross_points[i + 1], poly_center)){
                temp = cross_points[i];
                cross_points[i] = cross_points[i + 1];
                cross_points[i + 1] = temp;
            }
        }
    }

#ifdef DEBUG
    printf("cnt=%d\n", cnt);
    for (int i = 0; i < cnt; i++){
        printf("All cross point %d: (%.3f, %.3f)\n", i, cross_points[i].x, cross_points[i].y);
    }
#endif

    // get the overlap areas
    float area = 0;
    for (int k = 0; k < cnt - 1; k++){
        area += cross(cross_points[k] - cross_points[0], cross_points[k + 1] - cross_points[0]);
    }

    return fabs(area) / 2.0;
}

__device__ inline float iou_bev(const float *box_a, const float *box_b){
    // params box_a: [x, y, z, dx, dy, dz, heading]
    // params box_b: [x, y, z, dx, dy, dz, heading]
    float sa = box_a[3] * box_a[4];
    float sb = box_b[3] * box_b[4];
    float s_overlap = box_overlap(box_a, box_b);
    return s_overlap / fmaxf(sa + sb - s_overlap, 1e-8);
}

__global__ void boxes_overlap_kernel(const int num_a, const float *boxes_a, const int num_b, const float *boxes_b, float *ans_overlap){
    // params boxes_a: (N, 7) [x, y, z, dx, dy, dz, heading]
    // params boxes_b: (M, 7) [x, y, z, dx, dy, dz, heading]
    const int a_idx = blockIdx.y * THREADS_PER_BLOCK + threadIdx.y;
    const int b_idx = blockIdx.x * THREADS_PER_BLOCK + threadIdx.x;

    if (a_idx >= num_a || b_idx >= num_b){
        return;
    }
    const float * cur_box_a = boxes_a + a_idx * 7;
    const float * cur_box_b = boxes_b + b_idx * 7;
    float s_overlap = box_overlap(cur_box_a, cur_box_b);
    ans_overlap[a_idx * num_b + b_idx] = s_overlap;
}

__global__ void boxes_iou_bev_kernel(const int num_a, const float *boxes_a, const int num_b, const float *boxes_b, float *ans_iou){
    // params boxes_a: (N, 7) [x, y, z, dx, dy, dz, heading]
    // params boxes_b: (M, 7) [x, y, z, dx, dy, dz, heading]
    const int a_idx = blockIdx.y * THREADS_PER_BLOCK + threadIdx.y;
    const int b_idx = blockIdx.x * THREADS_PER_BLOCK + threadIdx.x;

    if (a_idx >= num_a || b_idx >= num_b){
        return;
    }

    const float * cur_box_a = boxes_a + a_idx * 7;
    const float * cur_box_b = boxes_b + b_idx * 7;
    float cur_iou_bev = iou_bev(cur_box_a, cur_box_b);
    ans_iou[a_idx * num_b + b_idx] = cur_iou_bev;
}

__global__ void nms_kernel(const int boxes_num, const float nms_overlap_thresh,
                           const float *boxes, uint64_t*mask){
    //params: boxes (N, 7) [x, y, z, dx, dy, dz, heading]
    //params: mask (N, N/THREADS_PER_BLOCK_NMS)

    const int row_start = blockIdx.y;
    const int col_start = blockIdx.x;

    // if (row_start > col_start) return;

    const int row_size = fminf(boxes_num - row_start * THREADS_PER_BLOCK_NMS, THREADS_PER_BLOCK_NMS);
    const int col_size = fminf(boxes_num - col_start * THREADS_PER_BLOCK_NMS, THREADS_PER_BLOCK_NMS);

    __shared__ float block_boxes[THREADS_PER_BLOCK_NMS * 7];

    if (threadIdx.x < col_size) {
        block_boxes[threadIdx.x * 7 + 0] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 0];
        block_boxes[threadIdx.x * 7 + 1] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 1];
        block_boxes[threadIdx.x * 7 + 2] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 2];
        block_boxes[threadIdx.x * 7 + 3] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 3];
        block_boxes[threadIdx.x * 7 + 4] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 4];
        block_boxes[threadIdx.x * 7 + 5] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 5];
        block_boxes[threadIdx.x * 7 + 6] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 6];
    }
    __syncthreads();

    if (threadIdx.x < row_size) {
        const int cur_box_idx = THREADS_PER_BLOCK_NMS * row_start + threadIdx.x;
        const float *cur_box = boxes + cur_box_idx * 7;

        int i = 0;
        uint64_t t = 0;
        int start = 0;
        if (row_start == col_start) {
          start = threadIdx.x + 1;
        }
        for (i = start; i < col_size; i++) {
            if (iou_bev(cur_box, block_boxes + i * 7) > nms_overlap_thresh){
                t |= 1ULL << i;
            }
        }
        const int col_blocks = DIVUP(boxes_num, THREADS_PER_BLOCK_NMS);
        mask[cur_box_idx * col_blocks + col_start] = t;
    }
}


__device__ inline float iou_normal(float const * const a, float const * const b) {
    //params: a: [x, y, z, dx, dy, dz, heading]
    //params: b: [x, y, z, dx, dy, dz, heading]

    float left = fmaxf(a[0] - a[3] / 2, b[0] - b[3] / 2), right = fminf(a[0] + a[3] / 2, b[0] + b[3] / 2);
    float top = fmaxf(a[1] - a[4] / 2, b[1] - b[4] / 2), bottom = fminf(a[1] + a[4] / 2, b[1] + b[4] / 2);
    float width = fmaxf(right - left, 0.f), height = fmaxf(bottom - top, 0.f);
    float interS = width * height;
    float Sa = a[3] * a[4];
    float Sb = b[3] * b[4];
    return interS / fmaxf(Sa + Sb - interS, 1e-8);
}


__global__ void nms_normal_kernel(const int boxes_num, const float nms_overlap_thresh,
                           const float *boxes, uint64_t*mask){
    //params: boxes (N, 7) [x, y, z, dx, dy, dz, heading]
    //params: mask (N, N/THREADS_PER_BLOCK_NMS)

    const int row_start = blockIdx.y;
    const int col_start = blockIdx.x;

    // if (row_start > col_start) return;

    const int row_size = fminf(boxes_num - row_start * THREADS_PER_BLOCK_NMS, THREADS_PER_BLOCK_NMS);
    const int col_size = fminf(boxes_num - col_start * THREADS_PER_BLOCK_NMS, THREADS_PER_BLOCK_NMS);

    __shared__ float block_boxes[THREADS_PER_BLOCK_NMS * 7];

    if (threadIdx.x < col_size) {
        block_boxes[threadIdx.x * 7 + 0] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 0];
        block_boxes[threadIdx.x * 7 + 1] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 1];
        block_boxes[threadIdx.x * 7 + 2] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 2];
        block_boxes[threadIdx.x * 7 + 3] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 3];
        block_boxes[threadIdx.x * 7 + 4] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 4];
        block_boxes[threadIdx.x * 7 + 5] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 5];
        block_boxes[threadIdx.x * 7 + 6] = boxes[(THREADS_PER_BLOCK_NMS * col_start + threadIdx.x) * 7 + 6];
    }
    __syncthreads();

    if (threadIdx.x < row_size) {
        const int cur_box_idx = THREADS_PER_BLOCK_NMS * row_start + threadIdx.x;
        const float *cur_box = boxes + cur_box_idx * 7;

        int i = 0;
        uint64_t t = 0;
        int start = 0;
        if (row_start == col_start) {
          start = threadIdx.x + 1;
        }
        for (i = start; i < col_size; i++) {
            if (iou_normal(cur_box, block_boxes + i * 7) > nms_overlap_thresh){
                t |= 1ULL << i;
            }
        }
        const int col_blocks = DIVUP(boxes_num, THREADS_PER_BLOCK_NMS);
        mask[cur_box_idx * col_blocks + col_start] = t;
    }
}

void boxesoverlapLauncher(const int num_a, const float *boxes_a, const int num_b, const float *boxes_b, float *ans_overlap){

    dim3 blocks(DIVUP(num_b, THREADS_PER_BLOCK), DIVUP(num_a, THREADS_PER_BLOCK));  // blockIdx.x(col), blockIdx.y(row)
    dim3 threads(THREADS_PER_BLOCK, THREADS_PER_BLOCK);

    boxes_overlap_kernel<<<blocks, threads>>>(num_a, boxes_a, num_b, boxes_b, ans_overlap);
#ifdef DEBUG
    cudaDeviceSynchronize();  // for using printf in kernel function
#endif
}

void boxesioubevLauncher(const int num_a, const float *boxes_a, const int num_b, const float *boxes_b, float *ans_iou){

    dim3 blocks(DIVUP(num_b, THREADS_PER_BLOCK), DIVUP(num_a, THREADS_PER_BLOCK));  // blockIdx.x(col), blockIdx.y(row)
    dim3 threads(THREADS_PER_BLOCK, THREADS_PER_BLOCK);

    boxes_iou_bev_kernel<<<blocks, threads>>>(num_a, boxes_a, num_b, boxes_b, ans_iou);
#ifdef DEBUG
    cudaDeviceSynchronize();  // for using printf in kernel function
#endif
}


void nmsLauncher(const float *boxes, uint64_t* mask, int boxes_num, float nms_overlap_thresh){
    dim3 blocks(DIVUP(boxes_num, THREADS_PER_BLOCK_NMS),
                DIVUP(boxes_num, THREADS_PER_BLOCK_NMS));
    dim3 threads(THREADS_PER_BLOCK_NMS);
    nms_kernel<<<blocks, threads>>>(boxes_num, nms_overlap_thresh, boxes, mask);
}


void nmsNormalLauncher(const float *boxes, uint64_t* mask, int boxes_num, float nms_overlap_thresh){
    dim3 blocks(DIVUP(boxes_num, THREADS_PER_BLOCK_NMS),
                DIVUP(boxes_num, THREADS_PER_BLOCK_NMS));
    dim3 threads(THREADS_PER_BLOCK_NMS);
    nms_normal_kernel<<<blocks, threads>>>(boxes_num, nms_overlap_thresh, boxes, mask);
}

iou3d_nms.h修改之后：

#ifndef IOU3D_NMS_H
#define IOU3D_NMS_H

#include <torch/serialize/tensor.h>
#include <vector>
#include <assert.h>
#include <cuda.h>
#include <cuda_runtime_api.h>

//int boxes_aligned_overlap_bev_gpu(at::Tensor boxes_a, at::Tensor boxes_b, at::Tensor ans_overlap);
int boxes_overlap_bev_gpu(at::Tensor boxes_a, at::Tensor boxes_b, at::Tensor ans_overlap);
int boxes_iou_bev_gpu(at::Tensor boxes_a, at::Tensor boxes_b, at::Tensor ans_iou);
int nms_gpu(at::Tensor boxes, at::Tensor keep, float nms_overlap_thresh);
int nms_normal_gpu(at::Tensor boxes, at::Tensor keep, float nms_overlap_thresh);

#endif

iou3d_nms_api.cpp修改之后：

#include <torch/serialize/tensor.h>
#include <torch/extension.h>
#include <vector>
#include <cuda.h>
#include <cuda_runtime_api.h>

#include "iou3d_cpu.h"
#include "iou3d_nms.h"


PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    //m.def("boxes_aligned_overlap_bev_gpu", &boxes_aligned_overlap_bev_gpu, "aligned oriented boxes overlap");
    m.def("boxes_overlap_bev_gpu", &boxes_overlap_bev_gpu, "oriented boxes overlap");
    m.def("boxes_iou_bev_gpu", &boxes_iou_bev_gpu, "oriented boxes iou");
    m.def("nms_gpu", &nms_gpu, "oriented nms gpu");
    m.def("nms_normal_gpu", &nms_normal_gpu, "nms gpu");
    m.def("boxes_aligned_iou_bev_cpu", &boxes_aligned_iou_bev_cpu, "aligned oriented boxes iou");
    m.def("boxes_iou_bev_cpu", &boxes_iou_bev_cpu, "oriented boxes iou");
}

cd 到clone好的OpenPCDet目录下，安装pcdet库：python setup.py develop

二、下载数据集

打开D:\PythonProject\PointPillars\OpenPCDet\data\kitti，创建testing和training

下载好zip文件之后，需要将数据集放在OpenPCDet/data/kitti文件夹下面，新建training、testing文件夹，将calib、velodyne、label_2、image_2文件夹解压，并把他们放到对应的training、testing文件夹中，testing文件夹没有label_2文件。之后我们可以看到文件夹的组织结构如下：

在完成数据集的准备之后我们再对数据集进行处理,从而生成kitti_gt_database中的的 .bin 格式的文件和.pkl 格式文件，处理的命令如下：

python -m pcdet.datasets.kitti.kitti_dataset create_kitti_infos tools/cfgs/dataset_configs/ kitti_dataset.yaml
报错：缺少av2的库

 下载库：pip install av2==0.2.0

继续预处理数据：

报错：缺少kornia库

下载库：pip install kornia, pip install kornia[x]

 要先下载好cuda版本！按照下面的提示！因为torch是1.10或者1.9的，cuda就用11.3就行！！下载好对应的cudann。然后安装对应的pytorch！最后安装支持cuda的torch！！！

CUDA、CUDNN在windows下的安装及配置_widows 安装cuda cudnn_花花少年的博客-CSDN博客全网最详细 | Windows 安装 TensorFlow2.0 GPU 详细教程https://blog.csdn.net/m0_37605642/article/details/98854753#:~:text=%E4%BA%8C%E3%80%81%E5%AE%89%E8%A3%85%E8%BF%87%E7%A8%8B%201%20%EF%BC%881%EF%BC%89%E5%AE%89%E8%A3%85CUDA%20%E5%AE%89%E8%A3%85%E8%B7%AF%E5%BE%84%EF%BC%8C%E9%80%89%E6%8B%A9OK%EF%BC%88%E5%AE%89%E8%A3%85%E5%AE%8C%E6%88%90%E5%90%8E%E8%B7%AF%E5%BE%84%E4%BC%9A%E8%87%AA%E5%8A%A8%E5%8F%98%E5%8C%96%EF%BC%8C%E6%89%80%E4%BB%A5%E8%BF%99%E9%87%8C%E8%B7%AF%E5%BE%84%E9%80%89%E6%8B%A9%E9%BB%98%E8%AE%A4%E5%B0%B1%E5%A5%BD%EF%BC%89%20%E8%87%AA%E5%AE%9A%E4%B9%89%E5%AE%89%E8%A3%85%20%E5%8F%96%E6%B6%88%E5%8B%BE%E9%80%89GeForce%20Experience%20...,...%204%20%EF%BC%884%EF%BC%89%E6%B5%8B%E8%AF%95CUDA%20CUDA%E7%9A%84%E7%89%88%E6%9C%AC%E6%98%AF10.1%20%EF%BC%885%EF%BC%89%E6%9F%A5%E8%AF%A2%E6%98%BE%E5%8D%A1%E7%AE%97%E5%8A%9B%20%E8%BF%9B%E5%85%A5%E7%9B%AE%E5%BD%95%20deviceQuery.exe%20成功预处理文件：

  三、训练

cd到tools下单卡训练：python train.py --cfg_file cfgs/kitti_models/pointpillar.yaml

报错：

环境应该不匹配，重装一下虚拟环境，这里的cumm包是跟着torch一起下载的，如果有问题，那只能重新来一次了：

打开Anaconda Prompt安装虚拟环境：

conda create -n OpenPCDet python==3.8

打开虚拟环境：

conda activate OpenPCDet

安装torch：

pip install torch==1.9.1+cu111 torchvision==0.10.1+cu111 torchaudio==0.9.1 -f https://download.pytorch.org/whl/torch_stable.html

cd 到clone好的OpenPCDet目录下，安装其他依赖库：

pip install -r requirements.txt

会报错：

win10没有sharedarray的包，自己装一个，先把git上的代码clone下来，然后里面文件张这样：

GitHub - imaginary-friend94/Shared-Array-for-Windows: Share numpy arrays between processes

  然后修改setup.py和shared_memory_python.cpp：参考：在Windows上安装openPCDet遇到的一些问题总结_iou3d_nms_cuda_努力学pc的博客-CSDN博客

修改后，编译：python setup.py develop

 现在虚拟环境里的包有：

cd 到clone好的OpenPCDet目录下，安装pcdet库：python setup.py develop

报错：

ImportError: DLL load failed while importing iou3d_nms_cuda :  The parameter is incorrect. 

应该是cuda版本和装的版本不匹配，把cuda卸掉重新装一次。这次装了11.3的版本。

ImportError: cannot import name 'iou3d_nms_cuda' from 'pcdet.ops.iou3d_nms' (unknown location)

还是包冲突的问题，重新把虚拟环境清空，然后再次走一遍上述流程：

在装pcdet的时候报错：CUDA\v11.3\include\crt/host_config.h(160): fatal error C1189: #error: 

根据下述网址，这个好像是因为底层c编译版本的某些东西不匹配，改一下c的版本区间：

#error: -- C1189: unsupported Microsoft Visual Studio version!icon-default.png?t=N3I4https://www.cnblogs.com/MK-XIAOYU/p/13413039.html

改完之后，报新的错误：

error STL1002: Unexpected compiler version, expected CUDA 11.6 or newer.

这个绕不过去了。。。。把cuda卸载再装到116版本，但是之前111和113都是可以的，这里为什么vc2019和cuda会冲突，我也没有找到匹配的答案。。。

装到116版本后：

openpcdet error: command 'E:\\NVIDIA GPU Computing Toolkit\\CUDA\\v11.6\\bin\\nvcc.exe' failed with exit code 2

不知道为啥，这里找不到cuda了，我自己的在命令行是可以拿到cuda的版本号的。

总体原因，还是虚拟环境的包冲突。把虚拟环境再重装一次。

sinounuo
关注
  • 0
    点赞
  • 踩
  • 15
    收藏
    觉得还不错? 一键收藏
  • 7
    评论
OPenPCDet windows流程及其问题
m0_46611008的博客
02-13 1769
首先的首先极其不推荐将OPenPCDet运行在Windows上，过程非常复杂，适配也不是很好，不推荐在windows下载并训练，本人做这个主要是方便再笔记本电脑上对实验结果进行整理，处理一些简单的推理评估等任务。如有必要请继续阅读：
OpenPCDet权重文件pointpillar_7728.pth
01-17
OpenPCDet权重文件pointpillar_7728.pth
Ubuntu18.04复现OpenPCDet过程(cuda11.3、pytorch1.10)
YYDS_COM的博客
03-18 1068
数据预处理的过程如果报错，检查training和testing文件中的image_2中的文件数量是否与training和testing文件下其他目录中的文件数量相同，如果不同，从另一个image_2中复制一部分补齐道另一个image_2中。注意：数据集的布置建议采用小编后边给出的形式，作者的布置方式新手可能不会操作，建立软连接哪一步可以直接忽略，直接按照小编后续给出的布置方式直接布置数据集。（2）末行模式下，按esc慢退、按两次esc快退、或者删除所有命令，可以回到命令模式。命令模式、编辑模式、末行模式。
基于OpenPCDet框架进行Pointpillars算法环境搭建并基于TensorRT和ROS部署
超爱吃小蛋糕的66的博客
04-29 1709
基于OpenDet进行训练Pointpillars算法，基于tensorrt-8.5进行部署并移植到ROS
windows下配置OpenPCDet
qq_61981718的博客
05-18 1195
2安装pytorchpip3 install torch==1.8.1+cu111 torchvision==0.9.1+cu111 torchaudio==0.8.1 -f https://download.pytorch.org/whl/torch_stable.html -i https://pypi.tuna.tsinghua.edu.cn/simple。4.克隆 OpenPCDet项目 git clone https://github.com/open-mmlab/OpenPCDet.git。
win10 OpenPCDet 训练KITTI以及自己的数据集
m0_68312479的博客
08-06 9618
win10 OpenPCDet 训练KITTI以及自己的数据集
复现pointpillars，anaconda
sinounuo的博客
05-09 427
解决方法一：将conda-forge添加到搜索路径上首先，当出现这种报错时，应该首先尝试使用以下命令将conda-forge channel添加到你的channel列表中：conda config --append channels conda-forge它告诉conda在搜索软件包时也要在conda-forge channel上查看。在此期间我遇到下不下来最后一个包google-sparsehash，参考以下内容，在conda的官网上找到自己包的下载指令，如果依旧下载不下来，就直接本地下载。
学习笔记|Point Pillars
onesideline的博客
06-05 511
Point Pillars接受点云作为输入，估计目标的3d框，其主要由三个阶段组成。1）将点云转换为稀疏伪图像的特征编码器网络2）一个2d卷积主干，将伪图像处理成高级表示3）检测头，用于检测和回归3d框。
second.pytorch:用于KITTI对象检测的PointPillars
05-11
欢迎来到PointPillars。 此存储库演示了如何通过对现有开放源代码库进行最小的要求更改，来在上重现（将在CVPR 2019上发布）的。 这不是官方的nuTonomy代码库，但可以用于匹配已发布的PointPillars结果。 警告：...
PointPillars
05-07
TensorFlow中的PointPillars Tensorflow中的Point PIllars 3D检测网络实施。 欢迎外部贡献，请分叉此存储库，并查看有关代码中可能需要改进的问题。安装从Kitti下载LiDAR，Calibration和Label_2 zip文件，并解压缩...
永恒之蓝复现kaliwin7
最新发布
07-19
永恒之蓝复现kaliwin7
论文复现-深度补全算法
03-10
1、传统深度不全算法复现 2、详细内容可见：https://blog.csdn.net/qq_43627520/article/details/123344654?spm=1001.2014.3001.5502 3、复现代码可直接运行、包含有测试用例、以及验证代码
GT1050 PointPillar成功运行
u010647296的博客
10-30 2116
目的 之前在感知这里主要是考传统的算法，但是遇到了瓶颈，计算效率提不上来，所以尝试一下其他的方法。目前看到pointpillar这种方法效果比较好，就打算尝试一下,跑通测试一下效果。 piontpillar 需要的环境 1做了一些尝试：在xavier上 需要nvidai驱动，cuda10.X+cudnn7.6.以及sensort。 实验条件：我手上有一台Xavier 和 一台联想拯救者Y7000（GTX1050），均安装了ubuntu18.04+ros meldic，其中Xavier已经安装了nvidai驱
【3D 目标检测模型部署】全链条打通！PointPillars 从模型到部署
qq_39967751的博客
07-14 7766
近几年点云目标检测一直是计算机视觉领域中很火的任务，各种方法也是层出不穷，从基于点云的PointNet、PointNet++，到使用体素方法的VoxelNet。这些网络取得了不错的检测性能，但同时也存在一些问题。PointNet系列的Point-based模型直接对点云进行处理，可以减少位置信息的损失，但同时也带来了巨大的计算资源消耗，使其很难做到实时。我们来看一下PointPillars中的PillarFeatureNet模块的具体操作，如下图所示。outshttps。......
OpenPCDet环境配置&&HVPR代码复现
weixin_44606139的博客
10-24 1108
OpenPCDet环境配置&&HVPR代码复现
OpenPCDet——环境配置和训练测试（升级系统、Nvidia驱动、cuda11.3、cudnn8.2）
xiaohu的博客
04-30 6642
文章目录NVIDIA 驱动升级CUDA安装 笔记本显卡驱动NVIDIA 418，cuda9.0 比较旧，和工作的环境不兼容，导致tensorRT装版本也不一致，决定把GTX 1060显卡升级到最新的驱动 NVIDIA 驱动升级 sudo apt-get remove --purge nvidia* sudo sh NVIDIA-Linux-x86_64-418.56.run --uninstall 禁用nouveau sudo gedit /etc/modprobe.d/blacklist.conf 在
点云目标检测——pointpillars环境配置与训练部署
dui845819593的博客
10-13 1450
如何在mmdet3d中使用自有数据集？- 知乎。
3d目标检测OpenPCDet安装
qq_43552048的博客
04-09 578
torch-1.9.0+cu111-cp38-cp38-linux_x86_64.whl和torchvision-0.10.0+cu111-cp38-cp38-linux_x86_64.whl右键复制下载链接，使用。最近在GPU服务器主RTX4090上部署了一下，折腾了一天发现就算4090安装cuda11.1按照上面RTX3060的安装步骤也无法安装成功，需要变通的地方在于torch。啥的说明安装成功了，可能还会出现缺少依赖包或Module的问题，用pip安装就行。，spconv需要安装2.x版本的直接。
安装OpenPCDet遇到的问题
K的学习之路
05-13 1015
环境: Item Content OS Ubuntu 18.04 Nvidia-driver 450.119.03 CUDA 10.1.243 大部分情况的错误为nvcc编译时，错误的flag项。 安装SpConv git clone https://github.com/traveller59/spconv.git --recursive sudo apt-get install libboost-all-dev python setup.py bdist_wheel cd .
openpcdet复现pointpillars
08-03
要复现OpenPCDet中的PointPillars算法，可以按照以下步骤进行操作： 1. 首先，克隆OpenPCDet的GitHub仓库到本地。可以使用以下命令进行克隆： ``` git clone https://github.com/open-mmlab/OpenPCDet.git ``` ...
评论 7
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值