📚 课程目录第2章 GPU计算基础

GPU计算基础：GPU架构、线程层次、内存层次、CUDA编程模型与第一个CUDA程序

在深度学习推理引擎的部署链路中，GPU 早已不是单纯的图形加速器，而是大规模并行计算的核心载体。TensorRT 之所以能实现毫秒级甚至微秒级的推理延迟，其底层依赖的正是对 GPU 硬件资源（SM、显存带宽、线程调度）的极致压榨。理解 GPU 的物理架构与 CUDA 的逻辑抽象模型，是编写高效 Kernel、诊断性能瓶颈、以及正确配置 TensorRT 引擎的基石。本章将从硬件架构出发，逐步拆解线程与内存的层次关系，并通过一个完整的向量加法案例，带你完成从编写、编译到运行的第一个 CUDA 程序。

上图清晰地展示了 GPU 计算的两层视角：物理硬件层由多个 SM（流式多处理器）构成，每个 SM 内部包含 CUDA 核心、共享内存、寄存器文件和 Warp 调度器；而 CUDA 逻辑层则通过 Grid、Block、Thread 三级结构组织并行任务。硬件与逻辑层之间并非一一对应，而是由 GPU 的硬件调度器动态地将 Block 分配到空闲的 SM 上执行。这种“逻辑并行”与“物理调度”的解耦，正是 CUDA 能够透明扩展至不同 GPU 型号的关键设计。

2.1 GPU 架构核心原理：从图形加速器到通用并行处理器

现代 GPU（如 NVIDIA Ampere、Hopper 架构）本质上是一个由数千个 ALU（算术逻辑单元）组成的吞吐量优先处理器。与 CPU 的“低延迟、大缓存、复杂控制”设计哲学不同，GPU 将芯片面积的绝大部分用于计算单元和带宽优化，而非分支预测与乱序执行。以 A100 GPU 为例，其内部包含 108 个 SM，每个 SM 拥有 64 个 FP32 CUDA 核心，总计 6912 个核心。这种大规模并行度使得 GPU 在矩阵乘法、卷积等数据并行任务上，能够达到数十 TFLOPS 的浮点算力。

理解 GPU 架构的关键在于三个核心概念：SIMT（单指令多线程）、Warp 调度与Occupancy（占用率）。SIMT 模型下，32 个线程组成一个 Warp，它们在同一时刻执行同一条指令，但操作各自的数据。Warp 是 GPU 调度的最小单位，SM 上的 Warp Scheduler 负责在多个 Warp 之间切换，以隐藏内存访问延迟。当某个 Warp 因等待全局内存数据而阻塞时，调度器会立即切换到另一个就绪的 Warp，从而保持计算单元持续忙碌。这种“延迟隐藏”机制是 GPU 高性能的核心，也是我们在编写 Kernel 时需要考虑的关键因素。

特性	CPU (Intel Xeon)	GPU (NVIDIA A100)
核心数	24-40 个高性能核心	6912 个 CUDA 核心
设计目标	低延迟、复杂逻辑、分支预测	高吞吐、数据并行、延迟隐藏
缓存策略	大容量 L2/L3 缓存 (数十MB)	小容量 L2 (40MB)，依赖共享内存
内存带宽	~200 GB/s (DDR5)	~2 TB/s (HBM2e)
调度单位	线程 (OS 调度)	Warp (硬件调度)

在深度学习推理场景中，TensorRT 正是利用了 GPU 的高吞吐特性。例如，在批量推理（Batch Inference）时，TensorRT 会将多个输入样本打包，让 GPU 的数千个核心同时处理不同样本的同一层计算。这种数据并行模式与 GPU 的 SIMT 架构天然契合，也是为什么 GPU 在推理任务中远优于 CPU 的根本原因。

2.2 CUDA 线程层次与内存层次：逻辑组织的双维度

CUDA 将并行任务抽象为三层结构：Grid（网格）→ Block（块）→ Thread（线程）。一个 Grid 包含多个 Block，每个 Block 包含多个 Thread。这种层次结构并非随意设计，而是与硬件资源紧密对应：Block 是分配给 SM 的最小单位，一个 SM 可以同时驻留多个 Block；Thread 则是实际执行计算的最小单位。在编程时，我们通过 blockIdx、threadIdx 和 blockDim 三个内置变量来索引每个线程的全局位置。

内存层次同样遵循由快到慢、由小到大的原则：寄存器（Registers）→ 共享内存（Shared Memory）→ L2 缓存 → 全局内存（Global Memory）。寄存器是每个线程私有的，访问延迟为 1 个时钟周期；共享内存是 Block 内所有线程共享的，延迟约为 20-30 个周期，容量有限（通常 48KB-100KB）；全局内存则是所有线程共享的，容量最大（数十 GB），但延迟高达 400-800 个周期。此外，还有只读的常量内存（Constant Memory）和纹理内存（Texture Memory），它们在特定场景下能提供缓存优化。

内存选择策略：在编写 Kernel 时，应优先使用寄存器存储频繁访问的局部变量；当 Block 内线程需要数据交换时，使用共享内存；全局内存则用于存储输入输出数据。避免在循环体内频繁访问全局内存，而应先将数据加载到共享内存或寄存器中，计算完成后再写回。这一策略在 TensorRT 的 Kernel 融合（Kernel Fusion）中同样适用——通过减少中间结果的全局内存读写，大幅提升推理速度。

