八. 实战：CUDA-BEVFusion部署分析-BEVFusion Initialization

前言

自动驾驶之心推出的 《CUDA与TensorRT部署实战课程》，链接。记录下个人学习笔记，仅供自己参考

本次课程我们来学习下课程第八章—实战：CUDA-BEVFusion部署分析，一起来学习 CUDA-BEVFusion 的初始化代码的实现

课程大纲可以看下面的思维导图

0. 简述

从这节开始我们进入到本次课程的最后几个部分，也就是 CUDA-BEVFusion 的代码解析，在 八. 实战：CUDA-BEVFusion部署分析-学习CUDA-BEVFusion推理框架设计模式 中我们有提到其中的接口模式以及类与类之间、命名空间与命名空间之间的调用关系，还有 CUDA-BEVFusion 中各个类负责的内容等等，那对 CUDA-BEVFusion 整个推理框架的设计熟悉之后，还是建议大家自己捋一遍代码之后再来看课程讲解会比较清晰

CUDA-BEVFusion 整体代码还是比较复杂的，这边会分几个部分来讲解，主要是初始化、相机内外参 update 以及 forward 过程等等，foward 中包括 camera 的 forward、lidar 的 forward、bevpool 的 forward 等等，这里的每个部分都会涉及到相关的 CUDA 核函数加速，会跟大家去分析 CUDA 核函数的设计理念是什么以及加速的技巧是什么

大家可能第一次接触多模态模型推理，对其中的一些坐标转换可能不是很熟悉，这边也会简单过一下，另外点云部分会把 voxelization 体素化这部分的 CUDA 加速实现和大家讲解

这节我们主要来讲解 CUDA-BEVFusion 的初始化，看看它具体都做了哪些工作

下面我们开始本次课程的学习🤗

1. 案例运行

在 八. 实战：CUDA-BEVFusion部署分析-环境搭建 文章中博主有详细介绍 CUDA-BEVFusion 的环境配置，这里就不再赘述了，不过九个月过去了 CUDA-BEVFusion 有了一些更新（截止到 2024/9/22），主要有以下几个点需要大家注意：

• libprotobuf-dev==3.6.1 条件去除，目前要求 protobuf >= 3.12.0 即可，在 CUDA-BEVFusion/src/onnx/onnx-ml.pb.h 头文件有提到
• layernorm plugin 支持，在 head model 替换为 head.bbox.layernormplugin.plan 后 bbox 回归分类准确率大大提升，对比图如下
• TensorRT-10.x 支持
• spconv-1.1.10 版本发布
• CUDA-11.x 编译选项需设置为 -std=c++14，CUDA-12.x 编译选项需设置为 -std=c++17

Note：这里韩君老师分析的代码是 2023 年 12 月中旬的代码，调试工具使用的是 vscode

head.bbox.plan 的推理结果如下：

head.bbox.layernormplugin.plan 的推理结果如下：

2. Initialization代码分析

2.1 create_core函数

我们先从 main.cpp 看起：

整个 main 函数其实非常简单，它对外只暴露了一些必须的接口，这也是它设计巧妙的地方，也是我们需要去学习的，那我们这节主要来分析 create_core 函数，来看看 CUDA-BEVFusion 的初始化都做了哪些事情

create_core 函数主要是对整个 BEVFusion 的各个部分进行初始化，包括 camera、lidar、bevpool 等等，通过调用接口类，创建实现类的实体，大体上包括两个部分：

• 对于非 DNN 部分，分配预处理所需要的 device 上的 global memory 和 host 上的 pinned memory
• 对于 DNN 部分，除了分配内存以外，还需要反序列化 plan，创建推理上下文，准备进行推理

我们进入到 create_core 函数中：

首先需要设置的是对于 6 个 camera 图像的初始化的参数，通过 NormalizationParameter 结构体对象来初始化，我们可以理解它其实就是对 camera backbone model 预处理的一些参数进行相关的设置

NormalizationParameter 结构体的定义如下：

enum class NormType : int { Nothing = 0, MeanStd = 1, AlphaBeta = 2 };
enum class ChannelType : int { Nothing = 0, Invert = 1 };
enum class Interpolation : int { Nearest = 0, Bilinear = 1 };

struct NormMethod {
  float mean[3];
  float std[3];
  float alpha, beta;
  NormType type = NormType::Nothing;
  ChannelType channel_type = ChannelType::Nothing;

