4、连续访问：线程ID与内存地址的映射关系
好，咱们接着聊。前面几章我反复强调「合并访问」这个概念，说白了就是让相邻的线程去访问相邻的内存地址。那问题来了——怎么保证相邻的线程恰好访问相邻的地址？
答案就在线程ID和内存地址的映射关系上。这个映射关系，我个人认为是全局内存优化的核心中的核心。你把这个搞明白了，性能优化就成功了一半。
4.1 线程ID的布局：你以为是线性的？
很多初学者会想当然地认为，线程ID就是0、1、2、3...这样排下去的。嗯，这个想法没错，但不够全面。
在CUDA里，线程是以block为单位组织的。每个block内部，线程又按三维网格排列：threadIdx.x、threadIdx.y、threadIdx.z。我刚开始学的时候也觉得三维有点多余，直到做图像处理才明白——二维索引天然对应像素坐标，太方便了。
但要注意，硬件执行时，这些三维线程会被线性化。线性化的顺序是：x维度变化最快，y维度次之，z维度最慢。也就是说，同一个warp里的32个线程，它们的threadIdx.x是连续的，threadIdx.y和threadIdx.z相同。
关键点：warp内的32个线程，threadIdx.x一定是连续的。这是合并访问的基础。
4.2 映射模式：一维、二维、三维
咱们来看最常见的几种映射模式。我建议你先把一维的搞透，再去看多维的。
4.2.1 一维映射：最直接的方式
假设你有一个长度为N的一维数组，每个线程处理一个元素。最简单的做法是：
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[tid];这个模式，线程ID和数组索引完全对应。同一个warp里的32个线程，tid是连续的，访问的data[tid]也是连续的。完美合并访问。
我在项目中遇到过一种情况：有人把代码写成这样——
int tid = threadIdx.x; // 忘了加blockIdx
float val = data[tid];结果所有block都在访问data[0]到data[31]，性能惨不忍睹。嗯，这种低级错误我当年也犯过。
4.2.2 二维映射：处理矩阵的常用手法
处理二维矩阵时，我们通常用二维线程块。假设矩阵大小是M行N列：
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * N + col; // 行主序存储
float val = data[idx];这里有个坑，你想想看——threadIdx.x变化最快，所以col变化最快。如果矩阵是行主序存储（C语言默认方式），那么col连续变化时，idx也是连续的。合并访问没问题。
但如果是列主序存储（比如Fortran风格），那就要小心了。我曾经接手过一个项目，矩阵是列主序的，但代码按行主序去访问，结果带宽利用率不到10%。
避坑指南：我曾经因为没搞清楚矩阵存储格式，浪费了整整两天调试性能。记住：让变化最快的线程维度对应变化最快的内存维度。
4.2.3 三维映射：立体数据处理
三维映射其实原理一样，只是多了一层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;
int idx = z * (gridDim.x * blockDim.x) * (gridDim.y * blockDim.y)
+ y * (gridDim.x * blockDim.x) + x;这里x变化最快，所以对应的内存地址变化也最快。只要保证x方向上的线程访问连续地址，合并访问就能成立。
4.3 常见错误：映射搞反了会怎样？
咱们来看一个反面教材。假设有一个二维矩阵，你想让每个线程处理一行：
// 错误示例
int row = threadIdx.x; // 用x表示行
int col = blockIdx.x; // 用blockIdx表示列
int idx = row * N + col;
float val = data[idx];同一个warp里，threadIdx.x从0到31，所以row从0到31。那么idx就是0*N+col, 1*N+col, 2*N+col... 你想想看，这些地址之间隔了N个元素！
N如果是1024，那相邻线程访问的地址就差了1024个float，也就是4096字节。这完全不符合合并访问的条件。性能？我测过，大概只有正确映射的1/10。
小技巧：判断映射是否正确，可以算一下warp内相邻线程的地址差。如果差值是1个元素（4字节），那就是合并访问。如果差值很大，赶紧改代码。
4.4 实战建议：如何设计映射关系
我个人习惯按照以下步骤来设计映射：
- 先确定数据布局：行主序还是列主序？一维还是多维？
- 再确定线程布局：用几维block？每个维度多大？
- 让threadIdx.x对应变化最快的维度：这是铁律，不要违反。
- 验证地址连续性：写个小kernel，打印warp内线程的地址，看看是不是连续的。
我记得有一次优化一个三维卷积的kernel，数据是(x, y, z)格式，z变化最快。我一开始用threadIdx.x对应x，结果性能上不去。后来改成threadIdx.x对应z，带宽利用率直接从30%飙到85%。
说白了，映射关系就是「对号入座」。你把线程ID和内存地址的对应关系搞对了，硬件就能高效地帮你搬运数据。搞反了？那硬件就只能一个一个地搬，效率自然上不去。
下一章咱们聊聊stride访问——就是那种不连续的、跳着走的访问模式。嗯，那个坑更多，但也更有意思。