2. NVLink带宽基准测试

任务目标

编写一个CUDA程序，通过NVLink点对点通信在两个GPU之间传输数据，测量实际带宽，并与PCIe带宽进行对比。程序需支持选择使用NVLink或PCIe路径，并输出可复现的带宽结果。

环境依赖

NVIDIA GPU（双卡，支持NVLink，如A100/A800/H800等）
CUDA Toolkit ≥ 11.0（建议12.x）
Linux x86_64（或WSL2）
GPU驱动 ≥ 450.80.02

实现原理与代码

使用CUDA的cudaMemcpyPeerAsync实现设备间直接拷贝（NVLink路径），使用cudaMemcpy经主机中转实现PCIe路径对比。通过cudaDeviceEnablePeerAccess启用UVA（统一虚拟地址）以支持点对点访问。带宽计算基于传输字节数除以实测耗时。

// nvlink_bw_test.cu
#include <cstdio>
#include <cstdlib>
#include <chrono>
#include <cuda_runtime.h>

#define CHECK_CUDA(call) do { \
    cudaError_t err = (call); \
    if (err != cudaSuccess) { \
        printf("CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
               cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

// 传输大小（字节）：从1MB到1GB，按2倍递增
const size_t SIZES[] = {
    1ULL<<20, 2ULL<<20, 4ULL<<20, 8ULL<<20,
    16ULL<<20, 32ULL<<20, 64ULL<<20,
    128ULL<<20, 256ULL<<20, 512ULL<<20,
    1ULL<<30
};
const int NUM_SIZES = sizeof(SIZES) / sizeof(SIZES[0]);

// 测量一次拷贝耗时（毫秒）
float measure_transfer(void* dst, int dst_dev, void* src, int src_dev,
                       size_t bytes, cudaStream_t stream) {
    cudaEvent_t start, stop;
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));

    CHECK_CUDA(cudaEventRecord(start, stream));
    CHECK_CUDA(cudaMemcpyPeerAsync(dst, dst_dev, src, src_dev, bytes, stream));
    CHECK_CUDA(cudaEventRecord(stop, stream));
    CHECK_CUDA(cudaStreamSynchronize(stream));

    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    return ms;
}

// 通过主机中转（PCIe路径）
float measure_pcie_transfer(void* dst, void* src, size_t bytes,
                            cudaStream_t stream) {
    void* host_buf = nullptr;
    CHECK_CUDA(cudaMallocHost(&host_buf, bytes));

    cudaEvent_t start, stop;
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));

    CHECK_CUDA(cudaEventRecord(start, stream));
    CHECK_CUDA(cudaMemcpyAsync(host_buf, src, bytes,
                               cudaMemcpyDeviceToHost, stream));
    CHECK_CUDA(cudaMemcpyAsync(dst, host_buf, bytes,
                               cudaMemcpyHostToDevice, stream));
    CHECK_CUDA(cudaEventRecord(stop, stream));
    CHECK_CUDA(cudaStreamSynchronize(stream));

    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    CHECK_CUDA(cudaFreeHost(host_buf));
    return ms;
}