  // out = (x * alpha - mean) / std + beta
  static NormMethod mean_std(const float mean[3], const float std[3], float alpha = 1 / 255.0f, float beta = 0.0f,
                             ChannelType channel_type = ChannelType::Nothing);

  // out = x * alpha + beta
  static NormMethod alpha_beta(float alpha, float beta = 0, ChannelType channel_type = ChannelType::Nothing);

  // None
  static NormMethod None();
};

struct NormalizationParameter {
  int image_width;
  int image_height;
  int num_camera;
  int output_width;
  int output_height;
  float resize_lim;
  Interpolation interpolation;
  NormMethod method;
};

NormalizationParameter 结构体成员变量：

• image_width，image_height：输入图像原始宽高，nuScenes 数据集中为 1600x900
• num_camera：相机数量，nuScenes 数据集中为 6 个提供 360° 视角
• output_width，output_height：图像归一化处理后的输出分辨率，BEVFusion 中为 704x256
• resize_lim：图像缩放的最小比例限制
• interpolation：图像缩放时的插值方式，枚举类，包括 Nearest 最近邻插值和 Bilinear 双线性插值
• method：归一化方法，使用 NormMethod 结构体，可以是 MeanStd、AlphaBeta 或者 None

从中我们可以知道预处理做了减均值除标准差、除以255、resize 等操作，那和 2D Detection Model 的预处理方式很像，都是一些常见的处理步骤，那本身 camera backbone 也是沿用的一些经典的网络模型例如 resnet、swin transformer

我们接着往下走：

这里主要设置对于 LiDAR 点云的 voxel 进行初始化的参数，num_feature 表示每个点云包含的信息， x, y, z, indensity, ring index，那其实我们实际使用过程中最后一个维度的信息即 ring index 其实是不需要的，我们只要用到前面四个维度的信息就行，大家如果对点云存储的数据比较熟悉的话可能会知道，之前杜老师也提到过，大家感兴趣的可以看看：AutoCV第八课：3D基础

VoxelizationParameter 结构体的定义如下：

nvtype::Int3 VoxelizationParameter::compute_grid_size(const nvtype::Float3 &max_range, const nvtype::Float3 &min_range,
                                                      const nvtype::Float3 &voxel_size) {
  nvtype::Int3 size;
  size.x = static_cast<int>(std::round((max_range.x - min_range.x) / voxel_size.x));
  size.y = static_cast<int>(std::round((max_range.y - min_range.y) / voxel_size.y));
  size.z = static_cast<int>(std::round((max_range.z - min_range.z) / voxel_size.z));
  return size;
}

struct VoxelizationParameter {
  nvtype::Float3 min_range;
  nvtype::Float3 max_range;
  nvtype::Float3 voxel_size;
  nvtype::Int3 grid_size;
  int num_feature;
  int max_voxels;
  int max_points_per_voxel;
  int max_points;

  static nvtype::Int3 compute_grid_size(const nvtype::Float3& max_range, const nvtype::Float3& min_range,
                                        const nvtype::Float3& voxel_size);
};

VoxelizationParameter 结构体成员变量：

• min_range，max_range：三维空间的最小和最大边界，它们共同定义了点云数据在体素化时被考虑的三维空间范围，任何超出这个范围的点云数据将被忽略。
• voxel_size：每个体素的尺寸，定义了体素在 x、y、z 三个方向上的大小。通常，较小的体素能提供更高的分辨率，但也会增加计算开销
• grid_size：通过 compute_grid_size 函数计算得到的三维向量，表示 x、y、z 三个方向上的体素网格大小
• num_feature：表示每个点云点的特征向量，通常 LiDAR 点云中每个点至少包含 3 个坐标 x、y、z，但有时也可能包含其他信息，如 intensity 强度等
• max_voxels：定义体素化处理时允许的最大体素数量
• max_points_per_voxel：每个体素中允许包含的最大点数
• max_points：体素化处理时允许处理的最大点云点数

VoxelizationParameter 结构体方法：

• compute_grid_size：静态方法，用来根据 min_range、max_range 和 voxel_size 计算体素网格的大小

配置的 voxelization 各个参数中 x 轴范围在 [-54m, 54m]，y 轴范围在 [-54m, 54m]，z 轴范围在 [-5m, 3m]，这样就定义了点云感兴趣区域的边界范围是 108x108x8 的一个三维空间。每个体素的大小是 [0.075m, 0.075m, 0.2m]，也就是一个 0.075x0.075x0.2 的小立方体

