3. CUDA编程模型：主机与设备、Kernel函数概念、线程层次结构、内存层次结构

好，咱们今天聊聊CUDA编程模型。说实话，很多初学者一上来就被各种概念砸晕了——主机、设备、Kernel、线程块、网格……我当年刚接触时也差不多。但别急，咱们一层层剥开来看。

3.1 主机与设备：CPU和GPU的分工

先搞清楚一个基本概念：在CUDA的世界里，主机（Host）指的是CPU及其内存系统，设备（Device）指的是GPU及其显存。说白了，CPU负责控制逻辑和串行任务，GPU负责大规模并行计算。

我个人习惯把CPU想象成项目经理，GPU就是一群埋头苦干的工程师。项目经理分配任务、检查结果，工程师们并行执行具体工作。嗯，这个比喻虽然简单，但很贴切。

关键点：主机端代码在CPU上执行，设备端代码在GPU上执行。两者通过PCIe总线通信，数据需要显式地在主机内存和设备显存之间拷贝。

我在项目中遇到过一个问题：有个同事把数据声明在主机端，然后在Kernel里直接访问，结果程序崩溃了。为什么？因为GPU不能直接访问主机内存（除非使用统一内存，那是后话）。所以记住：数据在哪，代码就在哪执行。

3.2 Kernel函数概念：GPU上执行的函数

Kernel函数，就是运行在GPU上的函数。它的定义方式很特别——用__global__修饰符声明。调用时用尖括号语法 <<<grid, block>>> 指定执行配置。

来看个最简单的例子：

// 定义一个Kernel函数
__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() {
    // ... 数据准备和拷贝 ...
    vecAdd<<<numBlocks, blockSize>>>(d_A, d_B, d_C, n);
    // ... 结果拷贝回主机 ...
}

你想想看，这个Kernel函数有什么特点？它没有返回值（void），而且每个线程执行相同的代码，但处理不同的数据——这就是所谓的单指令多线程（SIMT）模型。

小技巧：我习惯在Kernel函数里加一个边界检查，就像上面代码中的 if (i < n)。因为线程总数可能超过实际数据量，不加检查会访问越界内存，导致难以调试的错误。

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

这是CUDA最核心的概念之一。线程的组织方式分三层：

  • Thread（线程）：最小的执行单元，每个线程执行Kernel函数的一份实例
  • Block（线程块）：一组线程的集合，块内线程可以共享内存和同步
  • Grid（网格）：一组线程块的集合，一个Kernel对应一个Grid

为什么会这样设计？说白了，是为了适配GPU的硬件架构。一个Block内的线程会被调度到同一个流多处理器（SM）上执行，这样它们才能高效地共享数据。而不同Block之间是独立的，可以分配到不同的SM上并行执行。

我记得有一次调试一个图像处理程序，图像尺寸是1024x768。我一开始把每个像素映射到一个线程，但没注意线程块大小。结果性能很差。后来我调整了Block大小（256线程/块），Grid大小自动计算，性能提升了3倍。嗯，这里要注意：Block大小通常是32的倍数，因为warp（线程束）大小就是32。

层级 硬件对应 最大尺寸 通信方式
Thread CUDA Core 无限制（受资源限制） 寄存器、局部内存
Block SM（流多处理器） 1024线程（现代GPU） 共享内存、同步屏障
Grid GPU设备 2^31-1（一维） 全局内存

线程索引的计算公式也很重要：

// 一维情况
int tid = threadIdx.x + blockIdx.x * blockDim.x;

// 二维情况（图像处理常用）
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int idx = y * width + x;

避坑指南：我曾经犯过一个错误——在Kernel里直接用 threadIdx.x 作为全局索引，忘了加上 blockIdx.x * blockDim.x。结果所有Block的线程都从0开始索引，数据被反复覆盖。调试了一下午才发现。所以，全局索引一定要用公式计算。

3.4 内存层次结构：从快到慢

GPU的内存层次和CPU类似，但更强调带宽和并行访问。咱们从快到慢捋一遍：

  • 寄存器（Register）：最快，每个线程私有。容量极小（通常256KB/SM），但零延迟。
  • 共享内存（Shared Memory）：块内线程共享，可编程控制的缓存。延迟比全局内存低一个数量级。
  • L1/L2缓存：硬件自动管理，对程序员透明。L1和共享内存共用同一块物理存储。
  • 全局内存（Global Memory）：最大（GB级别），但延迟最高（几百个周期）。所有线程都能访问。
  • 常量内存（Constant Memory）：只读，有缓存机制。适合存储查找表等不变数据。
  • 纹理内存（Texture Memory）：专门为图形处理优化，支持插值和边界处理。

你想想看，为什么要有这么多层？说白了，就是用容量换速度。寄存器最快但最小，全局内存最慢但最大。写CUDA程序的核心就是：把频繁访问的数据放到快的存储里。

我个人的经验是：先默认用全局内存，等性能瓶颈出现时，再考虑把热点数据搬到共享内存。别一开始就过度优化，容易把代码搞复杂。

实战建议：矩阵乘法是理解内存层次结构的经典案例。朴素实现直接访问全局内存，性能很差。优化版本会使用共享内存做分块（tiling），把子矩阵加载到共享内存中，计算时反复重用。性能差距可以达到10倍以上。

最后说一句：内存层次结构不是死的。不同架构（比如Turing、Ampere、Hopper）的缓存大小、共享内存容量都不一样。写代码时最好查一下目标GPU的参数，或者用 cudaGetDeviceProperties 在运行时查询。

好了，这一章的内容就这些。下一章咱们聊聊CUDA编程模型进阶：同步与原子操作、流与事件、错误处理，到时候见。

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