第三节：cudaMemcpy详解——同步与异步传输的实战心法
各位同学，今天我们来聊聊 CUDA 编程里最基础、也最容易踩坑的一个函数——cudaMemcpy。
说实话，我见过太多新手甚至老手，在这个函数上翻车。有人觉得它就是个简单的内存拷贝，有人搞不清同步异步的区别，还有人因为没注意传输方向导致性能暴跌。今天我就把这些坑一个个给你填平。
3.1 三种传输模式：HostToDevice、DeviceToHost、DeviceToDevice
先看最基本的。cudaMemcpy 支持三种传输方向，用枚举值控制：
cudaMemcpyHostToDevice // 主机 → 设备
cudaMemcpyDeviceToHost // 设备 → 主机
cudaMemcpyDeviceToDevice // 设备 → 设备你可能会问：不就是拷贝数据吗，分这么细干嘛？
嗯，这里有个关键点——这三种传输走的硬件路径完全不同。
| 传输方向 | 硬件路径 | 带宽（典型值） | 延迟 |
|---|---|---|---|
| Host → Device | CPU内存 → PCIe → GPU显存 | 12~32 GB/s | 微秒级 |
| Device → Host | GPU显存 → PCIe → CPU内存 | 12~32 GB/s | 微秒级 |
| Device → Device | GPU显存内部DMA | 300~900 GB/s | 纳秒级 |
看到没？DeviceToDevice 的带宽是前两者的几十倍。我在项目中遇到过有人用 HostToDevice 做中间跳板来拷贝显存数据，结果性能直接掉了两个数量级。你想想看，本来几微秒能搞定的事，硬生生拖到几百微秒。
核心原则：能走 DeviceToDevice 就别走 HostToDevice + DeviceToHost。GPU 内部拷贝比过 PCIe 快得多。
3.2 同步传输：cudaMemcpy 的默认行为
默认情况下，cudaMemcpy 是同步的。什么意思？
就是 CPU 调用 cudaMemcpy 之后，会一直等在那里，直到数据拷贝完成才返回。说白了，CPU 被阻塞了。
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
// CPU 停在这里，直到数据全部拷完
printf("拷贝完成\n"); // 这行代码要等上面执行完才跑这种模式的好处是简单——你不需要操心同步问题。坏处也很明显：CPU 在等的时候啥也干不了，浪费了计算资源。
我个人习惯是：小数据量（几MB以内）用同步传输，因为开销不大，代码也清晰。但大数据量，尤其是几百MB甚至GB级别的传输，同步模式就是性能杀手。
3.3 异步传输：cudaMemcpyAsync 的正确打开方式
异步传输才是真正发挥 GPU 性能的关键。它允许 CPU 在数据传输期间继续做其他事。
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream);
// CPU 立即返回，不等待拷贝完成
do_some_cpu_work(); // 可以和传输并行执行
cudaStreamSynchronize(stream); // 需要时再同步这里有个坑，我必须要提醒你：异步传输要求主机内存必须是页锁定内存（pinned memory）。
警告：如果你用普通 malloc 分配的主机内存传给 cudaMemcpyAsync，它会静默地退化为同步传输。你以为是异步，实际上 CPU 还是被阻塞了。我曾经因为这个 bug 排查了整整两天。
正确的做法是用 cudaHostAlloc 分配：
float *h_a;
cudaHostAlloc(&h_a, size, cudaHostAllocDefault);
// 现在 h_a 是页锁定内存，可以用于异步传输3.4 流（Stream）与异步传输的配合
异步传输必须绑定到一个流上。流可以理解为 GPU 上的任务队列。
为什么需要流？因为你可以创建多个流，让数据传输和内核计算重叠执行：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在 stream1 中传输数据
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream1);
// 在 stream2 中执行内核
kernel<<<grid, block, 0, stream2>>>(d_b);
// 两个操作可以同时进行
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);你想想看，如果不用流，你得先等数据拷完，再跑内核。用了流，传输和计算就能并行。对于需要频繁传输数据的应用，这个优化能带来 30%~50% 的性能提升。
3.5 实战避坑指南
最后，我把自己这些年踩过的坑总结一下：
- 别用 cudaMemcpy 做 DeviceToDevice 的中间拷贝——直接传指针就行，或者用 cudaMemcpyDeviceToDevice
- 异步传输一定要用 pinned memory——否则它偷偷变同步，你还不知道
- 大块传输用异步，小块传输用同步——异步有额外开销，太小不值得
- 记得检查返回值——cudaMemcpy 返回 cudaError_t，别忽略它
小技巧：我习惯在调试阶段用 cudaMemcpy 的同步版本，因为出错时堆栈信息更清晰。上线前再改成异步版本做性能优化。
3.6 本章知识体系图
下面这张图帮你理清 cudaMemcpy 的核心脉络：
好了，cudaMemcpy 的核心内容就这些。记住一句话：方向选对，同步异步看场景，pinned memory 别忘掉。下一节我们聊聊 Unified Memory，那个又是另一番天地了。
想要更多资料，联系可可