CUDA批量内存复制性能对比：GFD与cudaMemcpyAsync深度评测

2026-05-27阅读 0热度 0

其他

在上一篇关于《Agent写的高性能Host-to-Device传输库》的文章发布后，有读者在评论区指出，CUDA 12.8引入了一个名为cudaMemcpyBatchAsync的新API。本文将通过详尽的基准测试，深入剖析其性能表现。所有测试代码均已同步更新至项目仓库。

TL;DR

cudaMemcpyBatchAsync的核心价值在于将多次独立的内存拷贝调用合并为单次API提交，从而显著削减了用户态到内核态的切换开销。然而，在处理海量、非连续的细粒度数据拷贝时，其底层实现机制——为每个条目顺序生成DMA命令——构成了性能瓶颈。与经过专门优化的传输方案相比，其吞吐量存在差距。在多GPU并行场景下，驱动锁的竞争问题会进一步放大这一劣势。

1. 什么是cudaMemcpyBatchAsync

作为CUDA 12.8版本的新特性，cudaMemcpyBatchAsync是一个专为批量异步内存拷贝设计的API。其解决的问题非常具体：传统上，将N个独立的CPU内存块传输至GPU需要调用N次cudaMemcpyAsync，每次调用约产生1-2微秒的固定开销。在大语言模型推理等需要处理数千个KV-cache块的场景中，这种累积开销将成为主要性能制约因素。

该API的解决方案是批量提交：将N个拷贝请求的参数数组打包，通过一次系统调用提交给GPU驱动，从而将N-1次的上下文切换开销降至最低。

// API 签名
cudaError_t cudaMemcpyBatchAsync(
    void* const* dstArray,          // N 个目标地址数组
    const void* const* srcArray,    // N 个源地址数组
    size_t* sizeArray,             // N 个拷贝大小数组
    size_t count,                  // 条目数 N
    cudaMemcpyAttributes* attrArray, // 属性数组(描述内存类型)
    size_t* attrIdxArray,          // 每条目对应的属性索引
    size_t numAttrs,               // 属性数量
    cudaStream_t stream            // CUDA 流
);

// cudaMemcpyAttributes 结构
struct cudaMemcpyAttributes {
    cudaMemcpySrcAccessOrder srcAccessOrder; // 源内存访问顺序
    unsigned int flags;                       // 保留标志(设为 0)
};

关键参数解析：

其内部工作机制如下：

核心优势： 将N次API调用合并为1次，消除了绝大部分的调用开销。
局限性： 驱动内部仍然需要为每个条目生成独立的Copy Engine (CE) DMA命令，并且这些命令是顺序执行的。这正是其性能潜力受限的根本原因。

2. 测试结果

我们通过基准测试获取了第一手性能数据。

测试环境

GPU： NVIDIA RTX PRO 5000 72GB (Blackwell, sm_120)
PCIe： Gen5 x16 (实测带宽约 53 GB/s)
CPU： 256核，2个NUMA节点
GFD配置： 15个gather工作线程，3个CE通道，5倍大页staging缓冲区
数据布局： Token以2倍步长分散在锁页CPU内存中
迭代： 每个配置运行50次，前15次作为预热

测试代码位于examples/04_benchmark.cu，运行./gfd_benchmark即可。

对比图表揭示了清晰的性能趋势：

在拷贝块数较少时，cudaMemcpyBatchAsync凭借其降低API调用开销的优势，表现尚可。但随着拷贝块数（N）的增加，其性能迅速达到瓶颈并趋于平缓。相比之下，GFD方案则能随着N的增加，持续利用更高的PCIe带宽，展现出优异的可扩展性。

性能差异源于底层实现路径的不同：

cudaMemcpyBatchAsync本质上仍为每个离散内存块生成独立的CE命令，并由硬件顺序执行。而GFD方案通过Gather线程先将离散的CPU数据汇聚到连续的大页缓冲区，再由CE执行高效的DMA传输，从而规避了离散拷贝的固有瓶颈。

3. 多GPU并行测试

单卡测试揭示了基础性能，多GPU并行场景则构成了真正的压力测试。我们设置每张GPU传输2048个4KB的块（总计8MB，以2倍步长分散）。测试代码见examples/05_multi_gpu_benchmark.cu，运行./gfd_multi_gpu_benchmark。