线程层次与内存层次的配合决定了 Kernel 的性能上限。例如，一个 Block 配置为 256 个线程时，每个 SM 可以驻留多个 Block（取决于寄存器与共享内存的使用量）。如果每个线程使用过多寄存器，SM 上可驻留的线程总数就会减少，导致 Occupancy 降低，无法有效隐藏内存延迟。因此，在编写 CUDA 程序时，需要权衡线程数、寄存器使用量与共享内存分配，以达到最佳的硬件利用率。

2.3 第一个 CUDA 程序：向量加法的完整实战

下面我们通过一个经典的向量加法案例，完整演示 CUDA 程序的编写、编译与执行流程。该案例将两个长度为 N 的浮点数组相加，输出结果数组。虽然功能简单，但它涵盖了 CUDA 编程的所有核心要素：内存分配、数据拷贝、Kernel 启动、线程索引计算与结果回传。

// vector_add.cu
#include <stdio.h>
#include <cuda_runtime.h>

// Kernel 函数：每个线程处理一个元素
__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int N = 1 << 20;  // 1,048,576 个元素
    size_t bytes = N * sizeof(float);

    // 1. 分配主机内存并初始化
    float *h_A = (float*)malloc(bytes);
    float *h_B = (float*)malloc(bytes);
    float *h_C = (float*)malloc(bytes);
    for (int i = 0; i < N; i++) {
        h_A[i] = rand() / (float)RAND_MAX;
        h_B[i] = rand() / (float)RAND_MAX;
    }

    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    // 3. 将数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    // 4. 配置 Kernel 启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    // 5. 启动 Kernel
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 6. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    // 7. 验证结果
    bool correct = true;
    for (int i = 0; i < N; i++) {
        if (fabs(h_C[i] - (h_A[i] + h_B[i])) > 1e-5) {
            correct = false;
            break;
        }
    }
    printf("Result: %s\n", correct ? "PASS" : "FAIL");

    // 8. 释放内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

编译与运行命令如下：

# 编译
nvcc -arch=sm_80 -o vector_add vector_add.cu
# 运行
./vector_add
# 输出
Result: PASS

在上述代码中，blockIdx.x * blockDim.x + threadIdx.x 是 CUDA 中最经典的线程全局索引计算公式。当 N 不是线程块大小的整数倍时，我们使用 if (idx < N) 进行边界检查，防止越界访问。这种“网格-块-线程”的映射方式，使得我们能够灵活地处理任意长度的数据。

调试技巧：在开发初期，建议使用 cuda-memcheck 工具检查内存访问错误，并使用 cudaError_t err = cudaGetLastError(); 捕获 Kernel 启动错误。此外，cudaDeviceSynchronize() 可以阻塞主机直到所有 Kernel 执行完毕，便于计时与调试。

2.4 常见误区与性能调优复盘

在初学 CUDA 时，开发者常陷入以下几个误区，这些误区同样会影响 TensorRT 的部署效果：

误区一：线程数越多越好。实际上，当线程数超过硬件可驻留的上限时，多余的线程会排队等待，反而增加调度开销。最佳线程块大小通常为 128-512，且应为 32 的倍数（Warp 大小）。
误区二：忽略内存合并访问（Memory Coalescing）。GPU 访问全局内存时，以 32 字节或 128 字节为粒度。如果相邻线程访问相邻地址（如 C[idx]），硬件会将多次访问合并为一次事务，效率极高；反之，若线程访问跳跃地址，则会产生多次事务，严重降低带宽利用率。
误区三：过度使用共享内存。共享内存虽然快，但容量有限。若每个 Block 使用超过 48KB 共享内存，会限制 SM 上驻留的 Block 数量，降低 Occupancy。在 TensorRT 中，算子融合时需谨慎评估共享内存的收益与代价。
误区四：忽略主机与设备间的数据传输开销。PCIe 带宽（约 32 GB/s）远低于设备内存带宽（约 2 TB/s）。在推理场景中，应尽量减少 Host-Device 拷贝次数，或使用 CUDA 流（Stream）与异步拷贝来重叠数据传输与计算。

以一个实际调优案例为例：某目标检测模型在 TensorRT 部署时，发现预处理（图像缩放与归一化）耗时占整体推理的 30%。通过分析，瓶颈在于预处理在 CPU 上执行，且每次推理都需将处理后的数据拷贝到 GPU。优化方案是：将预处理写成 CUDA Kernel，在 GPU 上直接读取原始图像并完成缩放与归一化，同时使用 CUDA 流将数据拷贝与 Kernel 执行重叠。优化后，预处理耗时降低至原来的 1/10，整体推理延迟下降 25%。这个案例说明，理解 GPU 架构与 CUDA 编程模型，不仅是编写 Kernel 的基础，更是定位和解决部署性能瓶颈的关键能力。

风险提示：在编写自定义 CUDA Kernel 时，务必注意浮点运算的精度问题。GPU 的 FP32 运算与 CPU 的 FP32 运算在舍入方式上可能存在细微差异，导致结果不完全一致。在验证结果时，应使用相对误差阈值（如 1e-5）而非相等。此外，对于涉及原子操作（Atomic Operation）的 Kernel，需警惕竞争条件，必要时使用 atomicAdd 等内置原子函数。

总结而言，GPU 计算的核心在于理解硬件架构的并行本质，并通过 CUDA 的逻辑抽象将其转化为可执行的 Kernel。线程层次决定了任务的划分粒度，内存层次决定了数据的访问效率，而两者的协同优化则是性能调优的永恒主题。掌握这些基础，你便具备了深入 TensorRT 内部机制、编写高效推理插件的扎实功底。