第2章 GPU架构基础

GPU架构基础：GPU并行计算原理、CUDA核心与Tensor Core、GPU内存层次结构

好，咱们今天聊聊GPU。说实话，在ADAS系统里，GPU是个绕不开的角色。很多人觉得GPU就是用来做图形渲染的，其实在自动驾驶领域，GPU的并行计算能力才是真正的杀手锏。我个人习惯把GPU比作一个「超级包工头」——它手下有成千上万个工人，每个工人只干一件简单的事，但架不住人多，一起干起来效率惊人。

GPU并行计算原理：从串行到并行的思维转变

CPU和GPU的设计哲学完全不同。CPU是「全能冠军」，什么活都能干，但一次只能干几件事。GPU是「专项选手」，只会干简单的数学运算，但一次能同时干成千上万件事。

为什么会这样？你想想看，CPU需要处理各种复杂的逻辑判断、分支预测、乱序执行，所以芯片面积大部分被控制单元和缓存占用了。而GPU呢？它把90%以上的晶体管都用来做计算单元，控制单元极其简单。说白了，GPU就是「以量取胜」。

我在项目中遇到过这样一个场景：处理一张1080p的图像，要做像素级的预处理。CPU跑起来要几十毫秒，而GPU只需要几毫秒。为什么？因为图像处理天然就是并行的——每个像素的计算互不依赖，完全可以同时进行。

GPU的并行计算模型叫做SIMT（单指令多线程）。什么意思？就是一条指令，控制多个线程同时执行。比如你要给1000个像素每个加1，GPU只需要发一条「加1」指令，1000个线程同时执行。CPU呢？得循环1000次。

核心要点：GPU并行计算的关键在于「数据并行」——把大任务拆成无数个小任务，每个小任务独立且相同，然后让成千上万个线程同时处理。

CUDA核心：GPU的「基本工人」

CUDA核心是NVIDIA GPU中最基本的计算单元。每个CUDA核心可以执行一个浮点或整数运算。一个现代GPU有几千个CUDA核心，比如NVIDIA Orin芯片有2048个CUDA核心。

但这里有个坑——CUDA核心并不是独立工作的。它们被组织成SM（Streaming Multiprocessor，流式多处理器）。每个SM包含一组CUDA核心、共享内存、寄存器等资源。我刚开始做CUDA编程时，以为线程越多越好，结果发现SM的资源是有限的，线程太多反而会因为资源竞争导致性能下降。

来看一个简单的CUDA代码示例，感受一下并行计算的样子：

// 一个简单的向量加法kernel
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

// 调用方式
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N);

你看，每个线程只处理一个元素，互不干扰。这就是GPU编程的基本思路——把问题分解成大量的小任务。

避坑指南：我曾经犯过一个错误——在kernel里用了太多if-else分支。GPU的SIMT模型最怕分支发散，同一个warp（32个线程）里如果线程走不同的分支，性能会急剧下降。所以写GPU代码时，尽量让所有线程走相同的路径。

Tensor Core：专为AI计算打造的「特种兵」

CUDA核心虽然厉害，但做矩阵乘法这种深度学习核心运算时，效率还是不够高。于是NVIDIA推出了Tensor Core——专门为矩阵运算设计的硬件单元。

Tensor Core能在单个时钟周期内完成一个4x4矩阵的乘加运算（D = A * B + C）。这比用CUDA核心实现快得多。我记得第一次在ADAS项目里用Tensor Core做神经网络推理时，速度提升了将近4倍。

Tensor Core支持多种精度：FP16、BF16、INT8、INT4等。在ADAS系统中，我们经常用INT8量化后的模型做推理，因为精度损失可以接受，但速度提升非常明显。

使用Tensor Core其实不需要手动写底层代码，CUDA的cuBLAS库和cuDNN库会自动调用。但如果你想手动控制，可以用下面的方式：

// 使用Tensor Core的矩阵乘法示例（通过cuBLAS）
cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N,
             m, n, k, &alpha,
             A, CUDA_R_16F, lda,
             B, CUDA_R_16F, ldb,
             &beta,
             C, CUDA_R_16F, ldc,
             CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP);

嗯，这里要注意：Tensor Core对数据对齐有要求，矩阵的维度最好是8的倍数。否则性能会大打折扣。

计算单元	适用场景	精度支持	相对性能
CUDA Core	通用并行计算、图像处理	FP32/FP64/INT32	1x（基准）
Tensor Core	深度学习推理/训练	FP16/BF16/INT8/INT4	4-8x（矩阵运算）

GPU内存层次结构：数据搬运的艺术

GPU编程中，最容易被忽视的就是内存管理。我见过太多人把时间花在优化计算上，结果瓶颈在内存带宽上。说白了，GPU计算就像做饭——计算单元是厨师，内存是食材仓库。厨师再快，如果食材搬不过来，也是白搭。

GPU的内存层次结构从快到慢、从大到小依次是：

寄存器（Register）：每个线程私有，速度最快，容量最小（每个线程最多255个32位寄存器）。
共享内存（Shared Memory）：同一个Block内的线程共享，速度接近寄存器，容量有限（通常几十KB）。
L1/L2缓存：自动管理，对程序员透明。
全局内存（Global Memory）：所有线程都能访问，容量最大（几GB到几十GB），但速度最慢（延迟几百个时钟周期）。
常量内存（Constant Memory）：只读，有缓存，适合存储常量数据。
纹理内存（Texture Memory）：专门为图像数据优化，支持插值等操作。

我建议你在写GPU代码时，始终记住一个原则：尽量减少全局内存的访问次数。能放共享内存的就放共享内存，能合并访问的就合并访问。

警告：全局内存的访问模式非常关键。如果线程访问的地址是连续的（合并访问），带宽利用率很高。如果是随机访问，带宽利用率可能只有1%。我曾经在项目中因为数据排列不当，导致全局内存带宽利用率只有5%，优化后提升到了80%以上。

来看一个利用共享内存优化的例子：

// 使用共享内存做矩阵转置
__global__ void transpose(float* in, float* out, int width) {
    __shared__ float tile[TILE_DIM][TILE_DIM];
    
    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;
    
    // 从全局内存加载到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 从共享内存写回全局内存（转置后）
    out[x * width + y] = tile[threadIdx.y][threadIdx.x];
}

你看，这里先用共享内存做了一次中转，避免了直接对全局内存进行非合并访问。性能提升非常明显。

总结一下

GPU架构的核心就三点：

并行计算原理：数据并行，SIMT模型，大量线程同时工作。
CUDA Core vs Tensor Core：前者是通用计算单元，后者是AI专用加速器。
内存层次结构：用好共享内存和寄存器，减少全局内存访问。

在ADAS系统中，GPU通常用来做图像预处理、特征提取、以及部分神经网络推理。但要注意，GPU的功耗比较高，在嵌入式平台上需要权衡性能和功耗。我一般建议：能用NPU做的推理尽量用NPU，GPU负责那些NPU不擅长的任务，比如图像拼接、光流计算等。

下一章我们会深入聊聊NPU的架构，看看它和GPU有什么不同，以及如何在ADAS系统中合理分配任务。

📚 课程目录

GPU架构基础：GPU并行计算原理、CUDA核心与Tensor Core、GPU内存层次结构

GPU并行计算原理：从串行到并行的思维转变

CUDA核心：GPU的「基本工人」

Tensor Core：专为AI计算打造的「特种兵」

GPU内存层次结构：数据搬运的艺术

总结一下