📚 课程目录第2章 GPU基础与CUDA编程模型

2. GPU基础与CUDA编程模型

核心原理：GPU为何适合并行计算

GPU（图形处理器）最初为图形渲染设计，其本质是大规模并行吞吐架构。与CPU追求低延迟、复杂控制逻辑不同，GPU将更多晶体管用于算术逻辑单元（ALU）和数据通路，以牺牲单核性能换取极高的并行吞吐量。现代GPU（如NVIDIA A100、H100）拥有数万个CUDA核心，可同时执行数万个线程，非常适合矩阵乘法、卷积等数据并行任务——这正是深度学习推理的主体运算。

理解GPU架构的关键在于流式多处理器（SM，Streaming Multiprocessor）。每个SM包含若干CUDA核心、共享内存、寄存器文件和调度器。以A100为例，其拥有108个SM，每个SM包含64个FP32 CUDA核心。线程实际执行时，以线程束（Warp）为单位——每32个线程组成一个warp，由调度器统一分发指令。这意味着，即便你启动了上万个线程，硬件也是按warp粒度（32线程）来调度执行的。

关键认知：GPU的并行度来自“大量线程 + 硬件自动调度”。程序员只需描述并行任务和线程组织方式，硬件负责将线程映射到SM上执行。但负载均衡问题恰恰源于此：如果不同warp的计算量差异巨大，先完成的SM会空闲等待，导致GPU利用率下降。

CUDA核心概念：线程层次与内存模型

CUDA（Compute Unified Device Architecture）是NVIDIA提供的并行计算平台。其编程模型将线程组织为三层结构：

Grid（网格）：一次内核启动（kernel launch）创建的所有线程集合，由多个Block组成。
Block（线程块）：一组线程，可包含1~1024个线程（取决于GPU架构）。同一Block内的线程可通过共享内存通信，并可进行同步（__syncthreads()）。
Thread（线程）：最小执行单元，每个线程有唯一的全局ID和局部ID。

内核启动时，程序员通过执行配置（execution configuration）指定Grid和Block的维度。例如：

// 启动一个包含 1024 个 Block、每个 Block 256 线程的内核
kernel_name<<<1024, 256>>>(args...);

// 线程全局ID计算方式
int global_tid = blockIdx.x * blockDim.x + threadIdx.x;

这种层次化设计使得GPU可以透明地扩展：小GPU上Block自动分配到少量SM上，大GPU上Block分配到更多SM上，无需修改代码。但这也意味着，Block是调度和负载均衡的基本单位——如果不同Block的计算量差异大，就会导致SM空闲。

内存层次：从寄存器到全局内存

CUDA内存模型是性能优化的核心。下表总结了各内存类型的特点：

内存类型	位置	作用域	延迟	容量	典型用途
寄存器	SM内部	单个线程	~1周期	每线程最多255个	局部变量、中间计算
共享内存	SM内部	Block内所有线程	~20周期	每SM 192KB（A100）	Block内数据复用、协作
L2缓存	芯片上	所有SM共享	~200周期	40MB（A100）	全局内存的缓存
全局内存	显存（HBM）	所有线程	~400-800周期	40GB/80GB	主数据存储
常量内存	显存（只读缓存）	所有线程	~400周期（缓存后快）	64KB	只读参数、权重

常见误区：初学者常将“共享内存”与“全局内存”混淆。共享内存在SM内部，延迟极低但容量有限；全局内存在显存中，容量大但延迟高。多卡推理中，显存带宽往往是瓶颈——例如A100的HBM带宽约2TB/s，但PCIe 4.0 x16带宽仅32GB/s，跨卡通信比卡内访问慢60倍以上。

方法流程：CUDA程序的基本开发步骤

分配主机内存：使用malloc或std::vector在CPU侧分配输入/输出数据。
分配设备内存：调用cudaMalloc在GPU上分配显存。
数据拷贝：使用cudaMemcpy将数据从主机传输到设备（H2D）。
编写内核函数：用__global__修饰符定义在GPU上执行的函数。
配置执行参数：确定Grid和Block维度，启动内核。
同步与拷贝回：调用cudaDeviceSynchronize等待内核完成，再将结果拷回主机（D2H）。
释放资源：调用cudaFree释放设备内存。

// 向量加法示例
__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];
}

int main() {
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(float));
    cudaMalloc(&d_b, n * sizeof(float));
    cudaMalloc(&d_c, n * sizeof(float));
    cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, n * sizeof(float), cudaMemcpyHostToDevice);
    
    vecAdd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);
    cudaDeviceSynchronize();
    
    cudaMemcpy(h_c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

案例与场景：多卡推理中的负载不均根源

在多卡推理系统中，负载不均的根源往往与CUDA线程组织方式密切相关。考虑一个典型的张量并行场景：将一个大矩阵乘法（如 Y = X @ W）切分到4张GPU上。每张卡负责计算输出的一部分列。如果W的列数不能被4整除，或者各卡负责的X分片大小不同，就会导致各卡的计算量不均衡。

更隐蔽的问题是warp发散（warp divergence）：当同一warp内的32个线程执行不同的分支路径时，GPU会串行执行每个分支，导致效率下降。例如，在推理的padding或mask操作中，如果某些线程处理有效数据、某些线程处理填充数据，就会产生warp发散。

  实战案例：某LLM推理系统使用4×A100，采用张量并行。发现GPU利用率分别为92%、88%、61%、45%。排查发现：

  ① 模型词表大小（32000）无法被4整除，导致最后一层输出头的计算分配不均；

  ② 注意力机制的padding mask导致部分warp执行空操作；

  ③ 各卡负责的KV cache大小不同，导致显存访问量差异。

  解决方案：调整切分策略，将词表补齐到32032（可被4整除），并优化mask逻辑减少warp发散。

误区与复盘：CUDA编程中的常见陷阱

误区1：线程越多越好。实际上，过多的线程会导致上下文切换开销和寄存器溢出（spill），反而降低性能。应通过occupancy计算找到最优线程数。
误区2：忽略内存合并访问（memory coalescing）。当同一warp的线程访问连续内存地址时，硬件可合并为一次内存事务；若访问不连续地址，则产生多次事务，带宽利用率骤降。
误区3：过度使用共享内存。共享内存虽快，但容量有限。若每个Block使用超过SM可用共享内存，内核将无法启动或导致occupancy下降。
误区4：忽视cudaMemcpy的同步开销。每次H2D/D2H拷贝都是同步操作，会阻塞CPU。应使用异步拷贝（cudaMemcpyAsync）和流（Stream）来重叠数据传输与计算。

复盘要点：在多卡推理中，负载不均的排查应从三个层面入手：计算层（各卡的计算量是否均衡）、内存层（各卡的显存访问量是否均衡）、通信层（各卡的通信量是否均衡）。CUDA编程模型提供了cudaEvent和nvprof/Nsight Systems等工具来量化各阶段耗时，帮助定位瓶颈。