nvbandwidth 深度解析：NVIDIA GPU 带宽测量工具全指南

nvbandwidth 深度解析：NVIDIA GPU 带宽测量工具全指南

一、工具概述

1.1 什么是 nvbandwidth？

nvbandwidth 是 NVIDIA 官方开发的一款专门用于测量 GPU 内存带宽的高性能工具。简单来说，它把“带宽到底跑到多少”这件事讲清楚，也让不同传输路径的性能差异一目了然。它能够精确量化各种数据传输场景下的带宽表现，包括：

• • PCIe 带宽：CPU 与 GPU 之间的数据传输
• • NVLink 带宽：GPU 与 GPU 之间的高速互联
• • IMEX 支撑的跨节点访问：多节点 GPU 内存访问场景

1.2 核心特点

当前版本：v0.8（基于 version.h）

二、环境搭建与使用说明

2.1 系统要求

软件依赖：

• • CUDA Toolkit: 11.x+（多节点需 12.3+）
• • NVIDIA Driver: 550+（多节点）
• • 编译器: GCC 7.x+（支持 C++17）
• • CMake: 3.20+
• • Boost: libboost-program-options-dev

注：驱动兼容性说明：CUDA Toolkit 版本必须与 NVIDIA 驱动版本兼容。例如：

• • 驱动 570.x 支持 CUDA ≤ 12.8
• • CUDA 13.x 需要更新的驱动版本
• • GDS (GPUDirect Storage) 用户可能需要特定 CUDA 版本

使用 nvidia-smi 查看驱动支持的 CUDA 版本。如需多版本共存，可安装多个 CUDA Toolkit 到不同目录（如 /usr/local/cuda-12.8 和 /usr/local/cuda-13.1），编译时指定路径：

# 使用特定 CUDA 版本编译
PATH=/usr/local/cuda-12.8/bin:$PATH cmake ..
make -j$(nproc)

硬件要求：

• • CUDA-enabled GPU
• • 兼容的 NVIDIA 显示驱动

2.2 安装步骤

Ubuntu/Debian：

# 安装依赖
apt install libboost-program-options-dev

# 或使用提供的脚本一键安装（会同时安装依赖并编译）
sudo ./debian_install.sh

Fedora：

sudo dnf -y install boost-devel

编译构建：

# 方式一：源码目录内构建（单节点版本）
cmake .
make -j$(nproc)

# 方式二：独立构建目录（推荐，保持源码整洁）
mkdir build && cd build
cmake ..
make -j$(nproc)

# 多节点版本（需要 MPI 环境）
cmake -DMULTINODE=1 .
make -j$(nproc)

构建输出示例：

-- The CUDA compiler identification is NVIDIA 13.1.115
-- The CXX compiler identification is GNU 13.3.0
-- Found Boost: /usr/lib/x86_64-linux-gnu/cmake/Boost-1.83.0
-- Detecting CUDA Arch to set CMAKE_CUDA_ARCHITECTURES
-- Configuring done
-- Generating done
-- Build files have been written to: /path/to/build

[  8%] Building CXX object CMakeFiles/nvbandwidth.dir/testcase.cpp.o
[ 33%] Building CUDA object CMakeFiles/nvbandwidth.dir/kernels.cu.o
...
[100%] Linking CXX executable nvbandwidth
[100%] Built target nvbandwidth

构建产物：

• • 可执行文件：./nvbandwidth（构建目录内）
• • 支持 CUDA 架构：70, 75, 80, 86, 89, 90, 100（CUDA 13.0+ 不再支持 52）

依赖检查：

# 检查 CUDA 版本
nvcc --version

# 检查 CMake 版本
cmake --version

# 检查 Boost 是否安装
dpkg -l | grep boost  # Debian/Ubuntu

2.3 命令行参数详解

./nvbandwidth -h

2.4 使用示例

运行所有测试：

./nvbandwidth

运行特定测试：

./nvbandwidth -t device_to_device_memcpy_read_ce

指定缓冲区大小和迭代次数：

./nvbandwidth -b 1024 -i 5

多节点测试（需先启动 IMEX 服务）：

# 启动 IMEX 服务
sudo systemctl start nvidia-imex.service

# 运行多节点测试
mpirun --allow-run-as-root --map-by ppr:4:node --bind-to core -np 8 \
  --hostfile /etc/nvidia-imex/nodes_config.cfg \
  ./nvbandwidth -p multinode

