4. NVLink延迟测量

任务目标

编写一个运行在DPU（BlueField）上的程序，通过NVLink与GPU通信，测量小消息（例如64字节）的往返延迟（RTT），并统计平均、最小、最大延迟。本案例聚焦于延迟测量本身，不涉及带宽优化。

环境与依赖

  • 硬件：NVIDIA BlueField-2/3 DPU + 配套GPU（如A100/H100），NVLink已连接
  • 软件：DOCA 2.x（含doca_comm库）、CUDA 12.x、GCC
  • 编译依赖：libdoca-comm-dev、libdoca-common-dev

实现原理

使用DOCA Comm的doca_comm_rdma_write实现GPU到DPU的单向写操作。为测量往返延迟，DPU收到GPU数据后立即通过doca_comm_rdma_write将相同数据写回GPU。GPU侧记录发送前和接收后的时间戳，差值即为一次往返延迟。重复N次后统计平均、最小、最大延迟。

关键完整代码

// latency_measure.c - DPU侧程序
#include <doca_comm.h>
#include <doca_log.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>

#define MSG_SIZE 64
#define ITERATIONS 1000

DOCA_LOG_REGISTER(LATENCY_MEASURE);

static double now_us(void) {
    struct timespec ts;
    clock_gettime(CLOCK_MONOTONIC, &ts);
    return ts.tv_sec * 1e6 + ts.tv_nsec / 1e3;
}

int main(int argc, char *argv[]) {
    if (argc != 3) {
        printf("Usage: %s <pcie_addr> <gpu_comm_addr>\n", argv[0]);
        return 1;
    }

    struct doca_comm *comm;
    struct doca_comm_peer *gpu_peer;
    struct doca_comm_rdma *rdma;
    struct doca_comm_rdma_peer *rdma_peer;
    char *local_buf, *remote_buf;
    doca_error_t result;

    // 初始化DOCA Comm
    result = doca_comm_init(&comm);
    if (result != DOCA_SUCCESS) {
        DOCA_LOG_ERR("Failed to init comm: %s", doca_error_get_descr(result));
        return 1;
    }

    // 创建本地端点和连接GPU
    result = doca_comm_peer_create(comm, argv[1], &gpu_peer);
    if (result != DOCA_SUCCESS) {
        DOCA_LOG_ERR("Failed to create peer: %s", doca_error_get_descr(result));
        return 1;
    }

    result = doca_comm_rdma_create(comm, &rdma);
    if (result != DOCA_SUCCESS) {
        DOCA_LOG_ERR("Failed to create rdma: %s", doca_error_get_descr(result));
        return 1;
    }

    // 注册内存
    local_buf = aligned_alloc(4096, MSG_SIZE);
    memset(local_buf, 0, MSG_SIZE);
    result = doca_comm_rdma_register_memory(rdma, local_buf, MSG_SIZE);
    if (result != DOCA_SUCCESS) {
        DOCA_LOG_ERR("Failed to register memory: %s", doca_error_get_descr(result));
        return 1;
    }

    // 连接GPU侧RDMA对端
    result = doca_comm_rdma_peer_create(rdma, gpu_peer, argv[2], &rdma_peer);
    if (result != DOCA_SUCCESS) {
        DOCA_LOG_ERR("Failed to create rdma peer: %s", doca_error_get_descr(result));
        return 1;
    }

    // 获取GPU侧远程缓冲区地址（由GPU程序通过参数传递）
    remote_buf = (char *)strtoull(argv[2], NULL, 16);

    // 等待GPU程序就绪
    printf("DPU ready, waiting for GPU...\n");
    fflush(stdout);
    char sync = 0;
    while (sync != 'G') {
        // 简单轮询同步：GPU会写一个字节到本地buf
        sync = local_buf[0];
    }

    // 延迟测量循环
    double min_lat = 1e9, max_lat = 0, sum_lat = 0;
    for (int i = 0; i < ITERATIONS; i++) {
        // 发送数据到GPU（触发一次往返）
        local_buf[0] = (char)i;
        result = doca_comm_rdma_write(rdma, rdma_peer, remote_buf, local_buf, MSG_SIZE);
        if (result != DOCA_SUCCESS) {
            DOCA_LOG_ERR("Write failed at iter %d: %s", i, doca_error_get_descr(result));
            return 1;
        }

        // 等待GPU写回（轮询本地buf变化）
        while (local_buf[0] != (char)i) {
            // busy wait
        }
    }

    // GPU侧会计算延迟并打印，DPU侧仅同步
    printf("DPU measurement complete. GPU will report statistics.\n");

    // 清理
    doca_comm_rdma_peer_destroy(rdma_peer);
    doca_comm_rdma_destroy(rdma);
    doca_comm_peer_destroy(gpu_peer);
    doca_comm_destroy(comm);
    free(local_buf);
    return 0;
}

