2、cudaMemcpy基础：同步拷贝的三种模式

说到 CUDA 编程，内存拷贝是绕不开的基本功。我刚开始接触 CUDA 时，觉得这玩意儿不就是把数据从 CPU 搬到 GPU 嘛，有啥好研究的？结果第一次跑大模型训练，发现光拷贝就占了将近一半的时间……嗯，从那以后我再也不敢小看这个 API 了。

今天咱们就聊聊 cudaMemcpy 的三种同步拷贝模式。说白了，就是数据在三种不同内存区域之间的搬运方式。

2.1 三种拷贝方向

cudaMemcpy 函数长这样：

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind);

最后一个参数 kind 决定了数据往哪搬。有三种选择：

• cudaMemcpyHostToDevice：从 CPU 内存搬到 GPU 显存
• cudaMemcpyDeviceToHost：从 GPU 显存搬回 CPU 内存
• cudaMemcpyDeviceToDevice：GPU 显存内部搬运

你可能会问：怎么没有 HostToHost？其实 CPU 内存之间拷贝用标准的 memcpy 就行，CUDA 不掺和这事儿。

2.2 HostToDevice：把数据喂给 GPU

这是最常用的模式。你在 CPU 上准备好了数据，想交给 GPU 去算，就得走这条路。

我个人习惯这样写：

float* h_data = (float*)malloc(N * sizeof(float));
float* d_data;
cudaMalloc(&d_data, N * sizeof(float));

// 初始化 h_data ...

cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

这里有个坑——内存对齐。我在项目中遇到过，有一次数据量特别大，我图省事用了 malloc 分配主机内存，结果拷贝速度慢得离谱。后来换成 cudaMallocHost 分配页锁定内存，速度直接翻了三倍。

2.3 DeviceToHost：把结果拿回来

GPU 算完了，你得把结果拿回 CPU 才能看。这就是 DeviceToHost 的活儿。

cudaMemcpy(h_result, d_result, N * sizeof(float), cudaMemcpyDeviceToHost);

你想想看，这个方向其实比 HostToDevice 更慢。为什么？因为 PCIe 总线是半双工的，虽然理论上双向带宽一样，但实际测试中 DeviceToHost 往往比 HostToDevice 慢 10%-20%。

我曾经踩过一个坑：在拷贝之前忘了调用 cudaDeviceSynchronize()。虽然 cudaMemcpy 本身是同步的，但它只保证拷贝完成，不保证前面的 kernel 执行完。结果我拿到的数据是上一轮计算的结果……排查了半天才发现。

2.4 DeviceToDevice：显存内部搬运

这个模式很多人用得少，但其实很实用。比如你要在 GPU 上做数据重排、合并多个缓冲区，或者实现一些复杂的访存模式。

cudaMemcpy(d_dst, d_src, N * sizeof(float), cudaMemcpyDeviceToDevice);

有意思的是，DeviceToDevice 拷贝不经过 PCIe 总线，完全在 GPU 内部完成。所以它的速度比前两种快得多——大概在 200-400 GB/s 的量级（取决于你的 GPU 型号）。

不过要注意：它走的是 GPU 的全局内存带宽，如果你频繁做这种拷贝，会影响 kernel 的执行。我建议只在初始化阶段或者数据整理时用，别在计算密集的循环里搞这个。

2.5 三种模式的性能对比

我整理了一张表，方便你直观对比：

拷贝模式	典型带宽	延迟	是否经过 PCIe	常见瓶颈
HostToDevice	10-12 GB/s	高（微秒级）	是	PCIe 带宽、主机内存类型
DeviceToHost	8-10 GB/s	高（微秒级）	是	PCIe 带宽、主机内存类型
DeviceToDevice	200-400 GB/s	低（纳秒级）	否	显存带宽、内存冲突

（以上数据基于 PCIe 3.0 x16 和 A100 级别 GPU，实际数值因硬件而异）

2.6 一张图看懂三种模式

下面这张 SVG 图展示了三种拷贝模式的完整数据流：

2.7 避坑指南

最后分享几个我踩过的坑，希望能帮你省点时间：

• 别用 malloc 分配大块数据：默认的 malloc 分配的是可分页内存，DMA 传输时需要先锁定页面，额外开销很大。用 cudaMallocHost 或 cudaHostAlloc 会好很多。
• 注意拷贝粒度：单次拷贝太小（比如几十个字节）会浪费 PCIe 的包传输开销。我一般建议单次拷贝至少 1MB 以上，否则不如用 kernel 自己搬。
• 别在 kernel 执行时做 D2D 拷贝：虽然 CUDA 支持并发，但 D2D 拷贝会占用显存带宽，影响 kernel 性能。如果非要同时做，考虑用流（stream）来隔离。
• 检查返回值：cudaMemcpy 返回 cudaError_t，别偷懒不检查。我曾经因为指针传错了，查了两天才发现是拷贝越界。

好了，三种同步拷贝模式就聊到这儿。记住一句话：数据怎么搬，决定了你的程序能跑多快。下一节咱们聊聊异步拷贝，那才是真正拉开性能差距的地方。