2、带宽测试工具介绍

做GPU性能优化，第一件事是什么？

我个人习惯，先摸清显存带宽的底牌。你想想看，如果连硬件能跑多快都不清楚，后面再怎么调优都是盲人摸象。这一章，我就带大家认识几款常用的带宽测试工具——从官方工具到我们自己写的微基准测试，挨个过一遍。

2.1 nvidia-smi：最基础的带宽观测工具

nvidia-smi 大家应该都不陌生。平时看显存占用、温度、功耗，都用它。但很多人不知道，它也能看带宽相关的指标。

嗯，这里要注意：nvidia-smi 不直接显示带宽利用率。它只能给你一个侧面参考。我在项目中遇到过好几次，有人看到 GPU-Util 100% 就以为带宽吃满了，结果一测实际带宽才用了 60%。

那怎么看带宽利用率呢？可以用 nvidia-smi dmon 或者 nvidia-smi pmon，它们能显示更细粒度的指标。比如 dmon 里的 mem_util 字段，就是显存控制器利用率——这个才接近真实的带宽占用。

# 查看实时显存控制器利用率
nvidia-smi dmon -s pucvmet -d 1

输出大概长这样：

# gpu   pwr gtemp mtemp    sm   mem   enc   dec  mclk  pclk
# Idx     W     C     C     %     %     %     %   MHz   MHz
    0   250    75    72    98    85     0     0  5000  1350

看到 mem 那列了吗？85% 说明显存控制器已经比较忙了。但注意，这也不是精确的带宽利用率——它反映的是控制器繁忙程度，不是实际数据传输量。

2.2 bandwidthTest：官方带宽测试程序

NVIDIA 官方提供了 bandwidthTest，在 CUDA 安装目录的 samples/1_Utilities/bandwidthTest 下。这个工具简单粗暴，直接测主机到设备、设备到主机、设备到设备三种传输模式的带宽。

我个人觉得，这是入门带宽测试最好的起点。为什么？因为它帮你屏蔽了底层细节，你只需要编译运行，就能拿到一个可信的基准值。

# 编译
cd /path/to/cuda-samples/Samples/1_Utilities/bandwidthTest
make

# 运行
./bandwidthTest

输出示例：

[CUDA Bandwidth Test] - Starting...
Running on...

 Device 0: Tesla V100-SXM2-32GB
 Quick Mode

 Host to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(GB/s)
   33554432                12.45

 Device to Host Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(GB/s)
   33554432                13.21

 Device to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(GB/s)
   33554432                898.76

看到没？Device to Device 的带宽是 898 GB/s，接近 V100 的理论带宽 900 GB/s。这说明显存本身没问题。但 Host to Device 只有 12 GB/s——这是 PCIe 3.0 x16 的瓶颈。

常用的参数组合：

参数	说明	示例
--memory=pinned	使用页锁定内存	必须加，否则结果不准
--mode=shmoo	测试不同传输大小的带宽	看带宽随传输量变化
--device=0	指定 GPU 设备	多卡环境必用
--htod / --dtoh	只测单向传输	节省时间

2.3 自定义CUDA微基准测试

官方工具虽好，但有时候不够灵活。比如我想测特定数据布局下的带宽，或者想看看不同访存模式对带宽的影响——这时候就得自己写微基准测试了。

说白了，自定义测试的核心就是：分配显存 → 启动 kernel → 计时 → 算带宽。我给大家一个最简版本：

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

__global__ void bandwidth_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 << 28;  // 256M 个 float
    size_t bytes = N * sizeof(float);

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

    // 预热
    bandwidth_kernel<<<(N+255)/256, 256>>>(d_in, d_out, N);
    cudaDeviceSynchronize();

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

    cudaEventRecord(start);
    bandwidth_kernel<<<(N+255)/256, 256>>>(d_in, d_out, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

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

    double bw = (bytes * 2) / (ms / 1000.0) / 1e9;  // 读+写 = 2倍
    printf("Bandwidth: %.2f GB/s\n", bw);

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

这个代码很简单，就是每个线程读一个 float，写一个 float。但注意：实际带宽要算读+写，所以是 bytes * 2。我刚开始写的时候忘了乘 2，结果带宽只有理论值的一半，查了半天才发现。

举个例子，跨步访问的 kernel 长这样：

__global__ void stride_kernel(float *in, float *out, int N, int stride) {
    int idx = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
    if (idx < N) {
        out[idx] = in[idx];
    }
}

你猜 stride=2 时带宽会掉多少？我实测过，在 V100 上 stride=2 时带宽降到连续访问的 50% 左右。stride=4 时只剩 25%。这就是 合并访问 的重要性——GPU 喜欢连续的数据，不喜欢跳着读。

2.4 工具对比与选择建议

说了这么多，到底该用哪个？我整理了一张表：

工具	适用场景	优点	缺点
nvidia-smi	快速查看显存状态	无需安装，实时监控	不直接显示带宽
bandwidthTest	获取标准带宽基准	官方可靠，一键运行	不能自定义访存模式
自定义微基准	深入分析访存行为	完全可控，灵活	需要编程，容易出错

我的建议是：先用 bandwidthTest 拿到基线，再用自定义测试深挖细节。比如你发现程序性能差，先用 bandwidthTest 确认硬件带宽没问题，然后写个微基准测试，看看是不是访存模式导致的。

2.5 本章小结

这一章我们聊了三类工具：

• nvidia-smi：日常监控，看个大概
• bandwidthTest：官方基准，快速定位硬件问题
• 自定义微基准：深入分析，找到访存瓶颈

记住一句话：没有万能工具，只有合适的场景。下一章我们会深入带宽测试的具体方法，包括如何设计测试用例、如何解读结果。但在此之前，我建议你先动手跑一遍 bandwidthTest，看看你的 GPU 到底能跑多快。

嗯，就到这里。有问题随时交流。