第二章：线程层次结构——Grid、Block、Thread 与 Warp 调度

好，咱们今天来聊聊 CUDA 里最核心的概念——线程层次结构。说实话，我刚接触 CUDA 那会儿，最头疼的就是搞不清 Grid、Block、Thread 到底是个什么关系。后来在项目里调优调得多了，才慢慢摸透了它们的脾气。

2.1 从硬件到软件：线程的“三级组织”

你可以把 GPU 想象成一个超大型的工厂。工厂里有多个车间（Grid），每个车间里有若干条生产线（Block），每条生产线上有多个工人（Thread）。工人是真正干活的人，生产线负责协调工人，车间则管理着整条生产线。

在 CUDA 里，Thread 是最小的执行单元。每个线程执行相同的核函数，但处理不同的数据。这就是所谓的 SIMT（单指令多线程）模型。

Block 是一组线程的集合。同一个 Block 里的线程可以共享一块内存（Shared Memory），也能通过同步操作来协作。我个人习惯把 Block 大小设为 256 或 512，这两个值在大多数 GPU 上表现都不错。

Grid 是所有 Block 的集合。一个核函数启动时，就对应一个 Grid。Grid 的维度决定了有多少个 Block 被调度。

一句话总结：Grid 包含 Block，Block 包含 Thread。你写核函数时，就是在定义这些线程如何协作。

2.2 一维、二维、三维线程索引

为什么要有维度？说白了，是为了方便处理不同结构的数据。

  • 一维索引：处理数组、向量这类线性数据。用 threadIdx.x 和 blockIdx.x 就能定位。
  • 二维索引：处理图像、矩阵。用 threadIdx.x、threadIdx.y 和对应的 block 索引。
  • 三维索引：处理体数据、3D 卷积。用得少，但需要时很顺手。

我记得有一次做图像滤波，数据是 1920x1080 的二维图像。我一开始傻乎乎地用一维索引，结果代码写得又臭又长。后来改成二维 Block，每个线程处理一个像素，代码瞬间清爽了。

// 一维索引示例
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;

// 三维索引示例
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int z = threadIdx.z + blockIdx.z * blockDim.z;
int idx = z * height * width + y * width + x;

我的小技巧：能用二维就别用一维，能用三维就别用二维。维度越高，索引计算越直观，代码可读性越好。

2.3 线程束（Warp）与调度

嗯，这里要注意。线程束是硬件层面的概念。GPU 执行指令时，不是一个个线程单独跑的，而是 32 个线程绑在一起，组成一个 Warp。这 32 个线程执行同一条指令，只是处理的数据不同。

为什么会这样？因为 GPU 的设计哲学就是“用更多的线程来隐藏延迟”。32 个线程一起跑，如果某个线程在等数据，其他线程可以继续干活。

我曾经在项目里遇到过一个坑：Block 大小设成了 100。你想想看，100 不是 32 的整数倍，最后一个 Warp 只有 4 个线程在干活，剩下 28 个线程空转。这浪费了多少计算资源！

避坑指南：Block 大小最好是 32 的整数倍。我一般用 128、256、512。别问我为什么不用 1024，因为有些老 GPU 不支持。

Warp 的调度是硬件自动完成的。每个 SM（流多处理器）里有多个 Warp Scheduler，它们负责把 Warp 分配给执行单元。你不需要手动干预，但需要理解它的行为——比如分支发散。

什么叫分支发散？就是同一个 Warp 里的线程走了不同的 if-else 分支。这时候 GPU 会串行执行所有分支，导致性能下降。我见过最夸张的例子，一个简单的 if 判断让性能掉了 50%。

// 不好的写法：分支发散
if (tid % 2 == 0) {
    // 偶数线程走这里
} else {
    // 奇数线程走这里
}

// 好的写法：用三元运算符或重新组织数据
result[tid] = (tid % 2 == 0) ? a[tid] : b[tid];

2.4 Occupancy 计算——你的 GPU 吃饱了吗？

Occupancy，说白了就是 GPU 的“上座率”。它表示每个 SM 上活跃的 Warp 数量与最大 Warp 数量的比值。Occupancy 越高，说明 GPU 越忙，理论上性能越好。

但 Occupancy 不是越高越好。我记得有一次调优，把 Occupancy 从 100% 降到 67%，性能反而提升了 30%。为什么？因为每个线程用了更多的寄存器，减少了寄存器溢出。

Occupancy 的计算公式其实不复杂：

Occupancy = (活跃 Warp 数) / (最大 Warp 数) * 100%

影响因素：
- 每个 Block 的线程数
- 每个线程使用的寄存器数
- 每个 Block 使用的共享内存大小
- GPU 的硬件限制（如 Max Threads per SM）

你可以用 NVIDIA 的 CUDA Occupancy Calculator 工具来算。或者直接在代码里调用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 函数。

GPU 型号 Max Warps per SM Max Threads per SM Max Blocks per SM
RTX 3090 64 1024 32
A100 64 2048 32
V100 64 2048 32

我的经验：Occupancy 在 50% 以上通常就够用了。别盲目追求 100%，关键是看实际吞吐量。用 ncu 或 nsys 跑一下，看看瓶颈在哪。

2.5 知识体系图

下面这张图是我自己画的，帮你理清本章的核心逻辑：

线程层次结构知识体系 Grid（网格） Block（线程块） Thread（线程） • 一维索引：threadIdx.x + blockIdx.x * blockDim.x • 二维索引：x, y 组合定位 • 三维索引：x, y, z 立体定位 Warp（32线程） Occupancy Grid → Block → Thread 三级结构，Warp 是硬件调度单位，Occupancy 衡量利用率

这张图把本章的核心概念串起来了。你从上往下看：Grid 包含多个 Block，Block 包含多个 Thread。Thread 是执行单元，32 个 Thread 组成一个 Warp，Warp 是硬件调度的基本单位。Occupancy 则反映了 SM 上活跃 Warp 的比例。

好了，这一章的内容就到这儿。线程层次结构是 CUDA 编程的基石，搞懂了它，后面的优化技巧才能用得上。下一章咱们聊聊内存层次结构——Global Memory、Shared Memory、Register 这些，到时候你会看到，内存访问模式对性能的影响有多大。

想要更多资料，联系可可

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