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

大家好，我是你们的老朋友。今天咱们来聊聊CUDA里最核心的概念——线程层次结构。说白了，就是线程(Thread)、线程块(Block)和网格(Grid)这三兄弟到底是怎么回事。

我记得刚接触CUDA那会儿，看了半天文档，脑子里还是一团浆糊。后来在项目里踩了几个坑，才真正搞明白。嗯，今天我就把这些经验分享给你。

一、为什么需要层次结构？

你想想看，GPU里有几千个核心。如果让程序员一个一个去管理这些核心，那不得累死？

所以NVIDIA想了个办法：把计算任务分层。就像公司组织架构一样——CEO管几个总监，总监管几个经理，经理管几个员工。这样管理起来就清晰多了。

在CUDA里，这个架构是这样的：

• 网格(Grid)：最顶层，代表整个任务
• 线程块(Block)：中间层，一组线程的集合
• 线程(Thread)：最底层，实际执行计算的单元

二、线程(Thread)——最小的执行单元

线程是CUDA程序里最基础的执行单位。每个线程都执行相同的核函数(Kernel)，但处理不同的数据。

怎么区分每个线程呢？靠的是内置变量：

• threadIdx.x：线程在块内的x方向索引
• threadIdx.y：线程在块内的y方向索引
• threadIdx.z：线程在块内的z方向索引

举个例子，一个线程块里有256个线程，那它们的threadIdx.x就是从0到255。

// 每个线程打印自己的ID
__global__ void printThreadId() {
    printf("Thread ID: %d\n", threadIdx.x);
}

我在项目中遇到过一个问题：新手容易把threadIdx和全局线程ID搞混。记住，threadIdx只是线程在块内的编号，不是在整个网格里的编号。

三、线程块(Block)——线程的集合

线程块是一组线程的集合。同一个块里的线程可以协作——它们能共享内存，也能同步执行。

线程块也有自己的索引：

• blockIdx.x：块在网格内的x方向索引
• blockIdx.y：块在网格内的y方向索引
• blockIdx.z：块在网格内的z方向索引

每个块的大小是有限制的。我记得老一点的GPU，一个块最多1024个线程。现在的安培架构也是这个限制。

线程块还有个重要属性——维度。你可以把块设成一维、二维或三维：

// 一维块：256个线程
dim3 blockDim(256);

// 二维块：16x16 = 256个线程
dim3 blockDim(16, 16);

// 三维块：8x8x4 = 256个线程
dim3 blockDim(8, 8, 4);

我个人习惯用二维块来处理图像。图像是二维的嘛，用二维块映射起来特别自然。

四、网格(Grid)——整个任务

网格是最高层次，它包含了所有线程块。网格也有维度，同样可以是一维、二维或三维。

网格的大小用gridDim来表示：

• gridDim.x：网格在x方向有多少个块
• gridDim.y：网格在y方向有多少个块
• gridDim.z：网格在z方向有多少个块

启动核函数时，你需要指定网格和块的大小：

// 定义网格和块的大小
dim3 gridDim(4, 4);    // 4x4 = 16个块
dim3 blockDim(16, 16); // 16x16 = 256个线程/块

// 启动核函数
myKernel<<<gridDim, blockDim>>>();

这段代码启动了16个块，每个块256个线程。总共就是16 × 256 = 4096个线程。

五、三者的关系——一张表说清楚

层次	内置变量	维度属性	典型大小
线程	threadIdx	blockDim	1~1024
线程块	blockIdx	gridDim	取决于GPU
网格	无	用户定义	取决于数据量

说白了，这三者的关系就是：网格包含块，块包含线程。每个线程通过blockIdx和threadIdx就能唯一确定自己在整个网格中的位置。

六、如何计算全局线程ID？

这是面试常考题，也是实际开发中必须掌握的技能。计算全局线程ID的公式：

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

// 二维情况
int tid_x = blockIdx.x * blockDim.x + threadIdx.x;
int tid_y = blockIdx.y * blockDim.y + threadIdx.y;
int tid = tid_y * (gridDim.x * blockDim.x) + tid_x;

// 三维情况类似，这里不展开了

我曾经在项目里犯过一个低级错误：忘了乘blockDim.x，结果所有线程都访问同一个数据。排查了半天才发现，真是血的教训。

七、实际应用中的选择策略

怎么选网格和块的大小？我一般这么考虑：

1. 先定块大小：通常选128或256。太小了浪费硬件，太大了受限制。
2. 再算网格大小：总线程数 = 块大小 × 网格大小，要大于等于数据量。
3. 考虑warp：块大小最好是32的倍数，这样不会产生浪费的线程。

举个例子，你要处理100万个元素：

int N = 1000000;
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize; // 向上取整

myKernel<<<gridSize, blockSize>>>(d_data, N);

这样算出来gridSize是3907。嗯，多出来的几个线程会在核函数里通过判断跳过。

八、避坑指南

最后，我总结几个常见的坑：

• 块内线程数别超1024——这是硬限制，超了直接报错
• 网格大小别太大——虽然理论上可以很大，但实际受GPU内存限制
• 别忘了处理边界——总线程数可能比数据量多，要用if判断
• 二维/三维索引别搞混——x、y、z的顺序要记清楚

我记得有一次，我把blockIdx.x和blockIdx.y搞反了。结果图像处理出来是转置的，找了半天才发现。从那以后，我每次写索引都会在纸上画一遍。

好了，线程层次结构就讲到这里。说白了就是三层嵌套：网格管块，块管线程。搞明白这个，后面的学习就顺了。下一章咱们聊聊warp——这个GPU调度的小秘密。