📚 课程目录第3章全局内存带宽测试

3. 全局内存带宽测试

3.1 为什么带宽测试是性能调优的第一步？

做GPU优化的人，第一个要面对的瓶颈就是全局内存带宽。我个人习惯，拿到一个新卡或者新架构，第一件事就是跑个带宽测试。为什么？因为很多kernel跑不快，不是计算不够快，而是数据喂不进去。

你想想看，GPU的计算单元动辄几千个，但显存带宽就那么几百GB/s到几TB/s。一旦访存模式不对，实际带宽可能连理论值的10%都跑不到。我曾在项目中遇到过，一个看起来计算密集的模型，优化完访存后速度直接翻了8倍——嗯，说白了，问题根本不在计算上。

核心概念：全局内存带宽 = 单位时间内从显存读取或写入的数据量。理论带宽由显存频率和位宽决定，但实际带宽取决于你的访存模式是否「友好」。

3.2 编写第一个带宽测试程序

我们先写一个最简单的版本。这个程序只做一件事：从全局内存中读取一个数组，然后写回另一个数组。注意，这里我刻意不做任何计算，就是为了纯粹测带宽。

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void bandwidth_test_kernel(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx];  // 简单的读-写操作
    }
}

int main() {
    int N = 1 << 24;  // 16M个float，约64MB
    size_t bytes = N * sizeof(float);

    float *d_in, *d_out;
    cudaMalloc(&d_in, bytes);
    cudaMalloc(&d_out, bytes);

    // 初始化数据
    float *h_in = (float*)malloc(bytes);
    for (int i = 0; i < N; i++) h_in[i] = 1.0f;
    cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice);

    // 配置kernel
    int threads = 256;
    int blocks = (N + threads - 1) / threads;

    // 计时
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start);
    bandwidth_test_kernel<<<blocks, threads>>>(d_in, d_out, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

    float ms;
    cudaEventElapsedTime(&ms, start, stop);

    // 带宽 = 总字节数 / 时间（秒）
    // 注意：读+写，所以总字节数是 2 * bytes
    float bandwidth = (2.0f * bytes) / (ms / 1000.0f) / 1e9f;
    printf("带宽: %.2f GB/s\n", bandwidth);

    cudaFree(d_in);
    cudaFree(d_out);
    free(h_in);
    return 0;
}

小提示：计算带宽时，读和写都要算进去。上面代码中，每个线程读一次、写一次，所以总数据量是2 * bytes。如果你只测读带宽，那就只算读的部分。

3.3 对齐访问：别让数据「错位」

对齐访问是什么意思？说白了，就是每个线程访问的地址，最好是它所在warp中所有线程地址的整数倍。GPU的L1/L2缓存和内存控制器，都是以128字节（cache line）为单位传输数据的。

我曾经踩过一个坑：有个同事写kernel时，故意让每个线程从偏移1字节的位置开始读。结果带宽直接掉到理论值的1/4。为什么？因为一个128字节的cache line，本来可以服务32个线程（每个读4字节），结果因为错位，一个cache line只能服务几个线程，剩下的都得重新请求。

警告：全局内存访问必须对齐到32字节（一个warp的访问粒度）。如果不对齐，硬件会自动拆分成多次访问，带宽直接腰斩甚至更低。

怎么保证对齐？很简单：

使用cudaMalloc分配的内存，默认是256字节对齐的
每个线程的访问地址 = base + threadIdx.x * sizeof(type)，只要base对齐，且sizeof(type)是2的幂，就自然对齐
避免使用char或short类型做全局内存访问（除非你清楚自己在做什么）

3.4 合并访问：让warp「抱团」读写

合并访问是GPU带宽优化的核心。一个warp有32个线程，如果这32个线程访问的地址是连续的，那么硬件可以合并成一次或几次大的内存事务。反之，如果地址是分散的，那就得拆成32次独立访问。

你想想看，32次访问 vs 1次访问，带宽差距就是32倍。我见过最夸张的例子：一个做粒子模拟的项目，因为粒子位置存储是随机的，导致访存完全无法合并，带宽利用率不到5%。后来改成按空间排序存储，带宽直接拉满。

来看一个对比实验：

// 合并访问版本（推荐）
__global__ void coalesced_read(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx];  // 连续地址：thread 0读in[0], thread 1读in[1]...
    }
}

// 非合并访问版本（不推荐）
__global__ void strided_read(float *in, float *out, int N, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx * stride < N) {
        out[idx] = in[idx * stride];  // 间隔stride个元素
    }
}

我建议你亲自跑一下这两个版本。当stride=1时，两个版本其实一样（合并）。当stride=2时，带宽直接掉一半。stride=4时，再掉一半。为什么会这样？因为每个cache line里只有一半的数据是需要的，另一半被浪费了。

3.5 实战：测量不同访存模式的带宽

我们来做一个完整的测试，看看不同模式下的实际带宽。我一般用这个测试来快速评估GPU的「健康状态」。

访存模式	代码示例	预期带宽（以RTX 4090为例）	说明
合并+对齐	out[i] = in[i]	~1000 GB/s	理想情况，接近理论值
合并+不对齐	out[i] = in[i+1]	~750 GB/s	偏移1个元素，仍可合并但cache line利用率下降
非合并（stride=2）	out[i] = in[i*2]	~500 GB/s	一半数据浪费
随机访问	out[i] = in[random[i]]	~100 GB/s	完全无法合并，带宽惨不忍睹

关键结论：合并访问 + 对齐访问 = 满带宽。任何偏离都会导致带宽下降。实际项目中，我们往往无法做到100%合并，但至少要保证主要数据路径是合并的。

3.6 知识体系：带宽测试的核心逻辑

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

3.7 避坑指南：我踩过的那些坑

做带宽测试这么多年，我总结了几条血泪教训：

不要用malloc分配的内存做测试。 我曾经犯过这个错——用host malloc分配，然后cudaMemcpy到device。结果测出来的带宽只有实际的一半，因为PCIe带宽成了瓶颈。
记得加cudaDeviceSynchronize()。 如果不加，你测的可能是kernel启动时间，而不是执行时间。嗯，这个坑我至少见过5个新手踩过。
数据量要够大。 至少几十MB，否则L2缓存会「作弊」，让你测出远超理论带宽的假数据。我一般用64MB以上。
注意warp divergence。 如果一个warp里有些线程不执行访存指令，那带宽利用率也会下降。虽然这不直接影响合并，但会降低有效带宽。

实用技巧：用nvidia-smi查看当前显存频率和温度。温度过高会导致显存降频，带宽直接打7折。我一般在测试前先跑个「热身」kernel，让GPU进入稳定状态。

3.8 小结

全局内存带宽测试，是GPU性能优化的基本功。记住三个关键词：对齐、合并、大事务。对齐保证地址正确，合并保证带宽利用率，大事务减少协议开销。下次你写kernel时，先问问自己：我的访存模式是合并的吗？地址对齐了吗？如果答案都是yes，那你的带宽大概率不会差。