GPU侧配套代码（延迟统计核心）

// gpu_latency.cu - GPU侧程序（简化核心逻辑）
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

#define MSG_SIZE 64
#define ITERATIONS 1000

__global__ void latency_kernel(char *buf, int iterations, double *results) {
    int tid = threadIdx.x;
    if (tid != 0) return;

    for (int i = 0; i < iterations; i++) {
        // 等待DPU写入
        while (buf[0] != (char)i) {}

        // 记录接收时间
        long long t1 = clock64();

        // 立即写回（模拟回显）
        buf[0] = (char)i;

        // 记录发送完成时间（实际延迟由DPU侧计算）
        long long t2 = clock64();
        results[i] = (double)(t2 - t1) / 1.5; // 转换为纳秒（假设GPU时钟1.5GHz）
    }
}

int main(int argc, char *argv[]) {
    // 初始化CUDA、分配共享内存、建立DOCA Comm连接（省略样板代码）
    // 关键：将共享内存地址作为remote_buf传给DPU程序

    double *d_results;
    cudaMalloc(&d_results, ITERATIONS * sizeof(double));

    // 启动内核
    latency_kernel<<<1, 1>>>(d_buf, ITERATIONS, d_results);
    cudaDeviceSynchronize();

    // 统计结果
    double *h_results = (double *)malloc(ITERATIONS * sizeof(double));
    cudaMemcpy(h_results, d_results, ITERATIONS * sizeof(double), cudaMemcpyDeviceToHost);

    double min_lat = 1e9, max_lat = 0, sum = 0;
    for (int i = 0; i < ITERATIONS; i++) {
        if (h_results[i] < min_lat) min_lat = h_results[i];
        if (h_results[i] > max_lat) max_lat = h_results[i];
        sum += h_results[i];
    }
    printf("Average latency: %.2f ns\n", sum / ITERATIONS);
    printf("Min latency: %.2f ns\n", min_lat);
    printf("Max latency: %.2f ns\n", max_lat);

    return 0;
}

编译与运行

  1. DPU侧编译： 
    gcc -o latency_measure latency_measure.c -ldoca_comm -ldoca_common
  2. GPU侧编译： 
    nvcc -o gpu_latency gpu_latency.cu -ldoca_comm -ldoca_common
  3. 运行（先启动GPU程序，再启动DPU程序）： 
    # 在GPU主机上
./gpu_latency <gpu_pcie_addr> <dpu_comm_addr> <shared_mem_addr>

# 在DPU上
./latency_measure <dpu_pcie_addr> <gpu_comm_addr>

验证方法

  • 程序正常退出且无错误日志
  • GPU侧输出平均、最小、最大延迟，数值在合理范围（NVLink小消息延迟通常为1-5微秒）
  • 最小延迟 ≤ 平均延迟 ≤ 最大延迟
  • 重复运行结果稳定（波动不超过20%）

常见问题

  • 连接失败：确认NVLink已启用（nvidia-smi nvlink -s），PCIe地址格式正确
  • 内存注册失败：确保使用aligned_alloc且大小是4KB对齐
  • 同步超时：检查GPU与DPU程序启动顺序，确保共享内存地址正确传递
← 上一章 📖 返回目录 下一章 →