《CUDA与TensorRT推理全景通识》

GPU计算基础：GPU架构、CPU与GPU对比、并行计算概念

在进入CUDA与TensorRT的世界之前，必须先建立对GPU硬件底层的直观认知。很多初学者直接上手写CUDA代码，却因为不理解线程调度和内存层次而写出低效程序。本节将拆解GPU的物理架构、与CPU的本质差异，以及并行计算的核心思维。

一、核心原理：GPU的硬件架构解剖

GPU并非简单的“很多个CPU核心”，而是一种大规模并行吞吐机。以NVIDIA Ampere架构为例，其硬件层级从大到小依次为：

  • GPU芯片（Die）：包含多个图形处理簇（GPC），每个GPC内部有若干流式多处理器（SM）。
  • 流式多处理器（SM）：这是GPU最核心的计算单元。一个SM内部包含：
    • 多个CUDA核心（负责整数/浮点运算）；
    • Tensor Core（专门加速矩阵乘加，对深度学习至关重要）；
    • 共享内存（Shared Memory）：SM内所有线程可访问的高速缓存；
    • 寄存器文件（Register File）：分配给线程的极速存储；
    • 调度器（Warp Scheduler）：负责将线程束（Warp）分发到执行单元。
  • 线程束（Warp）：GPU执行指令的最小单位，固定包含32个线程。这32个线程必须执行相同的指令（SIMT模式），如果出现分支分歧，则部分线程会被屏蔽，造成性能损失。
关键认知：CPU是“低延迟”设计，擅长复杂逻辑分支；GPU是“高吞吐”设计，用大量线程掩盖访存延迟。GPU的ALU（算术逻辑单元）数量远超CPU，但单个ALU频率较低。

下图展示了GPU内部从芯片到线程束的层级关系，以及内存层次结构：

GPU 芯片（Ampere 架构示例） GPC（图形处理簇） SM 0 CUDA Cores ×64 Tensor Cores ×4 Shared Mem 128KB Registers 64K×32bit 4 Warp Schedulers SM 1 CUDA Cores ×64 Tensor Cores ×4 Shared Mem 128KB Registers 64K×32bit 4 Warp Schedulers GPC（图形处理簇） SM 2 CUDA Cores ×64 Tensor Cores ×4 Shared Mem 128KB Registers 64K×32bit 4 Warp Schedulers SM 3 CUDA Cores ×64 Tensor Cores ×4 Shared Mem 128KB Registers 64K×32bit 4 Warp Schedulers 内存层次结构 寄存器（每线程私有） ~几KB/线程 共享内存（SM内） ~128KB/SM L2 缓存（芯片级） ~40MB 全局显存（VRAM / HBM） 几十GB，带宽可达 2TB/s 线程束（Warp） 32个线程组成1个Warp 同一Warp执行相同指令 分支分歧 → 串行执行 调度单位：Warp SM中最多常驻 64 Warps CPU vs GPU 关键差异 CPU：4-16核，单核强，延迟低 GPU：数千核，单核弱，吞吐高 CPU：适合分支、串行逻辑 GPU：适合数据并行、矩阵运算 CPU：大缓存，低带宽 GPU：小缓存，高带宽

二、CPU与GPU对比：设计哲学的分野

CPU与GPU的差异本质上是延迟优化与吞吐优化的取舍。CPU的晶体管大部分用于分支预测、乱序执行、大容量缓存，以降低单线程延迟；GPU则将晶体管用于大量ALU和较小的缓存，通过线程级并行隐藏访存延迟。

对比维度 CPU GPU
核心数量 4~16个高性能核心 数千个（如A100有6912个CUDA核心）
核心频率 3~5 GHz 1.4~2.0 GHz
内存带宽 ~50 GB/s（DDR5） ~2 TB/s（HBM2e）
缓存大小 L3可达32MB L2约40MB，无L3
适合任务 操作系统、逻辑控制、串行计算 矩阵乘法、图像处理、深度学习
编程模型 多线程（OS调度） SIMT（硬件调度Warp）
常见误区：很多人认为“GPU核心多，所以所有程序都会更快”。实际上，如果程序是串行的、分支密集的，或者数据量小无法填满GPU，GPU反而可能比CPU更慢。GPU的优势只有在大规模数据并行场景下才能发挥。