三、源码架构与原理分析

3.1 整体架构图

nvbandwidth
├── 主程序入口 (nvbandwidth.cpp)
│   ├── 命令行参数解析（Boost program_options）
│   ├── 测试用例管理与调度
│   └── 结果输出（Output / JsonOutput）
├── 测试用例层 (testcase.h / testcase.cpp)
│   ├── Testcase 基类定义
│   └── 辅助方法（allToOneHelper / latencyHelper 等）
├── CE 测试用例 (testcases_ce.cpp)
│   ├── HostToDeviceCE / DeviceToHostCE
│   ├── DeviceToDeviceReadCE / WriteCE
│   └── AllToOne / OneToAll 等聚合测试
├── SM 测试用例 (testcases_sm.cpp)
│   ├── HostToDeviceSM / DeviceToHostSM
│   ├── DeviceToDeviceReadSM / WriteSM
│   └── Latency 测试
├── 内存操作层 (memcpy.h / memcpy.cpp)
│   ├── MemcpyBuffer（内存缓冲区抽象）
│   ├── HostBuffer（主机内存，NUMA 亲和性）
│   ├── DeviceBuffer（设备内存，Peer Access）
│   ├── MemcpyOperation（拷贝操作编排）
│   └── MemcpyInitiator（CE/SM 拷贝发起者）
├── CUDA 内核层 (kernels.cuh / kernels.cu)
│   ├── simpleCopyKernel（简单拷贝，小缓冲区）
│   ├── stridingMemcpyKernel（跨步拷贝，大缓冲区优化）
│   ├── splitWarpCopyKernel（双向拷贝）
│   ├── ptrChasingKernel（延迟测量）
│   ├── spinKernel（阻塞同步）
│   └── memsetKernel / memcmpKernel（数据校验）
├── 输出层 (output.h / output.cpp)
│   ├── Output（文本输出）
│   └── JsonOutput（JSON 格式输出）
└── 多节点支持 (multinode_memcpy.h / multinode_memcpy.cpp)
    ├── MultinodeDeviceBuffer（多节点缓冲区抽象）
    ├── MultinodeMemoryAllocationUnicast（单播内存）
    ├── MultinodeMemoryAllocationMulticast（多播内存）
    └── NodeHelperMulti（MPI 同步辅助）

3.2 核心类设计

3.2.1 Testcase 基类

class Testcase {
 protected:
    std::string key;      // 测试用例唯一标识
    std::string desc;     // 测试描述

    // 静态过滤方法
static bool filterHasAccessiblePeerPairs();    // 检查是否有可访问的对等 GPU
static bool filterSupportsMulticast();          // 检查是否支持多播

    // 辅助方法
void allToOneHelper(...);    // AllToOne 测试辅助
void oneToAllHelper(...);    // OneToAll 测试辅助
void latencyHelper(...);     // 延迟测试辅助

 public:
    Testcase(std::string key, std::string desc);
    virtual ~Testcase() {}

    std::string testKey();
    std::string testDesc();

virtual bool filter(){ return true; }  // 过滤条件检查
virtual void run(unsigned long long size, unsigned long long loopCount)= 0;
};

所有具体测试用例都继承自 Testcase 基类，实现了自己的 run() 方法。filter() 方法用于检查当前系统是否满足测试条件（如是否有对等 GPU 可访问）。

3.2.2 内存缓冲区抽象

// 基类
class MemcpyBuffer {
 protected:
    void* buffer{};
    size_t bufferSize;
 public:
virtual int getBufferIdx() const= 0;
virtual CUcontext getPrimaryCtx() const= 0;
};

// 主机内存
class HostBuffer : public MemcpyBuffer {
    // 使用 cuMemHostAlloc 分配可分页锁定内存
    // 自动设置 NUMA 亲和性
};

// 设备内存
class DeviceBuffer : public MemcpyBuffer {
    int deviceIdx;
    CUcontext primaryCtx;
    // 使用 cuMemAlloc 分配设备内存
};

3.2.3 拷贝操作抽象

class MemcpyOperation {
    std::shared_ptr<NodeHelper> nodeHelper;
    std::shared_ptr<MemcpyInitiator> memcpyInitiator;

 public:
double doMemcpy(const MemcpyBuffer &src, const MemcpyBuffer &dst);
    std::vector<double> doMemcpyVector(...);
};