int main() {
    int dev_count = 0;
    CHECK_CUDA(cudaGetDeviceCount(&dev_count));
    if (dev_count < 2) {
        printf("需要至少2块GPU，当前检测到%d块\n", dev_count);
        return EXIT_FAILURE;
    }

    int dev0 = 0, dev1 = 1;
    CHECK_CUDA(cudaSetDevice(dev0));
    CHECK_CUDA(cudaDeviceEnablePeerAccess(dev1, 0));
    CHECK_CUDA(cudaSetDevice(dev1));
    CHECK_CUDA(cudaDeviceEnablePeerAccess(dev0, 0));

    // 检查是否支持点对点（NVLink）
    int can_access = 0;
    CHECK_CUDA(cudaDeviceCanAccessPeer(&can_access, dev0, dev1));
    if (!can_access) {
        printf("警告：GPU %d 和 %d 之间不支持点对点访问，"
               "NVLink路径将不可用\n", dev0, dev1);
    }

    cudaStream_t stream;
    CHECK_CUDA(cudaStreamCreate(&stream));

    printf("%-12s %-15s %-15s %-15s\n",
           "大小(MB)", "NVLink(GB/s)", "PCIe(GB/s)", "加速比");
    printf("------------------------------------------------------------\n");

    for (int i = 0; i < NUM_SIZES; i++) {
        size_t bytes = SIZES[i];
        double mb = (double)bytes / (1024.0 * 1024.0);

        // 分配设备内存
        void *d_src = nullptr, *d_dst = nullptr;
        CHECK_CUDA(cudaSetDevice(dev0));
        CHECK_CUDA(cudaMalloc(&d_src, bytes));
        CHECK_CUDA(cudaMemset(d_src, 0xAB, bytes));
        CHECK_CUDA(cudaSetDevice(dev1));
        CHECK_CUDA(cudaMalloc(&d_dst, bytes));

        // NVLink路径（点对点）
        float nvlink_ms = 0.0f;
        if (can_access) {
            nvlink_ms = measure_transfer(d_dst, dev1, d_src, dev0,
                                         bytes, stream);
        }

        // PCIe路径（经主机中转）
        float pcie_ms = measure_pcie_transfer(d_dst, d_src, bytes, stream);

        double nvlink_gbps = 0.0, pcie_gbps = 0.0;
        if (nvlink_ms > 0.0f) {
            nvlink_gbps = (double)bytes / (nvlink_ms / 1000.0) / 1e9;
        }
        if (pcie_ms > 0.0f) {
            pcie_gbps = (double)bytes / (pcie_ms / 1000.0) / 1e9;
        }

        double speedup = (pcie_ms > 0.0f && nvlink_ms > 0.0f) ?
                         pcie_ms / nvlink_ms : 0.0;

        printf("%-12.1f %-15.2f %-15.2f %-15.2f\n",
               mb, nvlink_gbps, pcie_gbps, speedup);

        CHECK_CUDA(cudaFree(d_src));
        CHECK_CUDA(cudaFree(d_dst));
    }

    CHECK_CUDA(cudaStreamDestroy(stream));
    CHECK_CUDA(cudaSetDevice(dev0));
    CHECK_CUDA(cudaDeviceDisablePeerAccess(dev1));
    CHECK_CUDA(cudaSetDevice(dev1));
    CHECK_CUDA(cudaDeviceDisablePeerAccess(dev0));

    return EXIT_SUCCESS;
}

编译与运行

nvcc -arch=sm_80 -O3 nvlink_bw_test.cu -o nvlink_bw_test
./nvlink_bw_test

注意：-arch需根据实际GPU架构调整（A100为sm_80，H100为sm_90）。

验证方法与判定标准

功能验证：程序正常输出表格，无CUDA错误；NVLink列在支持点对点的GPU上应显示非零带宽值。
带宽合理性：NVLink 3.0（A100）理论带宽为600GB/s（双向），实测单向通常为250～300GB/s；NVLink 4.0（H100）理论900GB/s，实测约400～450GB/s。PCIe 4.0 x16单向理论约32GB/s，实测约25GB/s。
加速比：NVLink带宽应显著高于PCIe（通常5～10倍），加速比列应明显大于1。
可重复性：多次运行结果波动应在±5%以内（排除系统负载干扰）。

常见问题

点对点访问失败：确认GPU拓扑为NVLink直连（运行nvidia-smi topo -m查看），且驱动支持UVA。
带宽偏低：检查是否使用了cudaMemcpyPeerAsync（同步版本会阻塞流水线）；确认GPU处于持久模式且未被其他进程占用。
编译错误：若使用CUDA 11.x，需将cudaMemcpyPeerAsync替换为cudaMemcpyAsync并显式指定cudaMemcpyDeviceToDevice（在启用P2P后两者等价）。

📚 课程目录