结果非常明确：cudaMemcpyBatchAsync在多卡场景下出现了严重的性能退化，扩展性几乎为零。而GFD方案则保持了近乎线性的性能提升。

性能瓶颈的根源在于驱动锁竞争。当8个线程（对应8张GPU）并发调用cudaMemcpyBatchAsync时，它们会竞争同一把全局驱动锁。每个线程持锁约717微秒以构建其2048个CE命令。平均而言，每张GPU需要等待大约3.5张其他GPU释放锁，导致总等待时间高达3.5 × 717 ≈ 2510微秒，这直接拖垮了整体吞吐量。

GFD方案实现线性扩展的关键在于其去中心化的架构设计，有效避免了锁竞争：

GFD为每个GPU实例分配了独立的Gather线程和CE通道资源。每个实例独立完成数据的汇聚与传输，实例间无资源争用，从而实现了近乎完美的并行扩展。

附录 A. cudaMemcpyBatchAsync 离散 H2D 批量拷贝示例代码

以下提供一个完整的代码示例，演示如何使用cudaMemcpyBatchAsync进行离散的主机到设备内存批量拷贝，便于读者理解和复现测试：

#include 
#include 
#include 

int main() {
    const int NUM_TOKENS = 2048;
    const size_t TOKEN_SIZE = 4096; // 4KB per token

    // 1. 分配 GPU 连续目标缓冲区
    char* gpu_buf;
    cudaMalloc(&gpu_buf, NUM_TOKENS * TOKEN_SIZE);

    // 2. 分配 CPU 锁页源缓冲区(模拟离散 KV-cache)
    char* cpu_buf;
    cudaMallocHost(&cpu_buf, NUM_TOKENS * TOKEN_SIZE * 2); // 2x stride 模拟离散
    // 填充测试数据
    for (int i = 0; i < NUM_TOKENS; i++) {
        memset(cpu_buf + i * TOKEN_SIZE * 2, i & 0xFF, TOKEN_SIZE);
    }

    // 3. 构建批量拷贝参数数组
    std::vector dsts(NUM_TOKENS);
    std::vector srcs(NUM_TOKENS);
    std::vector sizes(NUM_TOKENS);
    for (int i = 0; i < NUM_TOKENS; i++) {
        dsts[i] = gpu_buf + i * TOKEN_SIZE;          // GPU 连续排列
        srcs[i] = cpu_buf + i * TOKEN_SIZE * 2;      // CPU 离散(2x stride)
        sizes[i] = TOKEN_SIZE;
    }

    // 4. 设置内存属性
    cudaMemcpyAttributes attr = {};
    attr.srcAccessOrder = cudaMemcpySrcAccessOrderStream; // 源为锁页内存
    attr.flags = 0;
    // 所有条目共享同一个属性(索引 0)
    std::vector attrIdxs(NUM_TOKENS, 0);

    // 5. 创建流并执行批量拷贝
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    cudaMemcpyBatchAsync(
        (void* const*)dsts.data(),
        (const void* const*)srcs.data(),
        sizes.data(),
        NUM_TOKENS,
        &attr,
        attrIdxs.data(),
        1, // numAttrs = 1(只有一种属性)
        stream
    );
    cudaStreamSynchronize(stream);

    // 6. 验证
    std::vector verify(TOKEN_SIZE);
    cudaMemcpy(verify.data(), gpu_buf, TOKEN_SIZE, cudaMemcpyDeviceToHost);
    printf("第一个 token 首字节: 0x%02x (期望 0x00)\n", (unsigned char)verify[0]);
    cudaMemcpy(verify.data(), gpu_buf + 100 * TOKEN_SIZE, TOKEN_SIZE, cudaMemcpyDeviceToHost);
    printf("第 100 个 token 首字节: 0x%02x (期望 0x64)\n", (unsigned char)verify[0]);

    // 7. 清理
    cudaStreamDestroy(stream);
    cudaFree(gpu_buf);
    cudaFreeHost(cpu_buf);
    printf("批量拷贝完成: %d 个 token, 每个 %zu 字节\n", NUM_TOKENS, TOKEN_SIZE);
    return 0;
}