MemcpyInitiator 有两个主要实现：

• • MemcpyInitiatorCE：使用 cuMemcpyAsync 进行拷贝
• • MemcpyInitiatorSM：使用自定义 CUDA 内核进行拷贝

3.3 测量原理详解

3.3.1 Spin Kernel 机制

为了排除 CUDA 内核启动和流队列的调度开销，nvbandwidth 采用了 Spin Kernel 技术：

// 1. 重置阻塞变量
*blockingVarHost = 0;

// 2. 启动 Spin Kernel，GPU 在该流上等待信号
spinKernel(blockingVarHost, stream);  // GPU 自旋等待 blockingVarHost == 1

// 3. 在 Spin Kernel 运行期间，入队所有测量事件
//    这些操作会被 GPU 驱动排队，但尚未执行
cuEventRecord(startEvent, stream);
// ... 执行 loopCount 次拷贝操作 ...
cuEventRecord(endEvent, stream);

// 4. 释放信号，Spin Kernel 退出，测量真正开始
*blockingVarHost = 1;  // CPU 写入，GPU 检测到后退出 Spin Kernel

// 5. 等待所有操作完成
cuStreamSynchronize(stream);

预热机制：在正式测量前，会先执行 WARMUP_COUNT（默认 4 次）次预热拷贝，确保缓存和流水线处于稳定状态。

关键点：Spin Kernel 确保 CUDA Events 和所有拷贝操作在 GPU 上同时就绪后才真正开始执行，从而排除了 CPU 端入队开销对测量的影响。

流程图：

CPU Side                    GPU Side
─────────────────────────────────────────────────
spinKernel(blockingVar) ──→ 等待 blockingVar == 1
                            │
cuEventRecord(start) ─────→ 记录开始时间戳
                            │
memcpy operations ────────→ 执行拷贝
                            │
cuEventRecord(end) ───────→ 记录结束时间戳
                            │
*blockingVar = 1 ─────────→ Spin Kernel 退出
                            │
cuStreamSynchronize ──────→ 等待完成

3.3.2 带宽计算

基本公式：

带宽 (GB/s) = (拷贝数据大小 (字节) × 循环次数) / 耗时 (微秒) / 1000

推导过程：

1. 1. 总传输量 (字节) = 拷贝数据大小 × 循环次数
2. 2. 传输速率 (字节/秒) = 总传输量 / (耗时 × 10⁻⁶) = 总传输量 × 10⁶ / 耗时
3. 3. 传输速率 (GB/s) = 传输速率 (字节/秒) / 10⁹ = 总传输量 / 耗时 / 1000

CE 拷贝带宽：

// 使用 CUDA Events 计时，elapsedWithEventsInUs 单位为微秒
float timeMs;
cuEventElapsedTime(&timeMs, start, end);
double elapsedWithEventsInUs = timeMs * 1000.0;  // 毫秒转微秒
// bandwidth 单位为 B/s，后续除以 1e9 转换为 GB/s
unsigned long long bandwidth = (size * loopCount * 1000ull * 1000ull) / elapsedWithEventsInUs;

SM 拷贝带宽：
SM 拷贝会将数据大小截断以适应 GPU 的 SM 数量，确保每个线程处理的数据量一致：

// 实际拷贝大小计算
// 1. 获取 GPU 的 SM 数量和每块线程数
int numSm;
cuDeviceGetAttribute(&numSm, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, dev);
unsigned int totalThreadCount = numSm * numThreadPerBlock;  // numThreadPerBlock = 512

// 2. 对于大缓冲区（>= 64 MiB），使用优化内核并截断
if (size >= (smallBufferThreshold * _MiB)) {
    size_t sizeInElement = size / sizeof(uint4);  // 转换为 uint4 元素数
    // 截断为 totalThreadCount 的整数倍
    sizeInElement = totalThreadCount * (sizeInElement / totalThreadCount);
    return sizeInElement * sizeof(uint4);
}

截断原因：优化的 stridingMemcpyKernel 使用跨步访问模式，要求每个线程处理相同数量的元素，以达到最佳带宽性能。

3.4 CUDA 内核实现

3.4.1 简单拷贝内核