grid_size 是计算体素网格的大小，前面我们知道 x 轴长度为 108m，体素大小为 0.075m，则 x 轴上会有 108 / 0.075 = 1440 个体素，同理 y 轴上也会有 1440 个体素，z 轴上会有 8 / 0.2 = 40 个体素，这也就解释了输入的点云经过体素化处理后的 shape 是 1440x1440x40

每个体素中允许的最大点云点数为 10，可以处理的最大点云数为 300000，允许的最大体素数量是 160000，每个点云有 5 个特征，包括 x、y、z 以及反射强度和 ring index

大家可能对体素化可能不太了解，这边老师拿了一张示例图来讲解：

上面有一个小兔子它有无数个点组成，那我们可以做一个 Voxel Grid，把一个空间分割成好几个 voxel（即一个个立方体小块），每一个 voxel 中包含的点的数量各不相同，有的 voxel 可能包含多个点，有个 voxel 可能一个点都没有，voxelization 体素化会对每个 voxel 中的所有点求和取平均找出最具有代表性的一个点，这个其实就是 voxelization 所做的事情

voxelization 在 PCL 库中其实是有实现的，那这里 CUDA-BEVFusion 其实是做了一个 CUDA 加速，那这个我们后面也会跟大家去讲

我们接着看下一个部分：

这里设置针对点云的 3D Sparse Convolution Network 进行初始化的参数，也就是 LiDAR 支路的 backbone 网络

SCNParameter 结构体的定义如下：

enum class CoordinateOrder : int {
  NoneOrder = 0,
  XYZ = 1,  // BEVFusion
  ZYX = 2   // CenterPoint
};

// use model accuracy during SCN model inference.
enum class Precision : int { NonePrecision = 0, Float16 = 1, Int8 = 2 };

struct SCNParameter {
  VoxelizationParameter voxelization;
  std::string model;
  CoordinateOrder order = CoordinateOrder::XYZ;
  Precision precision = Precision::Float16;
};

SCNParameter 结构体成员变量：

• voxelization：VoxelizationParameter 类型的成员变量，存储体素化的相关参数，用于配置 LiDAR 点云的体素化过程
• model：string 类型，表示 ONNX 模型的路径
• order：CoordinateOrder 枚举类型，指定 LiDAR 点云坐标的排列顺序
• precision：Precision 枚举类型，定义推理时的计算精度

接着往下看

这里主要是 geometry 的设置，这个比较重要，它设置的是 camera->BEV 过程中所需要的参数，也就是 BEVPool 的 precomputation 所需要的参数。这个我们前面在讲 BEVPool 的优化方案时有提过，大家感兴趣的可以看看：八. 实战：CUDA-BEVFusion部署分析-学习BEVPool的优化方案

BEVPool 主要分为两个步骤，我们这边再来回顾下：

首先 BEVFusion 中的 camera backbone 模块有两个头，一个是 camera feature，另一个是 camera depth weights 即 depth 的概率图

得到这两个输出之后我们下一步要干什么呢？我们接下来要将 depth weights 和 camera feature 做一个内积，内积之后我们会得到一个 NxDxHxWxC 维度的特征图，由于这个特征图是具有 3D 信息的，所以论文中比较喜欢用 camera feature point cloud 来表示这个特征。得到这个“伪点云”特征后我们需要把它投影到 BEV 空间上，注意这里面的投影的方法其实是可以通过相机内外参算出来的

那原始的 NxDxHxWxC 大概是 200 万，而投影的 BEV 空间大概是 13 万，要把这 200 万个点给 project 投影到 13 万的 grid 上面去，这个计算量其实还是比较大的

因此 BEVFusion 和 CUDA-BEVFusion 对这一部分做了加速，主要考虑两个方面：

• precomputation
• interval reduction

如果说相机内外参固定，camera 在预测 depth 时 D 的数量固定，那么 bev grid 中每一个 grid 所对应的相机中的坐标也是固定的，不同场景的推理都是一样的，那么 camera 中的特征坐标和 bev 中的 grid 索引是可以提前计算出来的，没有必要去做坐标的映射转换等等操作，可以省去大量计算，这个其实就是 precomputation 预计算的加速

那上面提到的 geometry 其实设置的就是 precomputation 所需要的一些参数

GeometryParameter 结构体的定义如下：

struct GeometryParameter {
  nvtype::Float3 xbound;
  nvtype::Float3 ybound;
  nvtype::Float3 zbound;
  nvtype::Float3 dbound;
  nvtype::Int3 geometry_dim;  // w(x 360), h(y 360), c(z 80)
  unsigned int feat_width;
  unsigned int feat_height;
  unsigned int image_width;
  unsigned int image_height;
  unsigned int num_camera;
};