三、并行计算概念：从线程到网格

CUDA的并行模型是层级化的，理解这个层级是编写高效内核的前提。从下到上依次为：

  1. 线程（Thread）：最小的执行单元，每个线程有独立的寄存器和程序计数器。
  2. 线程块（Block）：一组线程（通常128~512个），块内线程可通过共享内存通信，并可通过__syncthreads()同步。
  3. 网格（Grid）：一组线程块，网格内的块之间不能直接通信，只能通过全局内存间接交互。

在硬件层面，线程块会被分配到SM上执行。一个SM可以同时驻留多个线程块，但受限于寄存器数量和共享内存大小。例如，一个SM有64K寄存器，如果每个线程用32个寄存器，则最多驻留2048个线程。

下面是一个典型的CUDA内核启动配置示例：

// 定义内核函数
__global__ void vectorAdd(float* A, float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// 主机端启动内核
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这段代码中，blockIdx.x表示当前线程块在网格中的索引，threadIdx.x表示线程在块内的索引，blockDim.x表示块内线程数。通过公式 blockIdx.x * blockDim.x + threadIdx.x 即可计算出全局线程ID。

四、案例与场景：矩阵乘法的并行化

矩阵乘法是深度学习的核心运算（全连接层、卷积本质都是矩阵乘）。假设我们要计算 C = A × B，其中A为M×K，B为K×N。最直观的并行策略是：每个线程计算C中的一个元素。

朴素实现如下：

__global__ void matMulKernel(float* A, float* B, float* C, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++) {
            sum += A[row * K + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

这种实现虽然正确，但性能很差，因为每个线程都直接访问全局内存，而全局内存延迟高达数百个时钟周期。优化方向包括：

  • 使用共享内存分块（Tiling）：将A和B的块加载到共享内存，减少全局内存访问次数。
  • 使用Tensor Core：调用wmma或mma指令，利用硬件矩阵乘加速单元。
  • 循环展开：减少循环开销，提高指令级并行。
性能对比参考：在A100上，朴素矩阵乘（FP32）约 10 TFLOPS，使用共享内存优化后可达 19 TFLOPS，使用Tensor Core（TF32）可达 156 TFLOPS。差距超过15倍，可见内存访问优化和专用硬件利用的重要性。

五、误区与复盘：并行编程的常见陷阱

初学者在理解GPU计算时，常犯以下错误，这里逐一复盘：

误区 正确理解
“线程越多越好” 线程过多会导致寄存器溢出（spill）到本地内存，反而变慢。需要平衡占用率与资源使用。
“共享内存无限大” 共享内存只有几十KB，需要精心管理，否则会限制SM上驻留的线程块数量。
“所有线程同时执行” Warp是调度单位，不同Warp的执行进度不同。块内同步用__syncthreads()，但块间无法同步。
“忽略分支分歧” 如果Warp内线程走不同分支，会串行执行所有分支路径，性能急剧下降。应尽量让Warp内线程统一路径。
“全局内存访问随意” 全局内存访问必须合并（coalesced），即相邻线程访问相邻地址，才能达到最大带宽。

复盘一个典型场景：假设一个Warp的32个线程分别访问数组的data[threadIdx.x * 2]，即间隔访问。此时硬件需要两次事务才能完成读取，带宽利用率只有50%。正确的做法是让线程访问data[threadIdx.x]，实现完全合并访问。

另一个常见问题是过度同步。有些程序员习惯在每个循环后都加__syncthreads()，这会导致线程等待，降低并行效率。同步只应在真正需要共享数据时使用。

最后，理解GPU的延迟隐藏机制至关重要。当某个Warp因访存而等待时，SM会立即切换到另一个就绪的Warp执行，从而保持ALU忙碌。因此，足够的线程数量（通常建议每个SM至少驻留1024个线程）是隐藏延迟的前提。如果线程太少，访存等待就会暴露，GPU利用率下降。

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