__global__ void simpleCopyKernel(unsigned long long loopCount, uint4 *dst, uint4 *src) {
    for (unsigned int i = 0; i < loopCount; i++) {
        const int idx = blockIdx.x * blockDim.x + threadIdx.x;
        size_t offset = idx * sizeof(uint4);
        uint4* dst_uint4 = reinterpret_cast<uint4*>((char*)dst + offset);
        uint4* src_uint4 = reinterpret_cast<uint4*>((char*)src + offset);
        // 使用全局缓存加载/存储指令
        __stcg(dst_uint4, __ldcg(src_uint4));
    }
}

3.4.2 跨步拷贝内核（优化版本）

__global__ void stridingMemcpyKernel(unsigned int totalThreadCount,
                                      unsigned long long loopCount,
                                      uint4* dst, uint4* src,
                                      size_t chunkSizeInElement) {
    // 每个线程处理一个跨步的数据块
    unsigned long long from = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned long long bigChunkSizeInElement = chunkSizeInElement / 12;
    dst += from;
    src += from;
    uint4* dstBigEnd = dst + (bigChunkSizeInElement * 12) * totalThreadCount;
    uint4* dstEnd = dst + chunkSizeInElement * totalThreadCount;

    for (unsigned int i = 0; i < loopCount; i++) {
        uint4* cdst = dst;
        uint4* csrc = src;

        // 手动展开，每次处理 12 个 uint4（192 字节）
        while (cdst < dstBigEnd) {
            uint4 pipe_0 = *csrc; csrc += totalThreadCount;
            // ... 省略中间 pipe_1 ~ pipe_10 ...
            uint4 pipe_11 = *csrc; csrc += totalThreadCount;

            *cdst = pipe_0; cdst += totalThreadCount;
            // ... 省略中间写入 ...
            *cdst = pipe_11; cdst += totalThreadCount;
        }

        // 处理剩余尾部
        while (cdst < dstEnd) {
            *cdst = *csrc; cdst += totalThreadCount; csrc += totalThreadCount;
        }
    }
}

优化要点：

1. 1. 循环展开：减少分支预测失败
2. 2. 流水线处理：先批量加载，再批量存储
3. 3. 跨步访问：充分利用内存带宽

3.4.3 双向拷贝内核（Split Warp）

__global__ void splitWarpCopyKernel(unsigned long long loopCount, uint4 *dst, uint4 *src) {
    for (unsigned int i = 0; i < loopCount; i++) {
        unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
        unsigned int globalWarpId = idx / warpSize;
        unsigned int warpLaneId = idx % warpSize;

        // 奇偶 Warp 交替拷贝方向
        if (globalWarpId & 0x1) {
            // 奇数 Warp：src -> dst
            dst_uint4 = dst + (globalWarpId * warpSize + warpLaneId);
            src_uint4 = src + (globalWarpId * warpSize + warpLaneId);
        } else {
            // 偶数 Warp：dst -> src
            dst_uint4 = src + (globalWarpId * warpSize + warpLaneId);
            src_uint4 = dst + (globalWarpId * warpSize + warpLaneId);
        }

        __stcg(dst_uint4, __ldcg(src_uint4));
    }
}

3.5 延迟测量原理

使用**指针追踪（Pointer Chase）**技术测量访问延迟：

struct LatencyNode {
    struct LatencyNode *next;
};

__global__ void ptrChasingKernel(struct LatencyNode *data, size_t size,
                                  unsigned int accesses, unsigned int targetBlock) {
    struct LatencyNode *p = data;
    if (blockIdx.x != targetBlock) return;

    // 顺序访问链表，产生依赖链
    for (auto i = 0; i < accesses; ++i) {
        p = p->next;  // 每次访问依赖上一次结果
    }

    // 防止编译器优化
    if (p == nullptr) __trap();
}

链表初始化：源码中使用**跨步模式（Stride Pattern）**初始化链表，跨步长度为 16：

