背景
近期我在学习点云目标识别技术,发现OpenPCDet模型支持多模型组合进行目标检测,功能相当实用。不过查阅相关资料时,我发现大多数教程都基于Ubuntu系统,Windows系统的部署指南非常有限,仅找到一篇相关文章OpenPCDet+Anaconda+Windows安装过程,这个教程提出了要修改相关文件已适配windows系统的编译,我是在这篇文章的指导下进行部署的。但是我在使用过程中发现修改的不是很彻底(或者不适合我的设备环境),于是在Deepseek的支持下,最终成功完成了在Windows环境下的部署工作。
写在前面:
1、我不保证此教程可以使每个人都不报错的安装成功,这显然是不可能的,但是大概步骤和修改的东西是差不多的,如果有问题,就问Deepseek!包好使!
2、本教程的主要贡献是如何修改源文件,对于前面的环境部署,我这写的不是很详细,CSDN中有很多相关教程,大家搜索即可。
一、环境搭建
主要安装:Python,Cuda,Visual studio 和Pytorch。
1、Conda创建python环境
这里我建议给点云任务单独创建一个环境,避免新装的包与之前的项目存在版本冲突,我安装的是python3.8版本。这里有两种方法,一个是conda创建新的环境,这里我就不详细展示了,这个网上有很多教程。另一个是克隆本机之前的环境并命名,因为我之前做视觉,已经安装过prtorch,cuda,opencv等很多包了,我不想再重新下载,于是就克隆过来了,后面有冲突的话,再相应调整即可。
克隆步骤如下:new_env_name是你新环境的名字,existing_env_name是现存的python3.8版本的环境名,在base端中执行下面命令即可。
conda create --name new_env_name --clone existing_env_name
2、Cuda安装和环境配置
我的显卡是Nvidia RTX4050,下载的是Cuda11.8;安装的是Visual studio2019;安装的cuDNNV8.6.0。这里我在下方引用一个前辈的帖子,大家可以按照他这个步骤进行操作,安装自己电脑合适的版本即可。
需要注意的是:
(1)在安装Visual studio时大家一定要勾选使用C++的桌面开发!!
(2)cuDNN的下载地址是:cuDNN Archive | NVIDIA Developer。在这里面找到符合你Cuda版本的cuDNN,进行下载,也可以直接在SCDN上搜相应的百度云资源。cuDNNV8.6.0的网盘连接如下:cuDNNV8.6.0。下载完后,记得把bin,include和lib中的文件分别复制到CUDA安装目录下的bin,include和lib中,Cuda安装目录一般在C盘,如下图所示。
3、Pytorch安装
在终端中输入nvcc-version,查看安装的cuda的编译版本版本
nvcc-version
我的是cuda11.8与之前安装的cuda驱动一致,如下图所示。
那么下面就要去PyTorch官网上下载对应的pytorch即可。点击页面最下方的previous versions of PyTorch按钮,查找自己的pytorch版本。我安装的命令如下。注意:一定要在该项目的python环境下安装!这个包很大,建议保持较好的网络环境。
conda install pytorch==2.0.1 torchvision==0.15.2 torchaudio==2.0.2 pytorch-cuda=11.8 -c pytorch -c nvidia
安装完成后,在pycharm的终端运行pip show torch结果如下。
这样大环境就配置完成了,环境配置过程中如果有不懂得可以搜相关教程,有很多很详细的。接下来,下载openPCDet项目并部署。
二、下载和部署openPCDet项目
1、下载openPCDet项目
在github中下载openPCDet项目,下载zip文件,解压后在pycharm中创建项目,编译环境选择之前创建好的点云处理新环境。也可以用git clone 将项目克隆到本地。
2、安装必要库
用pycharm打开requirements,需要指出的时,并不需要完全安装里面所写的库,比如不需要安装SharedArray库,后面把这个库给注释掉即可。下面是我python3.8环境下修改requirements后的库和安装的对应版本。库的版本不用和我一一对应,只要自己安装的库版本不冲突即可。
注释掉相关文件中需要的SharedArray库命令:
(1)setup.py文件
(2)其他文件
pcdet\datasets\augmentor\database_sampler.py(注释import SharedArray)
pcdet\datasets\waymo\waymo_dataset.py(注释import SharedArray)
pcdet\models\model_utils\mppnet_utils.py(注释from os import getgrouplist windows)
三、修改openPCDet源文件
这个环节是windows系统部署openPCDet项目的关键步骤,主要是涉及CUDA中Linux系统和windows系统对于数据类型支持问题。每个文件有许多要修改的东西,我就不一一列举了。我直接把我修改好的相关文件放上来,有需要的直接代替就行。
需要指出来:因为我是修改好了,后面记录的部署步骤,可能会遗忘一些东西,如果报错也不用怕,记住有Deepseek会帮忙,有问题直接问它即可!比问心一言好使多了!
1、修改ingroup_inds_kernel.cu文件和ingroup_inds.cpp文件
文件位置:pcdet\ops\ingroup_inds\src\
修改内容:两个文件都是将所有long改为为int32_t,所有unsigned long long 改为uint64_t
修改后的文件如下:需要的直接复制代替
ingroup_inds_kernel.cu文件
#include <assert.h>
#include <vector>
#include <math.h>
#include <stdio.h>
#include <torch/serialize/tensor.h>
#include <torch/types.h>
#include "cuda_fp16.h"
#define CHECK_CALL(call) \
do \
{ \
const cudaError_t error_code = call; \
if (error_code != cudaSuccess) \
{ \
printf("CUDA Error:\n"); \
printf(" File: %s\n", __FILE__); \
printf(" Line: %d\n", __LINE__); \
printf(" Error code: %d\n", error_code); \
printf(" Error text: %s\n", \
cudaGetErrorString(error_code)); \
exit(1); \
} \
} while (0)
#define THREADS_PER_BLOCK 256
#define DIVUP(m, n) ((m) / (n) + ((m) % (n) > 0))
// #define DEBUG
// #define ASSERTION
__global__ void ingroup_inds_kernel(
const int32_t *group_inds,
int32_t *out_inds,
int *ingroup_counter,
int N
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= N) return;
int32_t this_group_id = group_inds[idx];
int cnt = atomicAdd(&ingroup_counter[this_group_id], 1);
out_inds[idx] = cnt;
}
void ingroup_inds_launcher(
const int32_t *group_inds,
int32_t *out_inds,
int N,
int max_group_id
) {
int *ingroup_counter = NULL;
CHECK_CALL(cudaMalloc(&ingroup_counter, (max_group_id + 1) * sizeof(int)));
CHECK_CALL(cudaMemset(ingroup_counter, 0, (max_group_id + 1) * sizeof(int)));
dim3 blocks(DIVUP(N, THREADS_PER_BLOCK));
dim3 threads(THREADS_PER_BLOCK);
ingroup_inds_kernel<<<blocks, threads>>>(
group_inds,
out_inds,
ingroup_counter,
N
);
cudaFree(ingroup_counter);
#ifdef DEBUG
CHECK_CALL(cudaGetLastError());
CHECK_CALL(cudaDeviceSynchronize());
#endif
return;
}
ingroup_inds.cpp文件
#include <assert.h>
#include <torch/extension.h>
#include <torch/serialize/tensor.h>
#include <vector>
#define CHECK_CUDA(x) \
TORCH_CHECK(x.device().is_cuda(), #x, " must be a CUDAtensor ")
#define CHECK_CONTIGUOUS(x) \
TORCH_CHECK(x.is_contiguous(), #x, " must be contiguous ")
#define CHECK_INPUT(x) \
CHECK_CUDA(x); \
CHECK_CONTIGUOUS(x)
void ingroup_inds_launcher(
const int32_t *group_inds_data,
int32_t *out_inds_data,
int N,
int max_group_id
);
void ingroup_inds_gpu(
at::Tensor group_inds,
at::Tensor out_inds
);
void ingroup_inds_gpu(
at::Tensor group_inds,
at::Tensor out_inds
) {
CHECK_INPUT(group_inds);
CHECK_INPUT(out_inds);
int N = group_inds.size(0);
int max_group_id = group_inds.max().item().toLong();
int32_t *group_inds_data = group_inds.data_ptr<int32_t>();
int32_t *out_inds_data = out_inds.data_ptr<int32_t>();
ingroup_inds_launcher(
group_inds_data,
out_inds_data,
N,
max_group_id
);
}
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
m.def("forward", &ingroup_inds_gpu, "cuda version of get_inner_win_inds of SST");
}
2、修改iou3d_nms.cpp文件
位置:pcdet/ops/iou3d_nms/src/iou3d_nms.cpp
修改的内容很多,我建议直接用我修改后的!如果编译有报错,再询问deepseek,做相应修改。
(1)全局替换,所有long改为为int64_t,所有unsigned long long 改为uint64_t,修改remv_cpu函数。
(2)因为在 C/C++ 中,静态数组(Stack Array)的大小必须在编译时确定,而 col_blocks 是一个运行时计算的变量(通过 DIVUP 宏定义),因此无法直接用于静态数组的声明。需要使用 std::vector 替代原生数组,避免手动内存管理。如下:
(3)全局替换,将 tensor.type()
替换为 tensor.device()。
将 tensor.data<T>()
替换为 tensor.data_ptr<T>(),其中T是变量名称。
(4)修复 C2039 错误:"is_cuda": 不是 "c10::TensorOptions" 的成员
(5)修复C2664 错误:无法将 std::vector
转换为 void *
修改后的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 "iou3d_nms.h"
#define CHECK_CUDA(x) do { \
if (!x.device().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 boxesalignedoverlapLauncher(const int num_box, const float *boxes_a, const float *boxes_b, float *ans_overlap);
void boxesoverlapLauncher(const int num_a, const float *boxes_a, const int num_b, const float *boxes_b, float *ans_overlap);
void PairedBoxesOverlapLauncher(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_aligned_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: (N, 7) [x, y, z, dx, dy, dz, heading]
// params ans_overlap: (N, 1)
CHECK_INPUT(boxes_a);
CHECK_INPUT(boxes_b);
CHECK_INPUT(ans_overlap);
int num_box = boxes_a.size(0);
int num_b = boxes_b.size(0);
assert(num_box == num_b);
const float * boxes_a_data = boxes_a.data_ptr<float>();
const float * boxes_b_data = boxes_b.data_ptr<float>();
float * ans_overlap_data = ans_overlap.data_ptr<float>();
boxesalignedoverlapLauncher(num_box, boxes_a_data, boxes_b_data, ans_overlap_data);
return 1;
}
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_ptr<float>();
const float * boxes_b_data = boxes_b.data_ptr<float>();
float * ans_overlap_data = ans_overlap.data_ptr<float>();
boxesoverlapLauncher(num_a, boxes_a_data, num_b, boxes_b_data, ans_overlap_data);
return 1;
}
int paired_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: (N, 7) [x, y, z, dx, dy, dz, heading]
// params ans_overlap: (N, 1)
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);
assert(num_a == num_b);
const float * boxes_a_data = boxes_a.data_ptr<float>();
const float * boxes_b_data = boxes_b.data_ptr<float>();
float * ans_overlap_data = ans_overlap.data_ptr<float>();
PairedBoxesOverlapLauncher(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_ptr<float>();
const float * boxes_b_data = boxes_b.data_ptr<float>();
float * ans_iou_data = ans_iou.data_ptr<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_ptr<float>();
int64_t * keep_data = keep.data_ptr<int64_t>();
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);
// unsigned long long mask_cpu[boxes_num * col_blocks];
// unsigned long long *mask_cpu = new unsigned long long [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); // 163修改后动态分配内存
// 使用 remv_cpu.data() 获取指针
uint64_t* remv_cpu_ptr = remv_cpu.data();
for (int i = 0; i < col_blocks; i++) {
remv_cpu_ptr[i] = 0;
}
memset(remv_cpu.data(), 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.data()[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_ptr<float>();
int64_t * keep_data = keep.data_ptr<int64_t>();
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);
// unsigned long long mask_cpu[boxes_num * col_blocks];
// unsigned long long *mask_cpu = new unsigned long long [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); // 214修改后 动态分配内存
memset(remv_cpu.data(), 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.data()[j] |= p[j];
}
}
}
if ( cudaSuccess != cudaGetLastError() ) printf( "Error!\n" );
return num_to_keep;
}
3、修改iou3d_nms_kernel.cu文件
文件位置:pcdet/ops/iou3d_nms/src/iou3d_nms_kernel.cu
修改内容:
(1)修复CUDA设备代码中找不到标识符“EPS”的问题。
在 CUDA 内核文件(iou3d_nms_kernel.cu
)中显式定义 EPS
,确保其在设备端可见。在文件开头添加以下定义:
#define EPS 1e-8f // ✅ 使用宏定义确保设备端可见
如图所示
并将const float EPS注释掉,如图所示
(2)修改check_rect_cross函数名为check_rect_cross_cuda,避免与iou3d_nms中重复
修改后的文件:
四、修改ska_hash
库源码
PyTorch 2.0.1 内部依赖的 ska_hash 库版本与 MSVC 编译器存在兼容性问题。 PyTorch 2.0.1 官方预编译包未完全适配 Windows 下的所有 C++ 模板场景。 OpenPCDet 依赖的 spconv 库在编译时会报问题:warning C4624: “ska::detailv3::sherwood_v3_entry<std::pair<K,V>>”: 已将析构函数隐式定义为“已删除”。
解决方案:修改 ska_hash
库源码,显式定义析构函数以适配 MSVC。可能pytorch2.0.1删减了部分第三方库的完整实现,导致我无法在PyTorch安装目录中找到 ska_hash
头文件。因此我直接从源码仓库获取。
1、获取ska_hash
库源码
在openPCDet项目的终端,输入下面程序。会自动将仓库克隆到当前工作目录下的同名文件夹 flat_hash_map
中。
# 克隆 ska_hash 仓库到本地
git clone https://github.com/skarupke/flat_hash_map.git
然后修改flat_hash_map.hpp中的sherwood_v3_entry定义。在文件中ctrl+F查找sherwood_v3_entry,结果如下
将这个结构体用下面的代码替换,注意:替换的是整个结构体的定义(160行到198行的内容)
struct sherwood_v3_entry
{
sherwood_v3_entry() = default; // 显式声明默认构造函数
sherwood_v3_entry(int8_t distance_from_desired)
: distance_from_desired(distance_from_desired)
{
}
// 显式定义默认析构函数,确保正确处理联合体
~sherwood_v3_entry() = default;
static sherwood_v3_entry* empty_default_table()
{
static sherwood_v3_entry result[min_lookups] = { {}, {}, {}, {special_end_value} };
return result;
}
bool has_value() const { return distance_from_desired >= 0; }
bool is_empty() const { return distance_from_desired < 0; }
bool is_at_desired_position() const { return distance_from_desired <= 0; }
template<typename... Args>
void emplace(int8_t distance, Args &&... args)
{
new (std::addressof(value)) T(std::forward<Args>(args)...);
distance_from_desired = distance;
}
void destroy_value()
{
value.~T(); // 显式调用 T 的析构函数
distance_from_desired = -1;
}
int8_t distance_from_desired = -1;
static constexpr int8_t special_end_value = 0;
union { T value; }; // 联合体成员需要显式管理生命周期
};
2、创建两个文件夹
一是将修改后的 flat_hash_map.hpp
复制到 PyTorch 的第三方库目录,手动创建ska_hash文件夹
我的位置是:
C:\Users\zero9\.conda\envs\PCD_detect\Lib\site-packages\torch\include\third_party\ska_hash\
另外是在C盘创建如下的文件夹C:\patched_headers\ska_hash,并将修改的flat_hash_map.hpp文件放进去
3、编译
在pycharm的项目终端中输入下面命令,清理构建缓存后编译。
python setup.py clean --all
set INCLUDE=C:\patched_headers;%INCLUDE%
python setup.py build_ext develop
等待一段时间,没有如果出现报错(部分警告可以暂时忽略),有报错可以问deepseek!嘿嘿~。在终端最终显示下面内容,则表明安装成功!
Using c:\users\zero9\.conda\envs\pcd_detect\lib\site-packages
Finished processing dependencies for pcdet==0.6.0+0
五、部署成功验证
1、在点云环境中导入pcdet库
可以显示pcdet库的版本则表明成功部署openPCDet。
import pcdet
print(pcdet.__version__) # 应输出 0.6.0+0
2、运行 OpenPCDet 的测试脚本(如 demo.py
)
(1)权重和测试文件下载
从 OpenPCDet 的官方模型库下载预训练模型(如 pv_rcnn.pth
),并放置在 tools\weight
目录。在这我给出我下载好的模型权重链接:pv_rcnn_8369.pth。由于kitti数据集非常大,下载时间很久。这里先给出一个数据集文件000000.bin测试文件,用来跑demo。
(2)测试
在pycharm终端的openPCDet项目中,cd到tools文件夹,运行下面命令,出现下面结果说明openPCDet可以使用。
python demo.py --cfg_file cfgs/kitti_models/pv_rcnn.yaml --ckpt weight/pv_rcnn_8369.pth --data_path 000000.bin
OK,结束。下面就可以跑自己的数据和模型了。
如果觉得有帮助,可以动动小手指给我点点赞~