4. NVLink与PCIe混合通信性能对比

任务说明

本案例在包含NVLink与PCIe两种互联的混合拓扑中，分别使用两种路径完成相同大小的数据块传输，测量并对比带宽与延迟。目标是通过可复现的基准测试，量化NVLink相对PCIe的性能优势，并验证两种路径在混合拓扑中的实际可用性。

测试拓扑与数据布局

采用双GPU节点（GPU0与GPU1），两者同时具备NVLink直连和PCIe Switch上行连接。数据布局为：GPU0显存中生成源数据缓冲区，GPU1显存中分配目标缓冲区，传输方向固定为GPU0→GPU1。通过CUDA事件计时，分别记录NVLink路径（cudaMemcpyPeerAsync经NVLink）与PCIe路径（经PCIe Switch）的传输耗时。

代码实现

以下代码使用CUDA P2P API完成混合拓扑下的双路径传输基准测试。程序首先查询设备间P2P支持状态，然后分配显存缓冲区并填充源数据，分别通过NVLink（cudaMemcpyPeerAsync默认路径）与PCIe（通过cudaMemcpyPeerAsync配合cudaCtxResetPersistingL2Cache控制，或使用cudaMemcpyPeerAsync并指定非P2P访问方式）执行传输。为区分两条路径，NVLink路径使用设备间直接访问（需启用P2P），PCIe路径则通过禁用P2P后的staged传输（经主机内存中转）来模拟PCIe Switch路径。计时采用CUDA事件，带宽按传输字节数除以耗时计算。

// nvlink_pcie_bench.cu
// 编译: nvcc -o bench nvlink_pcie_bench.cu
// 运行: ./bench  (需双GPU且支持P2P)

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#include <chrono>

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

// 传输数据块大小 (字节)
const size_t DATA_SIZE = 256 * 1024 * 1024;  // 256MB

