第一章：CUDA编程模型概述

大家好，我是你们这期课程的主讲。在GPU计算这个领域摸爬滚打了十几年，我见过太多人一上来就写核函数，结果性能惨不忍睹。说白了，不理解硬件就写不出好代码。今天咱们就从最基础的开始聊——GPU架构、CUDA编程模型，还有那个最核心的概念：核函数。

1.1 GPU架构演进：从图形卡到通用计算利器

先说说GPU是怎么变成今天这个样子的。我刚开始做GPU计算那会儿，用的还是GTX 200系列，那时候CUDA刚出来没多久，文档都不全。你想想看，一个原本用来画三角形的芯片，突然要去做科学计算，这本身就是个疯狂的想法。

GPU架构的演进，说白了就是一条路：从专用走向通用。

  • 早期（2000年前后）：固定功能管线，只能干图形渲染的活。你想做点别的？没门。
  • 统一着色器架构（2006年）：NVIDIA的G80架构是个里程碑。顶点着色器和像素着色器合并了，变成了通用的计算单元。嗯，CUDA就是这时候诞生的。
  • Fermi架构（2010年）：我第一次用Fermi卡的时候，感觉像换了个时代。完整的缓存层次结构、ECC内存、真正的双精度性能。我记得当时做的一个流体模拟项目，从GTX 280迁移到GTX 480，性能直接翻了3倍。
  • Kepler/Maxwell/Pascal：每一代都在优化能效比和计算密度。我个人最喜欢Pascal，它把HBM2内存带到了消费级市场，带宽那个爽啊。
  • Volta/Turing/Ampere/Hopper：引入了Tensor Core、RT Core这些专用单元。说白了，GPU又开始"专用化"了——只不过这次是为AI和光线追踪服务。

核心要点：现代GPU本质上是大规模并行处理器阵列。它由几十个SM（Streaming Multiprocessor，流式多处理器）组成，每个SM里有几百个CUDA核心。你写的核函数，最终就是被这些CUDA核心并行执行的。

我给大家画个图，看看GPU的架构长什么样：

GPU架构示意图（简化版） 全局内存（Global Memory）— 显存，容量大但延迟高 L2缓存（所有SM共享） SM阵列（Streaming Multiprocessors） SM 0 CUDA Cores × 64 Shared Memory 48KB 寄存器 65536个 Warp调度器 SM 1 CUDA Cores × 64 Shared Memory 48KB 寄存器 65536个 Warp调度器 SM 2 CUDA Cores × 64 Shared Memory 48KB 寄存器 65536个 Warp调度器 SM ... ... ... ... ... 数据流向：全局内存 → L2 → SM内共享内存/寄存器

1.2 CUDA编程模型：主机与设备

CUDA编程模型，说白了就是CPU+GPU的协作模式。CPU叫"主机"（Host），GPU叫"设备"（Device）。主机负责控制逻辑、串行任务，设备负责大规模并行计算。

我刚开始教学生的时候，总有人问："为什么不能只用GPU？" 你想想看，GPU虽然计算快，但它的控制逻辑很弱，分支预测、乱序执行这些能力远不如CPU。所以合理的分工是：CPU做决策，GPU做苦力。

来看一个典型的CUDA程序流程：

  1. 分配设备内存：在GPU上开辟空间
  2. 数据拷贝：从主机内存拷贝到设备内存
  3. 调用核函数：让GPU干活
  4. 结果回传：从设备内存拷贝回主机内存
  5. 释放资源：清理战场

个人经验：我见过太多新手在数据拷贝这一步栽跟头。PCIe带宽再快，也赶不上显存带宽的零头。所以我的原则是：尽量减少主机和设备之间的数据传输。能一次传完就别分十次，能在GPU上算完就别回传。

1.3 核函数概念：GPU上的"轻量级线程"

核函数（Kernel）是什么？说白了，就是在GPU上运行的函数。你用__global__修饰符声明它，然后用尖括号语法<<<grid, block>>>调用它。

这里有个关键点：核函数不是串行执行的，而是被成千上万个线程同时执行。每个线程执行相同的代码，但处理不同的数据。这就是所谓的SIMT（Single Instruction, Multiple Threads）模型。

来看一个最简单的例子——向量加法：

// 核函数定义
__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 main() {
    // ... 数据准备和内存分配 ...
    
    // 调用核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // ... 结果回传和清理 ...
    return 0;
}

这段代码里，每个线程计算一个元素的和。threadIdx.x是线程在块内的索引，blockIdx.x是块在网格内的索引，blockDim.x是每个块的线程数。这三个变量加起来，就能唯一确定一个线程的全局ID。

避坑指南：我曾经在一个项目里，核函数跑出来结果全是错的。查了半天，发现是blocksPerGrid算少了，导致数组末尾的元素没被处理。记住：一定要做边界检查，就是那个if (i < N)。别偷懒，否则数据会教你做人。

1.4 线程层次结构：Grid、Block、Thread

CUDA的线程组织方式，我习惯用"军队"来类比：

  • Grid（网格）：整个任务，相当于一个集团军
  • Block（线程块）：一个作战单元，相当于一个师
  • Thread（线程）：单个士兵，执行具体计算

每个Block内的线程可以协作——它们共享一块Shared Memory，可以通过__syncthreads()同步。但不同Block之间的线程是独立的，不能直接通信。

为什么要有这个层次结构？说白了，是为了适配硬件。一个Block会被调度到同一个SM上执行，Block内的线程可以高效共享数据。而不同Block可以自由地分配到不同SM上，实现自动负载均衡。

层次 硬件对应 最大规模 通信方式
Thread CUDA Core 每个Block最多1024个线程 寄存器、局部内存
Block SM 每个Grid最多2^31-1个Block Shared Memory、__syncthreads()
Grid GPU设备 取决于显存大小 全局内存

核心要点：选择Block大小是个技术活。我个人习惯从256开始试，然后根据性能调优。太小了（比如32）会导致SM利用率不足，太大了（比如1024）又可能受限于寄存器数量。记住：Block大小最好是32的倍数，因为Warp（线程束）就是32个线程一组。

1.5 内存层次结构：速度与容量的博弈

GPU的内存系统，说白了就是用容量换速度。离计算单元越近，速度越快，但容量越小。

  • 全局内存（Global Memory）：容量最大（几个GB到几十GB），但延迟最高（几百个周期）。所有线程都能访问。
  • 共享内存（Shared Memory）：每个Block独享，容量小（几十KB），但延迟极低（几个周期）。我经常用它来做数据复用和线程间通信。
  • 寄存器（Registers）：每个线程私有，速度最快。但数量有限（每个SM几万个），用多了会 spill 到局部内存，性能暴跌。
  • 常量内存/纹理内存：有缓存加速的只读内存，适合存储查找表之类的数据。

嗯，这里要注意：全局内存的访问模式非常关键。如果线程访问连续地址（合并访问），带宽利用率能接近100%。如果随机访问，带宽可能掉到10%以下。我在优化矩阵乘法时，光是把数据布局从行主序改成列主序，性能就提升了5倍。

个人经验：调优GPU程序，80%的时间都在和内存打交道。计算本身很快，瓶颈往往在数据搬运上。我的建议是：先分析访存模式，再优化计算逻辑。用NVIDIA Nsight Compute工具看看L1/L2缓存命中率，命中率低于80%就要警惕了。

好了，第一章的内容就到这里。记住：理解硬件是写出高性能CUDA代码的前提。下一章我们会深入讨论线程同步和Warp调度，那才是真正考验功力的时候。

想要更多资料，联系可可

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