《CUDA与TensorRT入门知识框架》

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

核心原理:GPU为何而生

要理解GPU计算,首先需要回到GPU的诞生初衷。CPU(中央处理器)设计目标是低延迟——尽快完成单个复杂任务;而GPU(图形处理器)设计目标是高吞吐量——同时处理海量简单运算。这种本质差异决定了二者的硬件结构截然不同。

GPU的架构可以概括为“多核宽SIMD”模式。以NVIDIA Ampere架构为例,一块GA102芯片包含10752个CUDA核心,这些核心被组织为84个SM(流式多处理器),每个SM内部又分为4个处理块,每块包含32个CUDA核心。CUDA核心是执行浮点运算的最小单元,它们共享寄存器文件、共享内存和L1缓存。这种层级化设计让GPU能够以极低的调度开销并行执行成千上万个线程。

下图展示了GPU的典型硬件层级结构,以及线程在硬件上的映射关系:

GPU 芯片(以GA102为例) SM 0 处理块 0 32 CUDA核心 处理块 1 32 CUDA核心 处理块 2 32 CUDA核心 处理块 3 32 CUDA核心 SM 1 处理块 0 32 CUDA核心 处理块 1 32 CUDA核心 处理块 2 32 CUDA核心 处理块 3 32 CUDA核心 ······ L2 缓存 / 共享内存(40 MB) 所有SM通过L2缓存与全局内存交互 全局显存(GDDR6X,24 GB) 带宽高达 936 GB/s 线程映射关系 Grid(网格) 包含多个Block Block(线程块) 映射到SM,含多个Warp Warp(线程束) 32个线程同时执行 Thread(线程) 由CUDA核心执行 Block→SM

从图中可以看到,Grid → Block → Warp → Thread 构成了CUDA编程的线程层级模型。硬件调度器将Block分配到SM上执行,SM再以Warp(32线程)为最小调度单位。这种“线程块+线程束”的设计,使得GPU可以通过硬件线程切换来隐藏内存访问延迟——当一个Warp等待数据时,SM立即切换到另一个就绪的Warp,从而保持计算单元始终忙碌。

CPU与GPU对比:分工不同的“搭档”

CPU与GPU并非替代关系,而是各司其职的协作关系。下表从多个维度对比二者的核心差异:

对比维度 CPU(以Intel i9-13900K为例) GPU(以NVIDIA RTX 4090为例)
核心数量 24核心(8P+16E) 16384个CUDA核心
核心频率 5.8 GHz(单核睿频) 2.52 GHz(Boost)
内存带宽 89.6 GB/s(DDR5双通道) 1008 GB/s(GDDR6X)
缓存策略 大容量L3缓存(36MB),复杂分支预测 L2缓存(72MB),依赖线程切换隐藏延迟
设计目标 低延迟执行串行任务 高吞吐量执行并行任务
典型场景 操作系统、数据库、逻辑控制 矩阵运算、图像处理、深度学习

一个形象的比喻是:CPU像一位全能博士,能处理各种复杂问题,但一次只能专注几件事;GPU像一支千人计算兵团,每个人只会简单的加减乘除,但成千上万人同时动手,总量惊人。在实际系统中,CPU负责控制流和任务调度,GPU负责大规模并行计算,二者通过PCIe总线或NVLink互联,协同完成整体任务。

并行计算概念:从“排队”到“同时干”

并行计算的核心思想是将一个大问题拆解为多个可同时执行的小问题。在GPU语境下,主要涉及三个层次的并行:

  • 任务级并行:不同任务同时执行,例如同时处理多张图像的滤波。
  • 数据级并行:同一操作应用于大量数据元素,例如对100万个浮点数同时执行乘法。
  • 指令级并行:单条指令驱动多个数据单元运算,即SIMD(单指令多数据)模式。

GPU最擅长的是数据级并行。以向量加法为例,CPU的串行代码如下:

// CPU串行版本
for (int i = 0; i < N; i++) {
    C[i] = A[i] + B[i];
}

而CUDA的并行版本将循环拆分为N个线程,每个线程只计算一个元素:

// GPU并行版本(CUDA Kernel)
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

执行时,GPU创建N个线程,每个线程独立执行一次加法。由于线程数量远大于物理核心数,硬件会自动分批次调度,这就是“大规模线程 + 硬件调度”的并行模型。

案例:矩阵乘法的并行化改造

矩阵乘法是深度学习和科学计算中最基础的操作。假设我们要计算 C = A × B,其中A为M×K矩阵,B为K×N矩阵。CPU串行实现需要三重循环:

for (int i = 0; i < M; i++)
    for (int j = 0; j < N; j++) {
        float sum = 0;
        for (int k = 0; k < K; k++)
            sum += A[i*K+k] * B[k*N+j];
        C[i*N+j] = sum;
    }

这段代码的时间复杂度为O(M×N×K)。在GPU上,我们可以让每个线程负责计算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;
    }
}

假设M=N=K=1024,则总计算量为10亿次乘加。CPU单核需要约1秒,而RTX 4090的FP32算力约82.6 TFLOPS,理论上仅需约12毫秒——加速比接近100倍。实际应用中,通过共享内存优化、分块计算,还能进一步提升性能。

误区与复盘:并行计算的常见陷阱

误区1:线程越多越快。线程数量超过硬件可并发执行的极限后,增加线程只会增加调度开销。经验法则是:线程总数应为SM数量的整数倍,且每个Block的线程数建议为128或256的倍数(Warp大小的整数倍)。

误区2:忽略内存访问模式。GPU对连续内存访问(合并访问)效率极高,而对随机访问效率极低。例如,线程i访问A[i]是高效的,但访问A[random[i]]则会导致大量缓存未命中。

复盘要点:在编写CUDA程序时,应优先考虑数据布局和访问模式,其次才是线程组织。一个常见的性能分析流程是:先用ncu(NVIDIA Nsight Compute)分析内核的内存吞吐量和计算吞吐量,判断是内存瓶颈还是计算瓶颈,再针对性优化。

最后需要强调的是,GPU并行计算并非万能。对于分支密集、依赖性强、数据量小的任务,CPU反而更快。实际工程中,通常采用CPU+GPU异构计算:CPU执行串行控制逻辑,GPU执行大规模并行计算,二者通过异步传输重叠数据拷贝与计算,从而最大化整体吞吐量。