void Testcase::latencyHelper(const MemcpyBuffer &dataBuffer, bool measureDeviceToDeviceLatency){
    uint64_t n_ptrs = dataBuffer.getBufferSize() / sizeof(struct LatencyNode);

    if (measureDeviceToDeviceLatency) {
        // 设备侧链表：写入设备地址
        for (uint64_t i = 0; i < n_ptrs; i++) {
            struct LatencyNode node;
            size_t nextOffset = ((i + strideLen) % n_ptrs) * sizeof(struct LatencyNode);
            node.next = (struct LatencyNode*)(dataBuffer.getBuffer() + nextOffset);
            cuMemcpyHtoD(dataBuffer.getBuffer() + i * sizeof(struct LatencyNode),
                         &node, sizeof(struct LatencyNode));
        }
    } else {
        // 主机侧链表：写入主机地址
        struct LatencyNode* hostMem = (struct LatencyNode*)dataBuffer.getBuffer();
        for (uint64_t i = 0; i < n_ptrs; i++) {
            hostMem[i].next = &hostMem[(i + strideLen) % n_ptrs];
        }
    }
}

原理：通过构建一个固定跨步的链表，强制 GPU 进行串行内存访问，从而测量真实的访问延迟。固定跨步模式可以减少缓存局部性影响，使测量结果更具代表性。

注意：延迟测试固定使用 2 MiB 缓冲区，--bufferSize 参数对延迟测试无效。源码中会强制设置：

// nvbandwidth.cpp
if (test->testKey() == "host_device_latency_sm" || test->testKey() == "device_to_device_latency_sm") {
    test->run(2 * _MiB, loopCount);  // 固定使用 2MB 缓冲区
}

两种延迟测试的区别：

四、测试用例详解

4.1 单节点测试用例

Host ↔ Device 基础测试：

Host ↔ Device 双向测试：

Device ↔ Device 测试（需要 P2P 访问能力）：

聚合带宽测试：

延迟与本地拷贝测试：

说明：单节点版本共 35 个测试用例（索引 0-34），使用 ./nvbandwidth -l 查看完整列表。

4.2 双向测试原理

CE 双向测试：

• • 使用两个独立的 CUDA Stream
• • Host↔Device 双向：仅第一条拷贝（被测量的流）计入结果，另一条为干扰流量
• • Device↔Device 双向：仅测量流的带宽被报告，干扰流在相反方向同时运行但不被测量

SM 双向测试：

• • 使用 Split Warp 内核（splitWarpCopyKernel）
• • 奇数 Warp 执行正向拷贝，偶数 Warp 执行反向拷贝
• • Host↔Device 双向 SM 测试：报告的带宽值为测量流带宽的一半（代码中 getAdjustedBandwidth 返回 bandwidth / 2），用于估算双向同时传输时每个方向的有效带宽
• • Device↔Device 双向 SM 测试：输出三组数据
• • read1 / write1：第一个方向的带宽
• • read2 / write2：第二个方向的带宽
• • total：两个方向带宽之和

4.3 多节点测试

多节点版本需要：

1. 1. MPI 环境：用于跨节点进程通信
2. 2. IMEX 服务：NVIDIA Internode Memory Exchange Service
3. 3. NVLink 多节点部署：物理连接支持

多节点测试用例（共 13 个，需要在编译时启用 MULTINODE=1）：

关键类：

// 多节点设备缓冲区基类
class MultinodeDeviceBuffer : public MemcpyBuffer {
    int MPI_rank;  // 内存所在的节点（MPI rank）

    CUcontext getPrimaryCtx() const override;  // 返回本地 GPU 上下文
int getMPIRank() const override;            // 返回所属 MPI rank
};

// 单播内存（点对点访问）
// 内存物理上位于某个节点，其他节点可以远程访问
class MultinodeMemoryAllocationUnicast {
    CUmemGenericAllocationHandle handle;  // 内存分配句柄
    CUmemFabricHandle fh;                 // 跨节点共享句柄
    // 使用 cuMemExportToShareableHandle 导出，通过 MPI_Bcast 共享
};

// 多播内存（广播）
// 写入时自动传播到所有节点，适用于广播场景
class MultinodeMemoryAllocationMulticast {
    CUmemGenericAllocationHandle multicastHandle;  // 多播对象句柄
    CUmulticastObjectProp multicastProp;           // 多播属性
    // 使用 cuMulticastCreate 创建，cuMulticastBindMem 绑定
};

// 本地设备内存（仅单个节点可访问）
class MultinodeDeviceBufferLocal : public MultinodeDeviceBuffer {
    // 仅在 MPI_rank == worldRank 时分配实际内存
};

内存共享流程：

1. 1. 单播内存：
• • 源节点调用 cuMemCreate 分配内存
• • 通过 cuMemExportToShareableHandle 导出句柄
• • 使用 MPI_Bcast 广播句柄到其他节点