GeometryParameter 结构体成员变量：

• xbound、ybound、zbound：定义相机感知范围在三维空间中的边界
• dbound：定义了深度（距离）范围，用于处理相机到目标之间的深度信息
• geometry_dim：相机特征投影到 BEV 视角后的特征网格大小，定义了 camera bev feature 的 shape
• image_width、image_height：输入图像的宽度和高度
• feat_width、feat_height：定义了特征图的尺寸
• num_camera：多视角相机数量

配置的 geometry 各个参数中 x 方向和 y 方向范围都是 [-54m, 54m] 覆盖了 108m 的区域，并在该范围内以 0.3m 的刻度间隔进行离散化，z 方向范围是 [-10m, 10m]，不进行离散化仅分为一个区间，相机到目标的距离（深度）范围定义在 [1m, 60m]，步长为 0.5m

预处理后图像大小为 704x256，特征图 feature map 大小设置为 88x32，通过计算 704 / 88 = 8 以及 256 / 32 = 8 可知其经过了 8 倍下采样，相机个数设置为 6 个

相机 2D 图像特征投影到 BEV 特征的 grid size 大小设置为 360x360，每个 grid 用 80 维的特征通道表示，通过计算 108 / 0.3 = 360 也可知 grid 的大小为 360x360

我们接着往下面看：

TransBBox 其实是在 head 部分，这里设置对于最终在 3D 世界坐标下所绘制的 3D bbox（voxel）进行初始化

TransBBoxParameter 结构体的定义如下：

struct TransBBoxParameter {
  std::string model;
  float out_size_factor = 8;
  nvtype::Float2 voxel_size{0.075, 0.075};
  nvtype::Float2 pc_range{-54.0f, -54.0f};
  nvtype::Float3 post_center_range_start{-61.2, -61.2, -10.0};
  nvtype::Float3 post_center_range_end{61.2, 61.2, 10.0};
  float confidence_threshold = 0.0f;
  bool sorted_bboxes = true;
};

TransBBoxParameter 结构体成员变量：

• model：存储 bevfusion head 模型文件的路径
• out_size_factor：缩放因子，定义输出特征图的缩放比例。例如，值为 8 表示 8 倍下采样
• voxel_size：体素大小，定义体素在 x 和 y 方向的尺寸
• pc_range：点云范围，定义 x 和 y 方向的最小值，表示感知区域的边界
• post_center_range_start 和 post_center_range_end：分别定义后处理时边界框中心的最小和最大坐标范围（x、y 和 z 方向上的最小值和最大值）。这个范围通常用于过滤边界框，确保它们位于有效的检测区域内
• confidence_threshold：置信度阈值，通常用于在后处理时过滤低置信度的边界框
• sorted_bboxes：布尔值，表示在输出时是否对边界框进行排序，默认为 true

我们把上面一系列的参数初始化之后我们就可以把它给赋值到 CoreParameter 中去

CoreParameter 结构体的定义如下：

struct CoreParameter {
  std::string camera_model;
  std::string camera_vtransform;
  camera::GeometryParameter geometry;
  camera::NormalizationParameter normalize;
  lidar::SCNParameter lidar_scn;
  std::string transfusion;
  head::transbbox::TransBBoxParameter transbbox;
};

CoreParameter 结构体成员变量：

• camera_model：存储 camera 模型的路径
• camera_vtransform：存储 camera 视角转换的模型路径
• geometry：存储 BEVPool 的 precomputation 预计算需要的参数
• normalize：存储 camera 的归一化参数，用于对输入图像进行预处理操作
• lidar_scn：存储 LiDAR backbone 所需要的相关参数，主要涉及稀疏卷积网络（SCN）的配置
• transfusion：存储多模态融合模块的模型文件路径
• transbbox：存储 BEVFusion Head 模块的模型文件路径

最终我们调用了 bevfusion::create_core 返回了一个指向 Core 对象的智能指针，这里其实体现了 RAII 的接口封装模式。接口类只提供最小限度的接口，实现类通过接口类调用实现封装上的隐蔽

create_core 函数定义如下：

上面这种写法其实是 CUDA-BEVFusion 框架中常用的一个写法，它其实就是 RAII+接口类的一个实现方法，RAII 资源获取即初始化主要体现在我们创建了一个 instance，接着马上通过 init 函数做了一个初始化，最后返回这个 instance

