# CUDA 深入浅出（三）：什么是 Coalesced Load

前两篇围绕 Tiling 讲了一个问题：怎样减少 global memory 访问次数。

这篇看另一个维度：即使你必须访问 global memory，也要让访问地址排得整齐一点。

CUDA 里这个概念叫 coalesced memory access。放在 load 上，就是 coalesced load。

可以先用一句话理解：

一个 warp 里的相邻 thread，
最好访问 global memory 里的相邻地址。

这句话听起来像常识，但它直接影响 GPU 能不能把多次 load 合并成更少的 memory transaction。

# GPU 不是一个 thread 一个 thread 去读内存

CUDA 程序写起来像很多 thread 各自执行自己的代码。

硬件执行时，GPU 会把 thread 按 warp 组织起来。一个 warp 通常有 32 个 thread。

当这 32 个 thread 执行同一条 load 指令时，硬件会看它们要读哪些地址：

thread 0  -> address ?
thread 1  -> address ?
thread 2  -> address ?
...
thread 31 -> address ?

如果这些地址挨在一起，硬件可以把它们合并成少量 memory transaction。

如果这些地址散得到处都是，硬件就要发出更多 memory transaction。每个 transaction 可能只带回来一小部分真正有用的数据，剩下带宽被浪费掉。

coalesced load 关心的不是你读了多少个元素，而是一个 warp 的这些地址有没有排成连续的一段。

# 连续访问长什么样

假设一个 float 占 4 字节。

如果一个 warp 里的 32 个 thread 这样读：

int tid = blockIdx.x * blockDim.x + threadIdx.x;
float v = x[tid];

那么地址大概是：

thread 0  -> x[0]
thread 1  -> x[1]
thread 2  -> x[2]
...
thread 31 -> x[31]

换成字节地址：

base + 0
base + 4
base + 8
...
base + 124

32 个 float 正好是 128 字节。

硬件看到一段连续地址，就可以用很少的 transaction 把这段数据取回来。具体 transaction 大小会随架构和对齐情况变化，但方向很明确：连续、对齐、相邻 thread 读相邻元素，GPU 最喜欢这种模式。

# 跨步访问长什么样

再看另一个写法：

int tid = blockIdx.x * blockDim.x + threadIdx.x;
float v = x[tid * stride];

如果 stride=1024，一个 warp 的地址会变成：

thread 0  -> x[0]
thread 1  -> x[1024]
thread 2  -> x[2048]
...
thread 31 -> x[31744]

这些地址隔得很远。

硬件没法把它们合成一段连续读取。它只能发出更多 transaction。每个 thread 只要 4 字节，但一次 memory transaction 往往会搬回一整段 cache line 或 memory segment。

于是你会看到一个很别扭的现象：

程序只用了 32 个 float
硬件可能搬了很多段内存

你写的代码没有多读元素，硬件却付出了更高的访存代价。

# row-major 矩阵里的相邻地址

C/C++ 里的二维矩阵如果按 row-major 存储：

A[row][col] 的线性地址 = row * N + col

同一行里，相邻列是连续的：

A[row][0], A[row][1], A[row][2], A[row][3]

地址挨在一起。

同一列里，相邻行之间隔着一整行：

A[0][col], A[1][col], A[2][col], A[3][col]

地址间隔是 N。

这就是矩阵代码里最容易踩到的坑。沿着列访问，在数学上很自然；在 row-major 内存里，它是 stride 访问。

GPU 不关心你的公式长得漂不漂亮。GPU 只看一个 warp 执行 load 时，那 32 个地址有没有挨在一起。

# naive matmul 里的访问模式

还是看最朴素的矩阵乘法：

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

float sum = 0.0f;

for (int k = 0; k < N; ++k) {
    sum += A[row * N + k] * B[k * N + col];
}

C[row * N + col] = sum;

这里有一个常见安排：

threadIdx.x 对应 col
threadIdx.y 对应 row

也就是相邻 thread 负责相邻列。

在同一个 k 上，一个 warp 里的相邻 thread 读 B 时，地址通常像这样：

B[k][col + 0]
B[k][col + 1]
B[k][col + 2]
...

这是连续地址。B 的 load 很容易 coalesce。