// 执行一次传输并返回耗时(ms)
float run_transfer(int src_dev, int dst_dev, void* src, void* dst, 
                   size_t bytes, bool use_p2p) {
    // 设置当前设备为源设备
    CHECK_CUDA(cudaSetDevice(src_dev));
    
    // 根据路径类型启用/禁用P2P
    if (use_p2p) {
        // NVLink路径: 启用P2P访问
        int can_access = 0;
        CHECK_CUDA(cudaDeviceCanAccessPeer(&can_access, src_dev, dst_dev));
        if (!can_access) {
            printf("错误: GPU%d 无法P2P访问 GPU%d\n", src_dev, dst_dev);
            exit(EXIT_FAILURE);
        }
        CHECK_CUDA(cudaDeviceEnablePeerAccess(dst_dev, 0));
    } else {
        // PCIe路径: 禁用P2P, 强制走staged拷贝(经主机内存)
        // 若之前已启用, 先禁用
        cudaDeviceDisablePeerAccess(dst_dev);  // 忽略错误, 可能未启用
    }
    
    // 创建CUDA事件用于计时
    cudaEvent_t start, stop;
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));
    
    // 预热 (避免首次调用开销)
    CHECK_CUDA(cudaMemcpyPeerAsync(dst, dst_dev, src, src_dev, bytes));
    CHECK_CUDA(cudaDeviceSynchronize());
    
    // 正式计时
    CHECK_CUDA(cudaEventRecord(start));
    CHECK_CUDA(cudaMemcpyPeerAsync(dst, dst_dev, src, src_dev, bytes));
    CHECK_CUDA(cudaEventRecord(stop));
    CHECK_CUDA(cudaEventSynchronize(stop));
    
    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));
    
    // 清理: 若启用了P2P则禁用
    if (use_p2p) {
        CHECK_CUDA(cudaDeviceDisablePeerAccess(dst_dev));
    }
    
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    
    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;
    }
    
    // 使用设备0和设备1
    const int dev0 = 0, dev1 = 1;
    
    // 查询设备属性
    cudaDeviceProp prop0, prop1;
    CHECK_CUDA(cudaGetDeviceProperties(&prop0, dev0));
    CHECK_CUDA(cudaGetDeviceProperties(&prop1, dev1));
    printf("GPU0: %s\n", prop0.name);
    printf("GPU1: %s\n", prop1.name);
    
    // 检查P2P支持
    int can_access = 0;
    CHECK_CUDA(cudaDeviceCanAccessPeer(&can_access, dev0, dev1));
    printf("P2P支持 (GPU0→GPU1): %s\n", can_access ? "是" : "否");
    if (!can_access) {
        printf("错误: 当前GPU组合不支持P2P, 无法测试NVLink路径\n");
        return EXIT_FAILURE;
    }
    
    // 分配显存
    void *src = nullptr, *dst = nullptr;
    CHECK_CUDA(cudaSetDevice(dev0));
    CHECK_CUDA(cudaMalloc(&src, DATA_SIZE));
    CHECK_CUDA(cudaSetDevice(dev1));
    CHECK_CUDA(cudaMalloc(&dst, DATA_SIZE));
    
    // 初始化源数据 (在GPU0上填充)
    CHECK_CUDA(cudaSetDevice(dev0));
    unsigned char* h_data = (unsigned char*)malloc(DATA_SIZE);
    for (size_t i = 0; i < DATA_SIZE; i++) {
        h_data[i] = (unsigned char)(i % 251);
    }
    CHECK_CUDA(cudaMemcpy(src, h_data, DATA_SIZE, cudaMemcpyHostToDevice));
    free(h_data);
    
    // 目标缓冲区清零
    CHECK_CUDA(cudaSetDevice(dev1));
    CHECK_CUDA(cudaMemset(dst, 0, DATA_SIZE));
    
    // 执行多次测试取平均
    const int ITERATIONS = 5;
    float nvlink_ms = 0.0f, pcie_ms = 0.0f;
    
    printf("\n开始基准测试 (数据大小: %.0f MB, 迭代次数: %d)\n", 
           DATA_SIZE / (1024.0*1024.0), ITERATIONS);
    
    for (int i = 0; i < ITERATIONS; i++) {
        // NVLink路径
        float t1 = run_transfer(dev0, dev1, src, dst, DATA_SIZE, true);
        nvlink_ms += t1;
        
        // PCIe路径 (禁用P2P)
        float t2 = run_transfer(dev0, dev1, src, dst, DATA_SIZE, false);
        pcie_ms += t2;
        
        printf("迭代 %d: NVLink=%.3f ms, PCIe=%.3f ms\n", i+1, t1, t2);
    }
    
    nvlink_ms /= ITERATIONS;
    pcie_ms /= ITERATIONS;
    
    // 计算带宽 (GB/s)
    double nvlink_bw = (double)DATA_SIZE / (nvlink_ms / 1000.0) / 1e9;
    double pcie_bw = (double)DATA_SIZE / (pcie_ms / 1000.0) / 1e9;
    
    printf("\n========== 测试结果 ==========\n");
    printf("NVLink路径: 平均耗时 %.3f ms, 带宽 %.2f GB/s\n", nvlink_ms, nvlink_bw);
    printf("PCIe路径:   平均耗时 %.3f ms, 带宽 %.2f GB/s\n", pcie_ms, pcie_bw);
    printf("--------------------------------\n");
    printf("NVLink/PCIe 带宽比: %.2fx\n", nvlink_bw / pcie_bw);
    printf("NVLink/PCIe 延迟比: %.2fx (越小越好)\n", nvlink_ms / pcie_ms);
    
    // 验证数据完整性
    CHECK_CUDA(cudaSetDevice(dev1));
    unsigned char* verify = (unsigned char*)malloc(DATA_SIZE);
    CHECK_CUDA(cudaMemcpy(verify, dst, DATA_SIZE, cudaMemcpyDeviceToHost));
    
    bool ok = true;
    for (size_t i = 0; i < DATA_SIZE; i++) {
        if (verify[i] != (unsigned char)(i % 251)) {
            printf("数据校验失败 @ offset %zu\n", i);
            ok = false;
            break;
        }
    }
    free(verify);
    
    if (ok) {
        printf("数据完整性校验: 通过\n");
    } else {
        printf("数据完整性校验: 失败\n");
    }
    
    // 清理
    CHECK_CUDA(cudaSetDevice(dev0));
    CHECK_CUDA(cudaFree(src));
    CHECK_CUDA(cudaSetDevice(dev1));
    CHECK_CUDA(cudaFree(dst));
    
    return ok ? EXIT_SUCCESS : EXIT_FAILURE;
}

编译与运行

在具备双GPU且支持NVLink的服务器上执行以下命令：

nvcc -o bench nvlink_pcie_bench.cu
./bench

验证方法与判定标准

程序运行后应输出GPU型号、P2P支持状态、每次迭代的耗时以及最终的平均带宽对比。判定标准如下：

• 功能正确性：程序正常结束且输出“数据完整性校验: 通过”，说明两条路径均成功完成数据传输且数据无损坏。
• 路径区分有效性：NVLink路径（启用P2P）与PCIe路径（禁用P2P）的耗时存在显著差异，NVLink路径耗时明显更短。
• 性能对比合理性：在支持NVLink的平台上（如A100、V100等），NVLink带宽应显著高于PCIe Gen4 x16的理论带宽（约32GB/s），实测NVLink带宽通常可达PCIe的2~5倍。若NVLink带宽低于PCIe，需检查拓扑连接或驱动配置。
• 可重复性：多次运行结果应保持稳定，带宽波动不超过±10%。

注意：实际性能数值取决于具体GPU型号、NVLink版本（NVLink2/3/4）、PCIe代际及系统拓扑，本案例不预设具体数值，以实际运行输出为准。