4、计算单元详解（GPU）：GPU并行计算原理、GPU在感知与渲染中的应用、典型架构（如NVIDIA Ampere/Ada）、CUDA与Tensor Core

好，咱们今天聊聊GPU。说实话，在自动驾驶芯片里，GPU的地位有点像「万能胶」——哪里需要高性能并行计算，哪里就有它。我这些年经手的几个项目，从早期的Mobileye到后来的Orin，GPU的选型和优化永远是绕不开的坎。

4.1 GPU并行计算原理：别把它当CPU用

很多人刚接触GPU时，容易犯一个错——拿CPU的思维去理解GPU。CPU是「大力士」，一个核心能处理极其复杂的任务；GPU是「蚂蚁军团」，成千上万个简单核心一起上。

GPU的并行计算，核心就三个字：SIMT（单指令多线程）。什么意思？就是一条指令下去，几百上千个线程同时执行。你想想看，处理一张1920x1080的图像，每个像素点做同样的滤波操作，CPU得一个个来，GPU直接一波带走。

关键指标：

  • CUDA核心数：不是越多越好，但少了肯定不行
  • 显存带宽：GPU的命脉，我见过太多项目因为带宽瓶颈翻车
  • Occupancy（占用率）：SM上活跃的线程占比，直接影响吞吐

我在做第一代自动驾驶芯片选型时，就吃过亏。当时只看算力（TOPS），没仔细看显存带宽。结果跑BEV感知模型时，数据搬来搬去，GPU核心大部分时间在「等菜下锅」。嗯，从那以后，我选型必看带宽和算力的匹配度。

4.2 GPU在感知与渲染中的应用

自动驾驶里，GPU主要干两件事：感知和渲染。听起来风马牛不相及，但底层逻辑是一样的——大量数据并行处理。

4.2.1 感知：从像素到语义

感知这块，说白了就是让车「看懂」世界。摄像头拍到的原始图像，经过GPU加速的CNN（卷积神经网络），变成车道线、行人、车辆等语义信息。

我建议你重点关注这几个算子：

  • 卷积运算：GPU的强项，im2col+GEMM是经典套路
  • 池化/归一化：数据量小但频繁，注意访存优化
  • 后处理：NMS（非极大值抑制）在GPU上实现时，要小心线程同步

举个例子，YOLOv5的推理流程：

// 伪代码示意
cudaMemcpy(input, image, size, cudaMemcpyHostToDevice);
infer_kernel<<<grid, block>>>(input, output);
// 这里要注意grid和block的配置，我习惯让block为256
cudaMemcpy(result, output, size, cudaMemcpyDeviceToHost);

我曾经在一个项目里，因为block size设得太小（32），导致SM利用率只有30%。后来改成256，直接翻倍。你想想看，同样的硬件，换个参数就差这么多。

4.2.2 渲染：从数据到可视化

渲染这块，很多人觉得「不就是显示个画面吗？」其实没那么简单。自动驾驶的渲染包括：

  • 鸟瞰图（BEV）渲染：把多摄像头数据融合成俯视图
  • 3D场景重建：激光雷达点云+相机图像的融合渲染
  • HMI界面：仪表盘、中控屏的实时显示

渲染对GPU的要求和感知不一样。感知更看重计算吞吐，渲染更看重延迟。你想想看，如果驾驶员打方向盘，画面延迟了100ms，那体验得多糟糕。

我的经验：渲染管线里，最容易被忽视的是「帧同步」。我曾经遇到一个bug，GPU渲染线程和CPU控制线程不同步，导致画面撕裂。解决方案是用V-Sync或者双缓冲。

4.3 典型架构：NVIDIA Ampere vs Ada

说到GPU架构，NVIDIA是绕不开的。咱们重点看两代：Ampere（安培）和Ada Lovelace（阿达·洛芙莱斯）。

特性 Ampere（GA102） Ada（AD102）
CUDA核心 8704 16384
Tensor Core 272（第三代） 512（第四代）
显存带宽 760 GB/s 1008 GB/s
制程 8nm 4nm
关键特性 稀疏化支持 光流加速器、着色器重排序

从Ampere到Ada，最直观的变化是规模翻倍。但我觉得，真正重要的是架构微调：

  • SM（流式多处理器）结构优化：Ada的SM里，CUDA Core和Tensor Core的调度更灵活
  • 显存子系统改进：L2缓存从6MB增加到72MB（AD102），这对自动驾驶的BEV模型特别友好
  • 光流加速器：虽然主要面向游戏，但做目标跟踪时也能用上

我个人建议，如果做L2+级别的自动驾驶，Ampere够用；但要做L4，尤其是需要实时处理多传感器融合的，直接上Ada。

4.4 CUDA与Tensor Core

CUDA是NVIDIA的「护城河」。说白了，它就是一套让开发者能用C++写GPU程序的工具链。我刚开始学CUDA时，觉得不就是多线程编程吗？后来才发现，坑多着呢。

4.4.1 CUDA编程模型

核心概念就三个：

  • Grid：一个任务的总线程块集合
  • Block：一组线程，共享shared memory
  • Thread：最小的执行单元

写CUDA代码时，我习惯先画个图：

// 典型的CUDA kernel
__global__ void process_image(float* input, float* output, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (x < width && y < height) {
        // 每个线程处理一个像素
        int idx = y * width + x;
        output[idx] = some_operation(input[idx]);
    }
}

这里有个坑：线程发散。如果if-else分支太多，同一个warp里的线程走不同路径，性能会急剧下降。我见过一个项目，因为后处理里加了大量条件判断，性能直接腰斩。

避坑指南：我曾经在写NMS时，用了大量的原子操作（atomicAdd），结果发现GPU的原子操作效率极低。后来改成先局部排序再合并，性能提升了5倍。

4.4.2 Tensor Core：AI计算的「涡轮增压」

Tensor Core是NVIDIA专门为深度学习设计的矩阵计算单元。它支持混合精度计算（FP16/INT8/INT4），吞吐量是普通CUDA Core的几十倍。

为什么Tensor Core这么强？因为它做的是矩阵乘加（D = A * B + C），一次指令完成4x4矩阵的运算。而普通CUDA Core得一条条指令来。

在自动驾驶里，Tensor Core主要用在：

  • 卷积层加速：把卷积转成矩阵乘法，用Tensor Core算
  • 全连接层：BEV模型的Transformer里，QKV矩阵乘法全靠它
  • 量化推理：INT8推理时，Tensor Core的吞吐是FP16的两倍

我建议你使用cuBLAS或cuDNN库，别自己手写Tensor Core的调用。这些库已经优化到极致了。我自己试过手写，性能差了一倍不止。

实用技巧：用Tensor Core时，注意矩阵维度要对齐到16的倍数（FP16）或32的倍数（INT8）。不对齐的话，Tensor Core会退化成普通CUDA Core，性能损失很大。

最后说一句，GPU的优化是个「无底洞」。我做了这么多年，每次以为自己懂了，下一个项目又会有新坑。但核心思路不变：理解硬件特性，让软件去适配硬件，而不是反过来。你想想看，是不是这个理？

← 上一章 📖 返回目录 下一章 →