我们在框架中还可以看到很多类似的写法，包括 create_backbone、create_bevpool、create_transfusion 等等，那都是相同的写法，都是命名空间下对外部提供的的一个 create_xxx 接口函数，我们记住这个就好了

create_core 函数通过智能指针管理 Core 对象生命周期，并在初始化失败时安全地释放资源，那这里其实体现了RAII+接口模式

RAII 主要体现在 init 函数，通过 init 函数成功获取实例对象，如果初始化失败则返回的就是一个空指针，充分体现了资源获取即初始化的概念，此外使用 shared_ptr 智能指针来管理 CoreImplement 对象，在不再有任何引用时自动销毁该对象，这进一步强化了 RAII 模式

接口模式则是通过定义抽象接口来分离实现和使用，这里 Core 类定义了一个抽象接口，CoreImplement 类实现了这个接口

Core 是一个纯虚类（抽象类），定义了若干虚函数，用于表示 CUDA-BEVFusion 核心功能的接口：

class Core {
 public:
  virtual ~Core() = default;
  virtual std::vector<head::transbbox::BoundingBox> forward(const unsigned char** camera_images, const nvtype::half* lidar_points, int num_points, void* stream) = 0;
  virtual std::vector<head::transbbox::BoundingBox> forward_no_normalize(const nvtype::half* camera_normed_images_device, const nvtype::half* lidar_points, int num_points, void* stream) = 0;
  virtual void print() = 0;
  virtual void set_timer(bool enable) = 0;
  virtual void update(const float* camera2lidar, const float* camera_intrinsics, const float* lidar2image, const float* img_aug_matrix, void* stream = nullptr) = 0;
  virtual void free_excess_memory() = 0;
};

这里的纯虚函数（= 0）规定了 Core 对象的接口，使得使用 Core 的任何代码都不需要知道具体实现的细节，而只需要通过这个接口与之交互。这是一种典型的接口隔离原则，即调用者只需要知道接口，而不关心具体实现。

CoreImplement 类实现了 Core 接口，提供了具体的功能实现：

class CoreImplement : public Core {
  // 实现 Core 中定义的虚函数
};

通过继承 Core 类并实现其纯虚函数，CoreImplement 具备了完整的功能，并且可以通过 Core 接口进行访问

在 create_core 函数中，充分体现了 RAII 和接口模式的结合，通过 RAII 管理资源，通过接口模式进行功能的抽象和隔离：

std::shared_ptr<CoreImplement> instance(new CoreImplement());
if (!instance->init(param)) {
    instance.reset();
}
return instance;

• RAII 的内存管理使得 CoreImplement 对象的生命周期和资源管理绑定在一起，确保不发生内存泄漏
• 接口模式则通过返回 Core 的指针，实现了具体功能与调用方之间的解耦，调用者只需要与 Core 接口进行交互，而不需要知道背后的 CoreImplement 是如何实现的

那大家如果对 tensorRT_Pro 熟悉的话，可能知道 RAII+接口封装模式其实是杜老师在代码中常用的，杜老师在课程中也专门有提到过，大家感兴趣的可以看看：7.4.tensorRT高级(2)-使用RAII接口模式对代码进行有效封装

2.2 init函数

下面我们来看下 init 初始化都做了些什么

我们可以看到这里也是 RAII+接口封装模式的体现，这里面有很多 create 函数，从各个接口类/命名空间，我们可以访问到的只有 create_xxx，也就是所谓的“资源获取即初始化”。根据参数进行每一个部分的初始化，主要就是分配内存空间、创建 tensorrt engine 等等，做初始化的包括以下几个部分：

• create_camera_backbone：命名空间 camera
• create_bevpool：命名空间 camera
• create_vtransform：命名空间 camera
• create_transfusion：命名空间 fuser
• create_transbbox：命名空间 head
• create_scn：命名空间 lidar
• create_normalization：命名空间 camera
• create_depth：命名空间 camera
• create_geometry：命名空间 camera

既然初始化是做资源分配以及 engine 创建等工作，下面我们就来看看具体是怎么做的

我们先从 camera_backbone 看起：

这里又是之前看到的写法，RAII+接口类实现类，我们来看看 BackboneImplement 实现类中的 initialization 具体是怎么做的

首先通过 TensorRT 命名空间下的 load 函数加载 engine，这里主要是调用 TensorRT 的接口类实现推理引擎的初始化，我们从这里可以访问到的只有 load，避免了一些不可预测的行为

