4. 内存带宽瓶颈：Global Memory访问模式、合并访问与对齐、Bank Conflict

说到GPU性能优化，内存带宽往往是第一个拦路虎。

我见过太多人，核函数写得飞快，一跑起来却发现性能惨不忍睹。为什么？说白了，数据没喂饱计算单元。GPU的计算能力再强，如果内存访问跟不上，那就是典型的「等米下锅」。

4.1 Global Memory的访问模式

Global Memory是GPU中最大的存储空间，但也是延迟最高的。它的带宽虽然大，但有个前提——你得按规矩来访问。

我个人习惯把Global Memory想象成一个巨大的仓库。你每次去取货，如果取的是相邻货架上的东西，搬运工一趟就能搞定。如果你东取一个西取一个，搬运工就得来回跑断腿。

在CUDA里，这个「搬运工」就是内存控制器。它一次能搬运的数据量是固定的——通常是32字节、64字节或128字节。你想想看，如果每个线程只取4字节，但内存控制器一次搬了128字节，那剩下的124字节不就浪费了？

核心原则：让相邻线程访问相邻地址。这是Global Memory优化的第一性原理。

4.2 合并访问（Coalesced Access）

合并访问，说白了就是让一个warp里的32个线程，访问连续的内存地址。

我在项目中遇到过这样一个案例：一个简单的向量加法，性能死活上不去。我一看代码，发现线程索引和内存地址的映射搞反了。调整之后，带宽利用率从30%直接飙到了90%。

来看个对比：

// ❌ 非合并访问 —— 性能杀手
__global__ void bad_access(float *A, float *B, int N) {
    int tid = threadIdx.x;
    for (int i = 0; i < N / blockDim.x; i++) {
        int idx = i * blockDim.x + tid;
        B[idx] = A[idx] * 2.0f;
    }
}

// ✅ 合并访问 —— 高效
__global__ void good_access(float *A, float *B, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        B[tid] = A[tid] * 2.0f;
    }
}

为什么第一个不行？因为每个线程处理的是跨步的数据。线程0处理索引0、32、64...线程1处理索引1、33、65...你想想看，一个warp里的线程访问的地址根本不连续，内存控制器得发多少次请求？

第二个就好多了。线程0访问索引0，线程1访问索引1，线程2访问索引2...32个线程访问连续的32个地址，一次内存事务就搞定了。

小技巧：检查你的内存访问是否合并，最简单的方法就是看warp内线程的地址是否连续。如果地址跨度是1，那就是完美合并。跨度越大，性能越差。

4.3 对齐访问（Aligned Access）

对齐访问和合并访问是两兄弟，但很多人容易搞混。

对齐，指的是内存地址必须是某个值的整数倍。对于Global Memory，我建议地址至少按128字节对齐。

为什么会这样？因为GPU的内存事务是以128字节为单位的。如果你的数据起始地址不是128的倍数，那一次事务可能跨越两个128字节的边界，导致额外的内存请求。

我曾经调试过一个矩阵转置的核函数，性能始终达不到理论峰值。折腾了半天，发现是分配内存时没对齐。加上cudaMalloc的自动对齐特性后，问题迎刃而解。

注意：cudaMalloc默认分配的内存是256字节对齐的，但如果你自己用cudaMallocPitch或者手动计算偏移，就要格外小心对齐问题。

对齐访问的检查方法：

// 检查地址是否128字节对齐
if ((uintptr_t)ptr % 128 == 0) {
    printf("地址已对齐\n");
} else {
    printf("地址未对齐，性能可能受损\n");
}

4.4 Bank Conflict（存储体冲突）

说完了Global Memory，咱们聊聊Shared Memory。Shared Memory虽然快，但它有个坑——Bank Conflict。

Shared Memory被分成32个Bank，每个Bank的宽度是4字节。相邻的32位数据依次存放在不同的Bank中。比如地址0-3在Bank0，地址4-7在Bank1，依此类推。

当一个warp里的多个线程同时访问同一个Bank的不同地址时，就会发生Bank Conflict。这时候，这些访问会被串行化，性能直接打折扣。

我举个例子：

// ❌ 有Bank Conflict
__shared__ float shared[32];
int tid = threadIdx.x;
float val = shared[tid * 2];  // 线程0和线程16都访问Bank0

// ✅ 无Bank Conflict
float val = shared[tid];  // 每个线程访问不同的Bank

在第一个例子里，线程0访问shared[0]（Bank0），线程16访问shared[32]（也是Bank0）。两个线程抢同一个Bank，冲突了。

第二个例子就没事，每个线程访问自己的Bank，完美并行。

避坑指南：我曾经在实现一个归约操作时，用了步长为2的访问模式，结果Bank Conflict导致性能下降了40%。改成步长为1后，速度立马就上来了。

4.5 知识体系总览

下面这张图，是我自己总结的内存带宽优化知识体系。你可以把它当成一个检查清单：

4.6 实战建议

说了这么多，我总结几条实战经验：

先看Global Memory：用nvprof或Nsight Compute检查内存带宽利用率。如果低于80%，大概率是访问模式有问题。
再看Shared Memory：如果用了Shared Memory，检查有没有Bank Conflict。Nsight Compute会直接告诉你冲突次数。
最后看对齐：检查数据地址是否对齐。不对齐的话，考虑用cudaMallocPitch或者手动填充。

我的习惯：每次写完核函数，我都会用Nsight Compute跑一遍，重点关注L1/L2缓存命中率和内存带宽利用率。这两个指标能直接反映你的内存访问是否高效。

嗯，内存带宽这块就讲到这里。记住一句话：让数据流动起来，别让计算单元闲着。

📚 课程目录

4. 内存带宽瓶颈：Global Memory访问模式、合并访问与对齐、Bank Conflict

4.1 Global Memory的访问模式

4.2 合并访问（Coalesced Access）

4.3 对齐访问（Aligned Access）

4.4 Bank Conflict（存储体冲突）

4.5 知识体系总览

4.6 实战建议