线程层次与索引：gridDim、blockDim、threadIdx、blockIdx的使用与计算
说实话，刚接触CUDA那会儿，我花了不少时间才搞明白这些索引到底怎么算。你想想看，一个核函数里同时跑着成千上万个线程，每个线程怎么知道自己该处理哪块数据？靠的就是这套索引体系。
今天咱们就把gridDim、blockDim、threadIdx、blockIdx这四个家伙彻底讲透。我保证，看完这章你就能手写正确的索引计算。
一、线程层次结构：从大到小看
CUDA的线程组织方式，说白了就是三层结构：
- Grid（网格）：最顶层，包含多个Block
- Block（线程块）：中间层，包含多个Thread
- Thread（线程）：最底层，实际执行计算的最小单位
我在项目中遇到过不少新手，一上来就写threadIdx.x，结果数据全乱套了。为什么？因为他没搞清楚这个索引是相对于Block的，不是相对于整个Grid的。
核心概念：
gridDim：Grid的维度，即有多少个BlockblockDim：Block的维度，即每个Block有多少个ThreadblockIdx：当前Block在Grid中的索引threadIdx：当前Thread在Block中的索引
二、索引计算：一维到三维
我个人习惯把索引计算分成三种情况来记。这样不容易乱。
1. 一维情况（最常用）
处理一维数组时，比如向量加法，我们只需要x维度。
// 一维索引计算
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 完整示例：向量加法
__global__ void vecAdd(float *a, float *b, float *c, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
c[tid] = a[tid] + b[tid];
}
}
// 调用方式
int n = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(a, b, c, n);我的经验：这里有个小细节——边界检查if (tid < n)一定要加。因为blocksPerGrid是向上取整的，最后一个Block可能有多余的线程。不加的话，访问越界会出段错误。
2. 二维情况（图像处理常用）
处理图像矩阵时，用二维索引更直观。我记得第一次写图像模糊核时，就是用的二维索引，代码清晰多了。
// 二维索引计算
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
// 完整示例：矩阵转置
__global__ void transpose(float *in, float *out, int width, int height) {
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (col < width && row < height) {
int in_idx = row * width + col;
int out_idx = col * height + row;
out[out_idx] = in[in_idx];
}
}
// 调用方式
dim3 threadsPerBlock(16, 16);
dim3 blocksPerGrid((width + 15) / 16, (height + 15) / 16);
transpose<<<blocksPerGrid, threadsPerBlock>>>(in, out, width, height);3. 三维情况（体数据处理）
三维用得少一些，但做体数据渲染或3D卷积时会用到。原理一样，就是多了一层z维度。
// 三维索引计算
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
// 完整示例：三维数组初始化
__global__ void init3D(float *data, int dimX, int dimY, int dimZ) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < dimX && y < dimY && z < dimZ) {
int idx = z * dimX * dimY + y * dimX + x;
data[idx] = 1.0f;
}
}三、常见陷阱与避坑指南
做性能优化这么多年，我见过太多人在索引上栽跟头。下面这几个坑，你一定要注意。
陷阱1：混淆blockIdx和threadIdx
我曾经有个同事，写核函数时把blockIdx.x当成了全局线程ID来用。结果数据全乱了，排查了整整一天才发现问题。
记住：blockIdx是Block的索引，threadIdx是Thread在Block内的索引。两者必须组合使用才能得到全局索引。
陷阱2：忘记边界检查
这个坑我踩过不止一次。当数据量不是Block大小的整数倍时，最后一个Block会有多余的线程。不加边界检查，轻则数据错误，重则程序崩溃。
陷阱3：维度搞反
二维情况下，blockIdx.x对应列方向，blockIdx.y对应行方向。我见过有人把x和y搞反，结果矩阵转置变成了矩阵复制。
四、性能优化建议
索引计算虽然简单，但写法不同，性能也会有差异。这里分享几个我常用的优化技巧。
| 优化点 | 建议做法 | 原因 |
|---|---|---|
| Block大小 | 128-512个线程 | 太小则占用率低，太大则寄存器压力大 |
| 维度选择 | 优先使用一维 | 一维索引计算开销最小 |
| 内存访问 | 保证连续线程访问连续地址 | 满足合并访问条件，带宽利用率高 |
| 索引复用 | 将计算结果存为局部变量 | 避免重复计算，减少寄存器使用 |
我的习惯：写核函数时，我通常先把索引计算好，存成局部变量。比如：
int tid = blockIdx.x * blockDim.x + threadIdx.x;后面所有用到这个索引的地方，直接用tid。这样既清晰，又避免了重复计算。
五、实战：从索引到数据映射
光说不练假把式。咱们来看一个实际例子，把索引计算和数据映射串起来。
// 场景：对1000x1000的矩阵做逐元素平方
// 每个线程处理一个元素
__global__ void squareMatrix(float *data, int width, int height) {
// 1. 计算二维索引
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
// 2. 边界检查
if (col < width && row < height) {
// 3. 计算一维内存地址（行优先存储）
int idx = row * width + col;
// 4. 执行计算
data[idx] = data[idx] * data[idx];
}
}
// 调用参数
dim3 blockSize(32, 32); // 每个Block 1024个线程
dim3 gridSize(
(1000 + 32 - 1) / 32, // 32个Block在x方向
(1000 + 32 - 1) / 32 // 32个Block在y方向
);
squareMatrix<<<gridSize, blockSize>>>(d_data, 1000, 1000);你看，整个流程就是：线程索引 → 数据坐标 → 内存地址 → 计算。每一步都清晰明了。
嗯，到这里线程层次和索引计算就讲完了。说白了就是四个变量、一个公式、三个维度。你只要记住全局索引 = blockIdx * blockDim + threadIdx这个核心公式，再注意边界检查，基本就不会出错了。
想要更多资料，联系可可