3、GPU计算模型：CUDA编程模型基础、线程层次结构（Grid/Block/Warp）、内存模型（Global/Shared/Local）

聊到GPU计算，绕不开的就是CUDA。说实话，我刚开始接触CUDA时，觉得它就是个“C语言加几个关键字”。后来踩了不少坑才明白——CUDA编程模型的核心，其实是帮你把“并行”这件事，用人类能理解的方式组织起来。

你想想看，GPU里有几千个核心，怎么让它们协同工作？总不能一个核心写一段代码吧。CUDA的做法很聪明：它定义了一套线程层次结构，让你像管理军队一样管理这些计算单元。

3.1 线程层次结构：Grid、Block、Warp

先看这张图，我把它叫做“GPU线程的军事化管理体系”：

Grid（网格） Block (0,0) T(0,0) T(1,0) T(2,0) T(0,1) T(1,1) T(2,1) T(0,2) T(1,2) T(2,2) ... Warp（32线程一组） Shared Memory Block (1,0) （线程结构同上） Block (2,0) （线程结构同上） 一个Grid包含多个Block 一个Block包含多个Thread 32个线程组成一个Warp

这张图我画了好一会儿，就是想让你一眼看明白：Grid 是最大的容器，里面装着多个 Block；每个 Block 里又装着多个 Thread。而 Warp 是硬件真正执行时的最小单位——32个线程绑在一起，执行同一条指令。

核心概念速记：

  • Grid：一个 kernel 启动时创建的所有线程的集合。你可以把它想象成一个“任务大本营”。
  • Block：Grid 里的一个子任务组。同一个 Block 内的线程可以协作（通过 Shared Memory 和同步）。
  • Thread：最小的执行单元。每个线程执行相同的 kernel 函数，但处理不同的数据。
  • Warp：硬件调度单位。32个线程为一组，在 SM 上以 SIMT 模式执行。

3.2 线程索引：我怎么知道我是谁？

每个线程都需要知道自己处理哪部分数据。CUDA 提供了内置变量：

// 一个典型的 CUDA kernel
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    // 计算全局线程ID
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (tid < N) {
        C[tid] = A[tid] + B[tid];
    }
}

// 调用方式
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这里有个细节我当年栽过跟头：blockIdx、threadIdx 都是从0开始的。你算全局ID时，公式是 blockIdx.x * blockDim.x + threadIdx.x。别搞反了，我见过有人写成 threadIdx.x * blockDim.x + blockIdx.x，结果数据全乱套。

我的小技巧：调试时先用 printf 打印出每个线程的 blockIdx 和 threadIdx，确认索引映射正确。这招帮我省了不少排查时间。

3.3 内存模型：数据住哪儿？

GPU 的内存不像 CPU 那样“大一统”。它分了好几层，每层的速度、容量、作用都不一样。我习惯用“城市交通”来类比：

内存类型 位置 访问速度 容量 作用域
Global Memory 显存（DRAM） 慢（~400 GB/s） 大（GB级） 所有线程
Shared Memory SM 内部（SRAM） 快（~1.5 TB/s） 小（几十KB） 同一个 Block
Local Memory 显存（DRAM） 慢（同Global） 小（寄存器溢出时用） 单个线程
Registers SM 内部 极快 极小（每个线程有限） 单个线程

说白了，Global Memory 是“大马路”，什么车都能走，但堵车；Shared Memory 是“小区内部路”，只有本小区的人能用，但跑得快。

3.4 Shared Memory 实战：矩阵转置

我拿矩阵转置举个例子。直接操作 Global Memory 的话，访问模式不连续，性能很差。用 Shared Memory 做“中转站”就好很多：

__global__ void transpose(float *in, float *out, int width) {
    __shared__ float tile[TILE_DIM][TILE_DIM];
    
    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;
    
    // 协作加载到 Shared Memory
    if (x < width && y < width) {
        tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    }
    
    __syncthreads();  // 确保所有线程都加载完毕
    
    // 转置后写回
    x = blockIdx.y * TILE_DIM + threadIdx.x;
    y = blockIdx.x * TILE_DIM + threadIdx.y;
    
    if (x < width && y < width) {
        out[y * width + x] = tile[threadIdx.x][threadIdx.y];
    }
}

我曾经踩过的坑：忘记加 __syncthreads()！没有这个同步屏障，有些线程可能还没写完 Shared Memory，别的线程就开始读了。结果就是数据错乱，调试到怀疑人生。记住：同一个 Block 内的线程需要同步时，一定要用 __syncthreads()。

3.5 Warp 发散：为什么你的性能上不去？

Warp 是32个线程一起执行同一条指令。但如果遇到 if-else 分支，情况就麻烦了：

// 糟糕的写法：Warp 内分支发散
if (threadIdx.x % 2 == 0) {
    // 偶数线程走这里
    result = a + b;
} else {
    // 奇数线程走这里
    result = a * b;
}

// 好的写法：尽量让同一个 Warp 走同一个分支
int warpId = threadIdx.x / 32;
if (warpId % 2 == 0) {
    // 整个 Warp 都走这里
    result = a + b;
} else {
    // 整个 Warp 都走这里
    result = a * b;
}

为什么会这样？因为硬件是“串行化”执行分支的——先执行 if 分支（偶数线程干活，奇数线程闲着），再执行 else 分支（奇数线程干活，偶数线程闲着）。一个 Warp 里分支越多，性能越差。我见过有人写了个 switch-case 有8个分支，性能直接掉了60%。

避坑指南：

  • 尽量让同一个 Warp 内的线程走相同的执行路径
  • 分支条件最好基于 warpId 或 blockIdx，而不是 threadIdx 的低位
  • 如果无法避免分支，考虑用“掩码计算”代替条件判断

3.6 小结：记住这三句话

嗯，这一章内容不少，但核心就三句话：

  1. 线程层次：Grid 管 Block，Block 管 Thread，硬件按 Warp 调度。
  2. 内存层次：Global 大而慢，Shared 小而快，Local 是寄存器的“备胎”。
  3. 性能关键：减少 Warp 发散，多用 Shared Memory，别忘了同步。

我个人习惯在写每个 kernel 前，先在纸上画一下线程索引和内存访问模式。这步看似简单，但能避免80%的后期调试痛苦。你试试看？

公众号:蓝海资料掘金营，微信deep3321

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