4. 多GPU通信延迟测量与优化

任务说明

本案例目标是在多GPU环境中测量点对点通信延迟，识别影响延迟的关键因素，并通过调整传输数据块大小与使用CUDA事件计时来获得可复现的延迟基准。任务要求：编写一个可独立运行的CUDA程序，使用NVLink连接的两个GPU（或同一GPU上的两个上下文）完成一次固定大小的数据往返传输，测量端到端延迟，并对比不同数据块大小下的延迟变化趋势。

实现方案

采用CUDA的cudaMemcpyPeerAsync实现设备间直接拷贝，配合CUDA事件（cudaEvent）记录时间戳。为避免启动开销掩盖真实传输延迟，采用多次迭代取中位数的方式。程序接受两个参数：数据块大小（字节）和迭代次数，输出每次往返的平均延迟与标准差。

代码实现

// latency_measure.cu
// 编译: nvcc -arch=sm_70 latency_measure.cu -o latency_measure
// 运行: ./latency_measure [bytes] [iterations]
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#include <algorithm>
#include <vector>

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

// 返回单次单向延迟（微秒）
float measure_one_way_latency(void* send0, void* recv0,
                              void* send1, void* recv1,
                              size_t bytes, int dev0, int dev1) {
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    // 确保所有前置操作完成
    CUDA_CHECK(cudaDeviceSynchronize());

    CUDA_CHECK(cudaEventRecord(start, dev0));
    // 双向拷贝：GPU0→GPU1 和 GPU1→GPU0
    CUDA_CHECK(cudaMemcpyPeerAsync(recv1, dev1, send0, dev0, bytes));
    CUDA_CHECK(cudaMemcpyPeerAsync(recv0, dev0, send1, dev1, bytes));
    CUDA_CHECK(cudaEventRecord(stop, dev0));
    CUDA_CHECK(cudaEventSynchronize(stop));

    float ms = 0.0f;
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));

    // 双向总耗时 / 2 = 单次单向延迟（毫秒→微秒）
    return (ms * 1000.0f) / 2.0f;
}