那 load 函数这里也是一个 RAII+接口封装模式，调用 EngineImplement 中的 load 函数读取模型文件反序列化成 engine

那 load 完成之后我们就拿到了反序列化的 engine，通过 engine 我们可以获取到一些 dims 信息，这里的 static_dims 就是我们在创建 camera backbone 时 bindings 的维度，我们知道 camera 的 backbone 输入和输出都是两个，所以这里 static_dims 中的 0，1 是 camera 和 depth 的输入维度，2，3 是 backbone 输出得到的 camera 和 depth 的特征图维度

volumn 其实就是指输出的 camera 和 depth 中的元素个数分别为 NxCxHxW 和 NxHxWxD，用 accumulate 实现计算

最后我们可以得到 camera_shape_ 维度是 [N，C，D，H，W]，我们可以理解 camera_backbone 的初始化主要是做分配资源、反序列引擎以及设置一些成员变量，那设置的这些成员变量其实是跟 forward 有关的，我们后续 forward 时都是根据这些成员变量去推理的

这是 camera_backbone 的初始化，我们来看 bevpool 的初始化：

bevpool 的 initialize 主要是对一些成员变量进行初始化，分配资源

我们接着往下看 vtransform：

我们在 八. 实战：CUDA-BEVFusion部署分析-分析BEVFusion中各个ONNX 中有提到 vtransform 就是一个 DNN，要做的初始化也比较简单，和 camera_backbone 一样，通过 load 函数得到一个 engine，接着得到 engine 的输出维度，分配资源

接着看 transfusion：

transfusion 就是一个 Fusion 融合模块，它也是一个 DNN，它的初始化也是 load 一个引擎接着得到它的 output dims，计算 volume 最后 cudaMalloc 分配资源

这里有个点需要大家注意，我们从下面 Fusion 的 ONNX 中可以看出它有两个输入分别是 camera bev feature 和 lidar bev feature 以及一个输出即融合后的 bev feature

我们接着往下看 transbbox：

transbbox 其实就是 BEVFusion 中的 head 模块，它也是一个 DNN，通过 load 加载 engine，接着为 head 的输出做一个内存分配

下面来看 scn：

值得注意的是 SCN 这里面又套了一个初始化，这个初始化是 voxelization 体素化的初始化，我们先来看下 voxelization 的初始化：

这里的东西比较多，也主要是做了一个内存分配，因为体素化这边是做了一个 CUDA 加速，那体素化的 CUDA 加速我们后面会讲，其实会利用 hash_table 来保存每一个 atomic operator 的 id，

还有一点需要注意的是虽然 SCN 是一个 DNN，但是这边并不是通过 TensorRT::load 加载的，我们前面也提到过 spconv 的特殊性 TensorRT 无法推理，这里调用的 spconv 命名空间下的 load_engine_from_onnx 来反序列化 engine 的，这部分没有开源，我们这里看不见具体的实现

我们接着看 normalization：

那 normalization 部分不是一个 DNN，主要是 camera 图像的前处理部分的一些初始化，主要是对成员变量进行了一些初始化，由于前处理我们是放在 CUDA 上加速实现的，所以这里在初始化时还需要分配 device 内存

接着往下看 depth：

这里的 depth 主要是负责 lidar 点云到 camera 的投影，我们知道 camera backbone 有两个一个是 image 另外一个就是 depth，输入的 depth 维度是 1x1x256x704 也就是每个 camera 都有一个 256x704 大小的 depth image，那 depth image 就是深度图，把点云投影到图像上

那这个过程 CUDA-BEVFusion 框架中也是做了一些 CUDA 加速，意味着我们需要分配对应的 device 内存，包括相机的内外参等等，这个我们后面会跟大家去讲

我们接着往下看 geometry：

geometry 就是为了 camera 到 bev 中 BEVPool 的 precomputation 预计算加速，重点是 intervals 以及 interval_starts，这个部分比较重要

intervals 就是做一个规约，我们前面说 camera 的 200 万个点需要投影到 13 万的 bev grid 上面去，那也就是意味着 bev 中每一个 grid 会有多个 camera feature point cloud 的点，既然每个 grid 里面有多个点，我们需要做一个处理，怎么做呢？

如上图所示，我们可以将映射到每一个 grid 里的所有点进行相加，就可以只剩下一个点，也就是 1xC 个数据，那么结合 grid 的数量，最终形成的 bev feature 的大小就是 1xCxBHxBW