同一批 thread 读 A 时，很多 thread 会读同一个值：

A[row][k]
A[row][k]
A[row][k]
...

这不是连续的一段地址，但它有广播性质。现代 GPU 对这种同地址读取能处理得不错，cache 也会帮忙。

真正糟糕的情况通常来自另一种线程映射：相邻 thread 负责相邻行，而不是相邻列。

比如你把 threadIdx.x 用在 row 上：

int row = blockIdx.x * blockDim.x + threadIdx.x;
int col = blockIdx.y * blockDim.y + threadIdx.y;

这时相邻 thread 写不同的行、同一列。它们读或写 C[row][col] 时，地址会变成：

C[row + 0][col]
C[row + 1][col]
C[row + 2][col]
...

在 row-major 存储里，这就是 stride N 的访问。

一个简单规则可以先记住：

row-major 矩阵里，
尽量让 threadIdx.x 对应连续的 col。

这样 B 的读取、C 的写回，以及很多 tile load 都更容易 coalesce。

# Tiled matmul 里的 coalesced load

上一篇的 tiled kernel 里有两句 load：

As[ty][tx] = A[row * N + tile * T + tx];
Bs[ty][tx] = B[(tile * T + ty) * N + col];

先看 A。

同一个 warp 或 half-warp 里，如果 tx 连续变化，row 相同，那么这些 thread 读的是：

A[row][tile*T + 0]
A[row][tile*T + 1]
A[row][tile*T + 2]
...

这是同一行的连续元素。

再看 B。

col 通常由 blockIdx.x * T + tx 得到。如果 tx 连续变化，ty 相同，那么这些 thread 读的是：

B[tile*T + ty][col + 0]
B[tile*T + ty][col + 1]
B[tile*T + ty][col + 2]
...

这也是同一行的连续元素。

所以一个写得正常的 tiled matmul，会让 block 在把 A_tile 和 B_tile 从 global memory 搬进 shared memory 时，尽量使用 coalesced load。

这一步很关键。Tiling 让你少访问 global memory，coalescing 让你访问 global memory 时少浪费 transaction。

# 数一下浪费在哪里

假设一个 warp 读取 32 个 float。

理想情况下，这 32 个元素连续排列：

32 × 4 bytes = 128 bytes

硬件可以把这段连续数据用少量 transaction 取回来。

如果这 32 个元素相隔很远，比如 stride 很大，硬件可能要为许多 thread 各自取一段内存。每段里只有 4 字节被用上，其余字节被丢在路上。

从代码角度看，你读了同样的 32 个 float。

从内存系统角度看，coalesced 版本让带宽花在有效数据上，stride 版本让带宽花在散乱地址上。

这也是很多 CUDA profile 里会看到的问题：算术指令不多，global load 数量看起来也没夸张，但 kernel 还是慢。内存访问形状不好，带宽没有被用到刀刃上。

# coalesced load 和 tiling 的关系

Tiling 和 coalesced load 解决两个不同层面的问题。

Tiling 关心：

同一个数据能不能少从 global memory 搬几次？

Coalesced load 关心：

这一次必须从 global memory 搬数据时，
warp 里的地址能不能合并成少量 transaction？

所以你可以把它们连起来看：

先用 tiling 减少 global memory load 的元素数量。
再用 coalescing 提高每次 global memory load 的效率。

前者减少“去仓库的次数”。后者让每次去仓库时，大家拿的是同一排货架上的东西。

# 写 kernel 时先检查这几件事

看一个 CUDA kernel 的 global memory 访问，可以先问几个问题：

1. 一个 warp 里的 threadIdx.x 是否访问连续地址？
2. row-major 矩阵里，threadIdx.x 是否对应 col？
3. load A_tile 和 B_tile 时，tx 是否落在连续维度上？
4. 写回 C 时，相邻 thread 是否写相邻 C[row][col]？

如果答案经常是否定的，kernel 很可能在浪费 global memory 带宽。

coalesced load 不会改变矩阵乘法公式，也不会减少逻辑上的元素读取数量。它改变的是硬件取数据的方式。

CUDA 优化里很多细节都绕不开这条线：让 warp 的访问地址变得连续，让 global memory transaction 带回来的字节尽量都有用。