int main(int argc, char** argv) {
    size_t bytes = 1024;          // 默认1KB
    int iterations = 100;         // 默认100次

    if (argc >= 2) bytes = atoll(argv[1]);
    if (argc >= 3) iterations = atoi(argv[2]);
    if (bytes == 0 || iterations <= 0) {
        fprintf(stderr, "Usage: %s [bytes] [iterations]\n", argv[0]);
        return 1;
    }

    int dev0 = 0, dev1 = 1;
    CUDA_CHECK(cudaSetDevice(dev0));
    cudaDeviceProp prop0, prop1;
    CUDA_CHECK(cudaGetDeviceProperties(&prop0, dev0));
    CUDA_CHECK(cudaGetDeviceProperties(&prop1, dev1));

    // 检查设备间是否支持点对点访问（NVLink或PCIe）
    int canAccess = 0;
    CUDA_CHECK(cudaDeviceCanAccessPeer(&canAccess, dev0, dev1));
    if (!canAccess) {
        printf("Warning: P2P not supported between GPU%d and GPU%d, "
               "fallback to staged copy\n", dev0, dev1);
    } else {
        CUDA_CHECK(cudaSetDevice(dev0));
        CUDA_CHECK(cudaDeviceEnablePeerAccess(dev1, 0));
        CUDA_CHECK(cudaSetDevice(dev1));
        CUDA_CHECK(cudaDeviceEnablePeerAccess(dev0, 0));
    }

    // 分配缓冲区
    void *send0, *recv0, *send1, *recv1;
    CUDA_CHECK(cudaSetDevice(dev0));
    CUDA_CHECK(cudaMalloc(&send0, bytes));
    CUDA_CHECK(cudaMalloc(&recv0, bytes));
    CUDA_CHECK(cudaMemset(send0, 0xAB, bytes));
    CUDA_CHECK(cudaMemset(recv0, 0, bytes));

    CUDA_CHECK(cudaSetDevice(dev1));
    CUDA_CHECK(cudaMalloc(&send1, bytes));
    CUDA_CHECK(cudaMalloc(&recv1, bytes));
    CUDA_CHECK(cudaMemset(send1, 0xCD, bytes));
    CUDA_CHECK(cudaMemset(recv1, 0, bytes));

    // 预热
    for (int i = 0; i < 10; i++) {
        measure_one_way_latency(send0, recv0, send1, recv1, bytes, dev0, dev1);
    }

    // 正式测量
    std::vector<float> samples;
    samples.reserve(iterations);
    for (int i = 0; i < iterations; i++) {
        float us = measure_one_way_latency(send0, recv0, send1, recv1,
                                           bytes, dev0, dev1);
        samples.push_back(us);
    }

    // 统计：排序取中位数，计算P90
    std::sort(samples.begin(), samples.end());
    float median = samples[iterations / 2];
    float p90 = samples[(int)(iterations * 0.9)];
    float min_val = samples.front();
    float max_val = samples.back();

    // 输出结果
    printf("=== NVLink Latency Measurement ===\n");
    printf("GPU0: %s\nGPU1: %s\n", prop0.name, prop1.name);
    printf("P2P supported: %s\n", canAccess ? "yes" : "no");
    printf("Data size: %zu bytes (%.2f KB)\n", bytes, bytes / 1024.0);
    printf("Iterations: %d\n", iterations);
    printf("------------------------------------\n");
    printf("One-way latency (us):\n");
    printf("  median : %.2f\n", median);
    printf("  min    : %.2f\n", min_val);
    printf("  max    : %.2f\n", max_val);
    printf("  P90    : %.2f\n", p90);

    // 清理
    CUDA_CHECK(cudaSetDevice(dev0));
    CUDA_CHECK(cudaFree(send0));
    CUDA_CHECK(cudaFree(recv0));
    CUDA_CHECK(cudaSetDevice(dev1));
    CUDA_CHECK(cudaFree(send1));
    CUDA_CHECK(cudaFree(recv1));

    if (canAccess) {
        CUDA_CHECK(cudaSetDevice(dev0));
        CUDA_CHECK(cudaDeviceDisablePeerAccess(dev1));
        CUDA_CHECK(cudaSetDevice(dev1));
        CUDA_CHECK(cudaDeviceDisablePeerAccess(dev0));
    }

    return 0;
}

编译与运行

# 编译（需CUDA Toolkit，建议V10.0+）
nvcc -arch=sm_70 latency_measure.cu -o latency_measure

# 运行示例：测量1KB数据块延迟，迭代200次
./latency_measure 1024 200

# 测量不同大小数据块（字节），观察延迟变化
for size in 64 256 1024 4096 16384 65536; do
    ./latency_measure $size 200 | grep "median"
done

验证方法与判定标准

功能验证：程序应输出GPU型号、P2P支持状态、数据块大小和延迟统计值。若P2P不可用，程序会打印警告并自动回退到staged拷贝路径（此时延迟会显著升高，属正常现象）。

正确性判定：

程序能正常编译且无CUDA运行时错误；
输出中median值在合理范围：NVLink环境下，小数据块（≤1KB）单向延迟通常在1～5微秒；若超过20微秒，说明可能未走NVLink或系统负载过高；
连续运行两次相同参数，median值波动应小于15%，证明测量可重复；
增大数据块大小时，延迟应呈上升趋势（因传输时间增加），但小数据块（≤256B）延迟应基本持平（此时延迟主要由固定开销主导）。

优化对比：分别运行./latency_measure 64 200和./latency_measure 65536 200，记录两者的median值。若小数据块延迟明显高于大数据块延迟的1/10，说明固定开销（如API调用、事件同步）占主导，可考虑使用cudaMemcpyPeerAsync配合多流（cudaStream_t）重叠通信与计算来隐藏延迟。

📚 课程目录