比如图中有四个点落在 grid0，两个点落在 grid1，每个点有 C 维数据（80），总共有 N 个 grid，我们对每个 grid 的数据都做一个相加，最后每个 grid 可以得到一个 1xC 维的数据，这个就是 intervals 需要做的

对于求和的部分，我们可以使用 CUDA 加速，让每一个 thread 负责一个 grid 里的所有计算，包括求和、内积这些用一个 kernel 完成

理解了 intervals 之后我们还需要理解 interval_starts 是什么，这个也非常的重要，由于我们是让每一个 thread 负责处理一个 grid，所以我们需要知道每一个 grid 对应的起始的 intervals 是多少，这样我们可以直接通过索引去内存地址中找到对应的点的数据

这个部分有点类似于我们在 2D Detection 任务的 decode 解码中的实现，每个 bbox 假设是 N 维的，那么我们通过 blockDim、blockIdx 以及 threadIdx 就可以找到该线程的全局索引 idx，通过 idx 乘以 N 我们就可以在访问到内存中的第 N 个 bbox 的信息了

具体是不是这样的呢，这个就需要分析它具体的实现代码了，这个我们后续再详细分析，这里大家需要先有个基本概念

在 geometry 的初始化中有一些成员变量的初始化和内存的分配，我们这里可以先做简单了解，首先 numel_frustum_ 变量中的 frustrum 我们可以理解为视锥，对于 88x32 的feature map 大小中的每一个点，我们都有 D 个深度估计的值，其中 BEVFusion 中的 D 设置为 118，意味着 numel_frustum_ 大小为 88 * 32 * 118

大家可能对视锥这个概念不是很熟悉，这边老师也找了一张示例图来说明：

在 BEVFusion 中视锥（view frustum）是一个关键概念，用于将相机的二维图像信息映射到三维空间，以实现多视角特征的融合。如上图所示，视锥是从相机视点出发，沿着视线方向形成的一个锥形或棱锥形空间体积，它定义了相机所能“看到”的空间范围内，即相机的可视范围或视野，视锥由以下几个部分构成：(from ChatGPT)

• 视点（Camera Position）：视锥的顶点，代表相机的位置
• 视平面（View Plane）：视锥的底面，通常与相机的成像平面对应
• 视角（Field of View，FOV）：视锥的张角，决定了相机的广角程度
• 近平面和远平面（Near and Far Planes）：限制视锥的深度范围，只考虑这两个平面之间的物体

在 BEVFusion 中，视锥用于将二维图像特征映射到三维 BEV 空间，主要作用包括：

• 空间投影：通过视锥，可以将图像像素对应到三维空间中的射线。结合深度信息，可以确定每个像素在三维空间中的位置
• 特征融合：利用多个相机的视锥，将不同视角的图像特征投影到用一个 BEV 空间，实现多视角特征的融合
• 可视性判定：通过视锥，可以判断 BEV 空间中的某个点是否在相机的视野内，便于筛选有效的特征信息

视锥可以通过相机的内参矩阵和外参矩阵进行数学描述：

• 内参矩阵（Intrinsic Matrix）：描述相机的焦距、主点偏移等参数，决定了像素坐标与相机坐标系之间的转换关系
• 外参矩阵（Extrinsic Matrix）：描述相机在世界坐标系中的位置和朝向，决定了相机坐标系与世界坐标系之间的转换关系

利用这些参数，可以将图像中的像素坐标映射到三维空间中的射线，从而构建视锥，通过预测的离散深度概率可以将图像特征正确地投影到 BEV 空间，实际应用中的流程主要包括：

• 1. 图像特征提取：从相机图像中提取特征，得到二维的特征图
• 2. 构建视锥：利用相机的内外参和空间边界参数，构建每个相机的视锥
• 3. 特征投影：将二维图像特征通过视锥投影到三维 BEV 空间，得到对应的空间特征
• 4. 多视角融合：将多个相机的空间特征在 BEV 空间中进行融合，形成统一的鸟瞰视图
• 5. 后续处理：对融合后的 BEV 特征进行检测、分割等任务

视锥在 BEVFusion 中起到了桥梁的作用，将相机的二维图像特征与三维 BEV 空间联系起来。通过视锥，可以精确地将图像中地信息映射到三维空间，为多视角特征融合奠定了基础，这对于实现准确的环境感知和目标识别非常重要。

这里其实体现的就是 BEVFusion 中的视角转换模块的设计部分：

