第二章 CUDA编程模型：Kernel函数、线程层次结构、内存层次结构

好，咱们直接进入正题。CUDA编程模型，说白了就是NVIDIA给咱们准备的一套「并行计算工具箱」。你想想看，GPU里有几千个核心，怎么让它们听话地干活？靠的就是这套模型。我个人习惯把CUDA编程模型拆成三块来理解：Kernel函数怎么写、线程怎么组织、内存怎么管。这三块搞明白了，你就能写出能跑的CUDA程序了。

2.1 Kernel函数：GPU的入口函数

Kernel函数，就是跑在GPU上的函数。它跟普通的C函数最大的区别在于——前面多了个 __global__ 修饰符。嗯，这里要注意：__global__ 函数必须返回 void，而且不能是类的成员函数（除非是静态的）。

核心要点：Kernel函数是异步执行的。你调用它之后，CPU不会等它跑完就继续往下执行了。这一点坑过不少人，包括我。

来看个最简单的例子：

// 定义一个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];
    }
}

// 调用Kernel函数
int main() {
    // ... 分配内存、拷贝数据 ...
    vecAdd<<<gridDim, blockDim>>>(d_A, d_B, d_C, N);
    // CPU继续干别的事...
}

看到那个 <<<gridDim, blockDim>>> 了吗？这就是CUDA的「三括号语法」。左边是网格维度，右边是块维度。我刚开始学的时候老记反，后来干脆这么记：网格包块，块包线程。

个人经验：我在项目中遇到过一个问题——Kernel函数里用了 printf 想调试，结果啥也没打印出来。后来才发现，Kernel是异步的，程序退出了输出缓冲区还没刷新。解决办法是在Kernel调用后加 cudaDeviceSynchronize()。

2.2 线程层次结构：网格、块、线程

CUDA的线程组织方式，其实是个三层结构：

  • 网格（Grid）：最外层，包含多个线程块
  • 线程块（Block）：中间层，包含多个线程
  • 线程（Thread）：最内层，实际执行计算的最小单位

为什么会这样设计？说白了是为了适配GPU的硬件架构。一个线程块里的线程可以共享内存、同步执行，但不同线程块之间是独立的。你想想看，如果所有线程都能随便共享数据，那光同步开销就够你受的。

网格 (Grid) 线程块 0 blockIdx.x = 0 T0 T1 T2 ... Tn 线程块 1 blockIdx.x = 1 T0 T1 T2 ... Tn 线程块 2 blockIdx.x = 2 T0 T1 T2 ... Tn 更多线程块... 更多线程块... 更多线程块... 线程层次结构：Grid → Block → Thread 每个线程通过 threadIdx 和 blockIdx 唯一标识

每个线程都有自己唯一的ID，通过 threadIdx 和 blockIdx 来标识。比如 threadIdx.x 表示线程在块内的x方向索引，blockIdx.x 表示块在网格内的x方向索引。再加上 blockDim.x（块的大小），就能算出全局唯一的线程ID：

int globalThreadId = threadIdx.x + blockIdx.x * blockDim.x;

这个公式我建议你背下来。写Kernel函数时，90%的场景都要用它来定位数据。

避坑指南：我曾经在项目里把 blockDim 和 gridDim 搞混了。结果就是——线程ID算错了，数据全写到别人的位置上去了。调试了整整一个下午才发现。记住：blockDim 是每个块的大小，gridDim 是网格里有多少个块。

2.3 内存层次结构：从快到慢，从少到多

GPU的内存，说白了就是个「金字塔」结构。越靠近计算单元的内存越快，但容量越小；越远的内存越慢，但容量越大。咱们从上往下看：

内存类型 位置 访问速度 生命周期 作用域
寄存器 SM内部 ~1周期 线程内 单个线程
共享内存 SM内部 ~5周期 块内 线程块
L1缓存 SM内部 ~10周期 自动管理 SM内
L2缓存 芯片级 ~50周期 自动管理 全局
全局内存 显存 ~400周期 程序运行期 全局
本地内存 显存 ~400周期 线程内 单个线程

