在GPU编程中，我们常常会遇到一个令人困惑的现象：明明GPU拥有成百上千个计算核心，理论算力远超CPU，但实际运行时却常常跑不满，甚至比CPU还慢。
答案几乎总是：内存访问效率低下。
GPU的计算能力发展速度远远超过了内存带宽的提升速度，"内存墙"问题在GPU上比在CPU上更加突出。而内存合并访问（Memory Coalescing）是解决GPU内存性能问题的第一把钥匙，也是所有CUDA程序员必须掌握的核心优化技术。
本文将从最底层的DRAM工作原理讲起，一步步带你理解内存合并的本质，以及如何在实际编程中应用它来大幅提升程序性能。
一、DRAM的本质：为什么它天生就慢？
要理解内存合并，我们首先必须理解DRAM本身的工作方式。很多人以为内存是一个可以随机访问、速度均匀的存储设备，但实际上，DRAM的设计充满了妥协，它的性能特性非常特殊。
1.1 电荷共享：DRAM读取的基本原理
DRAM的每个存储单元由一个晶体管和一个极小的电容组成，电容中存储的电荷量代表了数据：有电荷表示"1"，没有电荷表示"0"。
读取DRAM单元的过程基于电荷共享（Charge Sharing）原理：
- 晶体管门（字线，Word Line）打开，释放存储单元中极其微小的电荷量
- 这一点点电荷必须把很长的位线（Bit Line）所具有的大电容的电压抬升到足够高的水平
- 灵敏放大器（Sense Amplifier）检测到这个电压变化，判断出存储的是0还是1
- 读取完成后，灵敏放大器将数据回写到存储单元（因为读取是破坏性操作，电容中的电荷已被耗散）
这个过程可以用一个非常形象的类比来理解：
一个人在长长的走廊一端拿着一小杯咖啡，另一个人在走廊另一端，靠飘过去的香味来判断咖啡的口味。
- 小杯子 = DRAM 单元里的微量电荷
- 长走廊 = 长长的位线
- 香味 = 微弱的电信号
信号要传很远、还很弱，当然慢、难识别。这就是 DRAM 访问延迟高的根本原因。
1.2 容量与速度的永恒矛盾
理论上，我们可以通过在每个存储单元中使用更大、更强的电容来加快读取过程。电容越大，电荷越多，信号越强，读得就越快。
然而，DRAM 的发展方向恰恰相反。为了在每个芯片上存储更多比特数据，单元中的电容尺寸一直在不断缩小，电荷存储能力也随之变弱。
这就是为什么这么多年过去，DRAM 的容量翻了无数倍，但访问延迟依然很慢的原因。
1.3 DRAM的唯一优势：突发传输
DRAM 虽然随机访问很慢，但它有一个非常重要的特性：连续访问速度很快。
每次访问 DRAM 中的一个地址时，会一并访问包含该地址在内的一段连续地址。DRAM 芯片中有许多并行工作的灵敏放大器，每个灵敏放大器负责读取这段连续地址中某一个比特的内容。
一旦被灵敏放大器检测完成，所有这些连续地址的数据就可以高速传输给处理器。这些被一起访问并传输的连续地址单元，被称为 DRAM 突发（Burst）。
- 如果应用程序能够集中使用这些突发传输来的数据，DRAM 可以提供极高的带宽
- 如果应用程序进行随机访问，每次都要重新等待 DRAM 启动，速度会变得极慢
DRAM 的设计哲学就是：牺牲随机访问性能，换取极高的连续访问带宽。
二、CUDA内存合并：让DRAM发挥最大性能
CUDA 编程模型的核心是大量线程并行执行。而内存合并访问，就是让这些线程的访存行为能够完美匹配 DRAM 的突发传输特性。
2.1 什么是内存合并访问？
内存合并访问利用了一个基本事实：一个 warp 中的所有线程（通常为32个），在任意时刻都执行同一条指令。
当一个 warp 中的所有线程执行加载指令时，硬件会检测它们是否访问连续的全局内存地址。如果是，硬件会将所有这些访问合并（Coalesce）为一次对连续 DRAM 地址的集中访问。
以下是一个简单示例，展示合并与非合并访问的区别：
// ✅ 合并访问：连续线程访问连续地址
__global__ void coalescedRead(float* input, float* output, int N) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N) {
// 线程0访问input[0]，线程1访问input[1]，线程2访问input[2]...
// 32个线程访问32个连续地址 → 一次合并的内存事务
output[tid] = input[tid] * 2.0f;
}
}
// ❌ 非合并访问：连续线程访问跳跃地址
__global__ void stridedRead(float* input, float* output, int N, int stride) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid * stride < N) {
// 线程0访问input[0]，线程1访问input[stride]，线程2访问input[2*stride]...
// 地址分散 → 触发多次内存事务，性能大幅下降
output[tid] = input[tid * stride] * 2.0f;
}
}
对于合并访问，一个 warp 的32个线程只需要一次内存事务（假设数据对齐）；而对于跨步为32的非合并访问，最坏情况下需要32次独立的内存事务，带宽利用率接近 1/32。
2.2 行优先存储的影响
要学会有效使用合并访问硬件，我们必须先搞清楚多维数组在内存中是如何排列的。
C 语言和 CUDA 中的多维数组，都是按照行优先（Row-Major）规则存放在线性地址空间中的。 行优先指的是：先存完一整行，再存下一行。
比如二维数组 A[2][3]：
逻辑布局：
A[0][0] A[0][1] A[0][2]
A[1][0] A[1][1] A[1][2]
线性内存布局（行优先）：
地址: 0 1 2 3 4 5
数据: A[0][0] A[0][1] A[0][2] A[1][0] A[1][1] A[1][2]
|<-------- 第0行 ------->| |<-------- 第1行 ------->|
这里有一个非常重要的结论：
- 同一行中所有相邻的元素，在地址空间中都是连续存放的
- 同一列的元素，在线性内存中相隔了一整行的距离（Width 个元素）
这对 CUDA 合并访问有着决定性的影响：
| 访问模式 | 内存地址是否连续 | 是否合并 | 性能 |
|---|---|---|---|
| warp 线程按行访问（相邻线程访问同行相邻元素） | ✅ 连续 | ✅ 合并 | 快 |
| warp 线程按列访问（相邻线程访问同列相邻元素） | ❌ 跨步 Width | ❌ 不合并 | 慢 |
2.3 矩阵乘法中的访存分析
矩阵乘法 C = M × N 是 GPU 编程中最常见的操作之一，也是内存合并访问的经典案例。
在标准的朴素实现中，每个线程负责计算输出矩阵 C 的一个元素：
__global__ void matMulNaive(float* M, float* N, float* C, int Width) {
// 线程的行列坐标
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < Width && col < Width) {
float sum = 0.0f;
for (int k = 0; k < Width; k++) {
// 访问 M[row][k]：行主序，row固定，k递增 → 同一warp内row相同，k相同
// 访问 N[k][col]：行主序，k固定，col递增 → 同一warp内col连续 ✅ 合并
sum += M[row * Width + k] * N[k * Width + col];
}
C[row * Width + col] = sum;
}
}
对同一个 warp（threadIdx.y 相同，threadIdx.x 连续）中的线程分析访存：
-
矩阵 N 的访问：
N[k * Width + col]，其中k、Width对所有线程相同，col由blockIdx.x * blockDim.x + threadIdx.x决定，连续线程的col连续 → 地址连续，访问合并 ✅ -
矩阵 M 的访问：
M[row * Width + k]，其中row、Width、k对同一 warp 的所有线程（threadIdx.y相同）完全相同 → 所有线程访问同一地址（广播），同样高效 ✅ -
矩阵 C 的写入：
C[row * Width + col]，同上，col连续 → 写入合并 ✅
⚠️ 注意：若将线程组织方式改为
row = blockIdx.x * blockDim.x + threadIdx.x（即 x 维度对应行），则对 C 和 N 的访问将变为列方向跨步访问，导致非合并。正确的线程组织至关重要。
三、Corner Turning：解决列优先矩阵的访存难题
在实际应用中，我们经常会遇到矩阵存储布局不那么理想的情况。最常见的就是需要访问矩阵的转置。
3.1 问题：列优先矩阵的非合并访问
在线性代数中，我们经常需要同时使用一个矩阵的原始形式和转置形式。为了避免同时创建并存储两份数据，一个常见的做法是：只创建一种形式，当需要转置形式时，只需交换行下标和列下标来访问原始矩阵即可。
在 C 语言中，这相当于把转置后的矩阵看作是原始矩阵的列优先（Column-Major）布局。列优先指的是：先存完一整列，再存下一列。
列优先布局示意（Fortran/MATLAB 风格）：
逻辑布局: 线性内存地址:
A[0][0] A[0][1] 地址 0: A[0][0]
A[1][0] A[1][1] 地址 1: A[1][0] ← 同一列的元素在内存中相邻
A[2][0] A[2][1] 地址 2: A[2][0]
地址 3: A[0][1]
地址 4: A[1][1]
地址 5: A[2][1]
对于列优先存储的矩阵，如果 warp 中的连续线程沿行方向访问（即访问同一行的连续列元素），会发生什么？
// 访问列优先矩阵的第 row 行（等价于访问行优先矩阵的第 row 列）
// element[row][col] 在列优先内存中的地址 = col * Height + row
int col = blockIdx.x * blockDim.x + threadIdx.x; // 连续线程的 col 连续
float val = colMajorMatrix[col * Height + row];
// 连续线程访问地址: row, row+Height, row+2*Height, ...
// 地址步长 = Height（通常为数百到数千）→ 严重非合并 ❌
连续线程访问地址间距为 Height（整行长度），导致严重的非合并访问，性能会下降一个数量级。
3.2 解决方案：共享内存中转
面对无法合并的访存模式，我们有三种主要的优化策略：
- 调整线程与数据的映射方式
- 重新排列数据自身的存储布局（数据预处理）
- 以合并访问的方式在全局内存和共享内存之间传输数据，然后在共享内存中执行不友好的访存模式
第三种策略是最常用、最有效的，它利用了共享内存的一个关键特性：共享内存采用 SRAM 技术，不需要合并访问。
不管数据是列优先布局还是行优先布局，一旦把输入分块加载到共享内存之后，每个线程访问它需要的数据时，几乎不会有性能损失。
3.3 Corner Turning（转角优化）的实现
这种优化方法被称为 Corner Turning（转角优化），专门用于解决列优先矩阵的访存问题。
它的核心思想非常简单：
- 让连续的线程去加载输入分块同一列的连续元素（而不是同一行的连续元素）
- 由于矩阵是列优先存储的，同一列中的连续元素在内存中是相邻的
- 因此，连续的线程加载内存中相邻的元素，保证了内存访问可以合并
- 将数据存入共享内存（此时数据在共享内存中呈"转置"布局）
- 在共享内存中，线程可以按行访问数据，没有性能损失
下面是一个完整的 Corner Turning 示例，将列优先矩阵（等价于行优先矩阵的转置）与行优先矩阵相乘 C = A^T × B，其中 A 以行优先存储（访问时以列优先方式读取以模拟转置）：
#define TILE_SIZE 32
// C = A^T × B，其中 A 存储为行优先（按列访问模拟转置）
__global__ void matMulCornerTurning(
float* A, float* B, float* C,
int M, int K, int N // C(M×N) = A^T(M×K) × B(K×N)，A原始形状为K×M
) {
__shared__ float tileA[TILE_SIZE][TILE_SIZE]; // 存放 A^T 的分块
__shared__ float tileB[TILE_SIZE][TILE_SIZE]; // 存放 B 的分块
int row = blockIdx.y * TILE_SIZE + threadIdx.y; // C 的行索引
int col = blockIdx.x * TILE_SIZE + threadIdx.x; // C 的列索引
float sum = 0.0f;
for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
// === 加载 A^T 的分块（Corner Turning 的关键步骤）===
// A 原始存储为行优先，形状 K×M
// A^T[row][k] = A[k][row]，在行优先内存中地址为 k * M + row
//
// 普通加载方式（非合并）：
// tileA[threadIdx.y][threadIdx.x] = A[(t*TILE_SIZE + threadIdx.x) * M + row]
// → threadIdx.x 连续，但 * M 导致地址跨步，非合并 ❌
//
// Corner Turning（合并加载）：
// 交换 threadIdx.x 和 threadIdx.y 的角色来加载
int a_col = blockIdx.y * TILE_SIZE + threadIdx.x; // 对应A原始矩阵的列（A^T的行）
int a_row = t * TILE_SIZE + threadIdx.y; // 对应A原始矩阵的行
if (a_row < K && a_col < M)
// A[a_row][a_col]，行优先，a_row固定，a_col=blockIdx.y*TILE+threadIdx.x连续
// → 连续线程访问连续地址 ✅ 合并
tileA[threadIdx.x][threadIdx.y] = A[a_row * M + a_col];
else
tileA[threadIdx.x][threadIdx.y] = 0.0f;
// === 加载 B 的分块（标准合并加载）===
int b_row = t * TILE_SIZE + threadIdx.y;
int b_col = blockIdx.x * TILE_SIZE + threadIdx.x;
if (b_row < K && b_col < N)
// B[b_row][b_col]，行优先，b_row固定，b_col连续 → 合并 ✅
tileB[threadIdx.y][threadIdx.x] = B[b_row * N + b_col];
else
tileB[threadIdx.y][threadIdx.x] = 0.0f;
__syncthreads();
// === 在共享内存中计算（无需担心访存模式）===
for (int k = 0; k < TILE_SIZE; k++) {
// tileA[k][threadIdx.y] 访问 tileA 的列 → 共享内存，无惩罚
sum += tileA[k][threadIdx.y] * tileB[k][threadIdx.x];
}
__syncthreads();
}
if (row < M && col < N)
C[row * N + col] = sum;
}
Corner Turning 的核心在于加载 tileA 时将 threadIdx.x（warp 内连续变化）对应到全局内存中连续变化的维度（矩阵的列），从而保证合并访问。代价是数据在共享内存中以转置形式存储，但这完全没有问题，因为 SRAM 支持任意访问模式。
四、用生活类比理解内存合并
内存合并的概念虽然听起来有些抽象，但我们可以用生活中非常熟悉的例子来理解它。
4.1 堵车类比
内存合并的主要优势，是通过将多次内存访问合并为一次，来减少全局内存的访问流量。这就像缓解高速公路的交通拥堵。
- 全局内存带宽 = 高速公路
- 每次内存访问 = 一辆车
- 不合并访问 = 一堆车挤一条路 → 大堵车
- 合并访问 = 几辆车拼成一辆大巴 → 路不堵、跑得快
当拥堵发生时，每辆车的通行时间都会大幅增加。同样，当内存访问拥堵时，每个线程的访存延迟也会大幅增加，程序整体性能暴跌。
4.2 拼车类比
减少堵车的关键是少开车、多拼车。同理，减少内存拥堵的关键是少做零散访问、多做合并访问。
内存合并 = 线程拼车
拼车需要想要一起搭车的人互相迁就，并且达成一个共同的上下班时间。如果两个人的作息时间非常相似，他们可以很轻松地同乘一辆车上下班。但如果一个人白天睡觉晚上上班，另一个人晚上睡觉白天上班，他们就根本无法协调出共同时间一起开车上下班。
这对应到 CUDA 中就是：
- 作息相似的人 = 能合并的访问（连续线程访问连续地址）
- 作息不同的人 = 不能合并的访问（连续线程访问跨步地址）
想要实现内存合并访问，线程必须像作息一致的拼车伙伴一样：同时、连续、整齐地访问内存，硬件才能把它们"拼成一车"一起传输。
五、性能对比：合并 vs 非合并
通过 NVIDIA Nsight Compute 等工具，可以量化合并访问对性能的影响。下面的示例展示如何用 CUDA 事件计时对比两种访问模式：
#include <cuda_runtime.h>
#include <stdio.h>
#define N (1 << 24) // 16M 个 float，共 64MB
// 合并访问：stride = 1
__global__ void accessCoalesced(float* data, float* result, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) result[tid] = data[tid] + 1.0f;
}
// 非合并访问：stride = 32（warp 大小）
__global__ void accessStrided(float* data, float* result, int n, int stride) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int idx = (tid % stride) * (n / stride) + (tid / stride);
if (idx < n) result[tid % n] = data[idx] + 1.0f;
}
int main() {
float *d_data, *d_result;
cudaMalloc(&d_data, N * sizeof(float));
cudaMalloc(&d_result, N * sizeof(float));
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
int threads = 256;
int blocks = (N + threads - 1) / threads;
// 测试合并访问
cudaEventRecord(start);
for (int i = 0; i < 100; i++)
accessCoalesced<<<blocks, threads>>>(d_data, d_result, N);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms_coalesced;
cudaEventElapsedTime(&ms_coalesced, start, stop);
// 测试非合并访问（步长32）
cudaEventRecord(start);
for (int i = 0; i < 100; i++)
accessStrided<<<blocks, threads>>>(d_data, d_result, N, 32);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms_strided;
cudaEventElapsedTime(&ms_strided, start, stop);
printf("合并访问时间: %.2f ms\n", ms_coalesced);
printf("非合并访问时间: %.2f ms\n", ms_strided);
printf("性能比（非合并/合并）: %.1fx 慢\n", ms_strided / ms_coalesced);
cudaFree(d_data);
cudaFree(d_result);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return 0;
}
在典型的现代 GPU 上，步长为32的非合并访问比合并访问慢约 10~30 倍，具体数值取决于 GPU 架构和内存系统。
六、总结：内存优化的核心原则
通过本文的讲解，我们可以总结出 GPU 内存优化的几个核心原则：
- 理解底层硬件特性：DRAM 天生随机访问慢、连续访问快，所有优化都必须围绕这个特性展开
- 优先保证全局内存合并访问：这是对性能影响最大的优化，没有之一
- 利用共享内存作为中转：将全局内存的非合并访问转换为合并访问，再在共享内存中进行任意访问
- 避免随机访问全局内存：随机访问会完全浪费 DRAM 的带宽优势
-
注意线程组织方向：
threadIdx.x在 warp 内连续变化，应始终对应全局内存中连续的维度
内存合并访问是 CUDA 优化的入门第一课，但也是最重要的一课。掌握了它，你就掌握了 GPU 性能优化的钥匙。在实际编程中，永远先问自己一个问题：
我的 warp 线程访问的地址是连续的吗？
如果答案是肯定的，那么你的程序已经成功了一半。
附录A：关键概念速查表
| 概念 | 核心含义 |
|---|---|
| 电荷共享 | DRAM 读取的基本原理，小电容驱动大位线，导致信号微弱，延迟较高 |
| DRAM 突发 | DRAM 批量读取连续地址的机制，是 DRAM 高带宽的唯一来源 |
| Warp | CUDA 中同时执行同一指令的32个线程，是内存合并的基本单位 |
| 内存合并 | 将 warp 内多个连续地址的访问合并为一次 DRAM 突发传输 |
| 行优先 | C/CUDA 默认的多维数组存储方式，先存完一行再存下一行 |
| 列优先 | 先存完一列再存下一列，按行访问会导致非合并访问 |
| Corner Turning | 用共享内存中转解决列优先矩阵访存问题的优化技术 |
| 共享内存 | 采用 SRAM 技术，无需合并访问，速度极快但容量小（通常 48KB~128KB/SM） |
附录B：常见访存模式性能速查
| 访存模式 | 是否合并 | 典型相对性能 |
|---|---|---|
| 连续访问（stride=1） | ✅ 是 | 1x（基准） |
| stride=2 | 部分 | ~2x 慢 |
| stride=32（warp 大小） | ❌ 否 | ~10–30x 慢 |
| 完全随机访问 | ❌ 否 | ~10–30x 慢 |
| 广播（所有线程访问同一地址） | N/A（硬件优化） | 接近 1x |