在上面 BEVFusion 的框图中我们可以看到 6 个不同视角的图像经过一个 2D Backbone 之后提取图像特征，那这个 2D Backbone 可以是 ResNet、Swin Transformer 等等都行，提取完特征之后我们需要通过视角转换模块把 Camera Feature 投影到 BEV 上，那这里其实会出现一个问题，我们在柒柒老师的课程 二. BEV感知算法基础模块讲解 中讲过，2D 到 3D 的投影其实是一条射线，图像中的每一个像素与空间中的一条射线是对应的，是点线的对应关系，那具体投影到射线上的哪一个点呢，我们无从得知，假设 BEV Grid 的大小是 180x180，那么每一个 grid 对应的是哪个 camera 上的哪一个坐标呢？这个其实是大家需要考虑的事情

另外我们都知道 camera 传感器拍摄的图片其实是没有深度信息的，我们从相机的小孔成像原理中可以知道沿着光线的角度去看一个物体的时候，它其实无论是在远处还是在近处，在图片上看到的都是一个点，那这个就导致图片的深度信息其实是丢失的，那么如果深度信息是丢失的话，我们要怎么去学习它，我们要怎么去把这个深度信息给拿到之后去做一个 BEV 上的投影呢，这个其实在 LSS 这篇文章里面就有讲。

首先第一步先做离散的深度分布，将像素的特征按照概率分布可以映射到不同的深度值上，那这个其实是对 3D 空间特征利用图像特征去进行一个补全的一个过程，当补完之后我们再拍扁到 BEV 空间去做预测，那这就是一个完整的流程

那 BEVFusion 中一共有 6 个视锥，可以拼接得到全景图，6 个视锥组合在一起形成一个 geometry，geometry 的大小就是 6 * 88 * 32 * 118

geometry 初始化最后调用了一个核函数 create_frustum_kernel：

static __global__ void create_frustum_kernel(unsigned int feat_width, unsigned int feat_height, unsigned int D,
                                             unsigned int image_width, unsigned int image_height, float w_interval,
                                             float h_interval, nvtype::Float3 dbound, float3* frustum) {
  int ix = cuda_2d_x;
  int iy = cuda_2d_y;
  int id = blockIdx.z;
  if (ix >= feat_width || iy >= feat_height) return;

  unsigned int offset = (id * feat_height + iy) * feat_width + ix;
  frustum[offset] = make_float3(ix * w_interval, iy * h_interval, dbound.x + id * dbound.z);
}

这个 kernel 核函数是用来创建 frustum（视锥），frustum 中一共有 88x32x118 个数据，每一个数据都是 float3，分别代表的是原图大小中的 x 坐标，y 坐标以及以 0.5m 为刻度的深度

这个就是 geometry 的初始化，内容比较多但是也比较重要

整个初始化的最后就是一些成员变量的赋值以及内存的分配

至此，CUDA-BEVFusion 的 initialize 初始化部分就分析完成了，主要是对 CoreImplement 中的 init 函数进行分析

总结

本次课程我们学习了 CUDA-BEVFusion 中的初始化部分，主要是通过 RAII 接口模式的方式对各个模块进行初始化，包括 camera_backbone、bevpool、scn 等等，其中比较重要的是 geometry 部分，它主要是为 BEVPool 的 precomputation 预计算初始化一些参数。大部分的代码分析工作都是 ChatGPT 协助完成的，博主这边笔记呈现效果可能不是很好，需要大家自己耐心调试分析

OK，以上就是第 9 小节有关 CUDA-BEVFusion 中初始化的代码分析的全部内容了，下节我们来讲下坐标系转换以及 BEVPool 中的 precomputation 具体是怎么做的，敬请期待😄

下载链接

• 论文下载链接【提取码：6463】
• 数据集下载链接【提取码：data】
• 代码和安装包下载链接【提取码：cuda】

参考

• 八. 实战：CUDA-BEVFusion部署分析-环境搭建
• 八. 实战：CUDA-BEVFusion部署分析-学习BEVPool的优化方案
• 八. 实战：CUDA-BEVFusion部署分析-分析BEVFusion中各个ONNX
• 八. 实战：CUDA-BEVFusion部署分析-学习CUDA-BEVFusion推理框架设计模式
• AutoCV第八课：3D基础
• 7.4.tensorRT高级(2)-使用RAII接口模式对代码进行有效封装
• 二. BEV感知算法基础模块讲解
• ppt画深度学习网络图
• PPT（Visio）如何画立方体及多个小立方体堆叠效果 - 科研绘图