你想想看，如果每个线程都去全局内存读数据，那400个周期的延迟谁受得了？所以优化的核心思路就是——让数据尽量待在快的内存里。

2.4 共享内存：线程块的「小食堂」

共享内存是我最喜欢用的优化手段。它位于SM内部，同一个线程块里的所有线程都能访问。说白了就是个「小食堂」——大家把菜（数据）从大仓库（全局内存）搬过来，在食堂里现炒现吃，省得来回跑。

声明共享内存用 __shared__ 修饰符：

__global__ void sharedMemExample(float* input, float* output, int N) {
    __shared__ float cache[256];  // 共享内存数组
    
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    
    // 把数据从全局内存搬到共享内存
    if (tid < N) {
        cache[threadIdx.x] = input[tid];
    }
    __syncthreads();  // 同步！确保所有线程都搬完了
    
    // 在共享内存上做计算
    // ... 比如做归约操作 ...
    
    // 把结果写回全局内存
    if (tid < N) {
        output[tid] = cache[threadIdx.x];
    }
}

个人经验：我在做矩阵乘法优化时，用共享内存把全局内存访问次数减少了90%以上。但要注意——共享内存只有几十KB，别塞太多东西。我曾经试图在一个块里放1024个float，结果编译通过了，运行时直接报错。后来查文档才知道，每个SM的共享内存上限是48KB（老架构）或96KB（新架构）。

2.5 同步与原子操作

多线程编程最头疼的就是同步问题。CUDA里提供了两种同步机制：

  • __syncthreads()：块内同步，确保所有线程都执行到同一位置
  • 原子操作：比如 atomicAdd()，保证对同一内存地址的读写不被中断

嗯，这里要特别提醒：__syncthreads() 只能在同一个线程块内使用。不同块之间没法同步——这是硬件决定的。如果你需要跨块同步，那就得用全局内存加原子操作，或者干脆换个算法。

避坑指南：我曾经在条件分支里用了 __syncthreads()，结果程序直接卡死了。为什么？因为有些线程走了if分支，有些走了else分支，它们永远等不到对方。记住：__syncthreads() 必须保证所有线程都能执行到，不能在分支里用。

2.6 内存访问模式：合并访问

全局内存的访问效率，很大程度上取决于你的访问模式。GPU喜欢「连续访问」——也就是相邻的线程访问相邻的地址。这叫「合并访问」（Coalesced Access）。

来看个对比：

// ✅ 合并访问（高效）
float val = data[threadIdx.x];  // 线程0读data[0]，线程1读data[1]...

// ❌ 非合并访问（低效）
float val = data[threadIdx.x * 64];  // 线程0读data[0]，线程1读data[64]...

为什么？因为GPU一次从全局内存读取128字节（32个float）。如果线程0到31连续访问data[0]到data[31]，一次内存事务就搞定了。但如果跳着访问，就得多次内存事务，效率直接打折扣。

我个人习惯在写Kernel函数之前，先在纸上画一下线程和数据的对应关系。确保访问模式是连续的，这能省下不少优化时间。

2.7 小结

CUDA编程模型的核心就这些：Kernel函数是入口，线程层次结构是组织方式，内存层次结构是数据通路。搞明白这三块，你就能写出正确的CUDA程序了。至于性能优化——那是后面章节的事。

记住一句话：写CUDA程序，本质上就是在管理「数据在哪里」和「谁在计算它」。把这两个问题想清楚，代码自然就写对了。

核心要点回顾：

  • Kernel函数用 __global__ 修饰，返回void，异步执行
  • 线程层次：Grid → Block → Thread，用 threadIdx 和 blockIdx 定位
  • 共享内存是块内的高速缓存，用 __shared__ 声明
  • 合并访问是全局内存优化的关键
  • 同步用 __syncthreads()，注意不能在分支里用

想要更多资料，联系可可

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