2. cudaMemcpy基础：同步拷贝的三种模式
好，咱们今天聊聊 cudaMemcpy。这个函数可以说是 CUDA 编程里最常用的 API 之一，没有它，CPU 和 GPU 之间就是「老死不相往来」的状态。
我个人习惯把 cudaMemcpy 叫做「同步拷贝三兄弟」——HostToDevice、DeviceToHost、DeviceToDevice。名字虽然长，但逻辑其实很简单：数据从哪来，到哪去。
2.1 API 签名与参数解析
先看函数签名，这是咱们吃饭的家伙：
cudaError_t cudaMemcpy(
void* dst, // 目标地址
const void* src, // 源地址
size_t count, // 拷贝字节数
cudaMemcpyKind kind // 拷贝方向
);四个参数，一个返回值。嗯，这里要注意：dst 和 src 的类型都是 void*，意味着你可以传任意类型的指针。但千万别传错了方向——我曾经在项目里把 dst 和 src 写反了，结果调试了整整一个下午才发现。
返回值 cudaError_t 也很关键。我建议每次调用后都检查一下返回值，或者用 cudaGetLastError() 兜底。你想想看，如果拷贝失败了，后续计算全白搭。
2.2 三种拷贝模式详解
2.2.1 HostToDevice（主机到设备）
这是最常用的模式。CPU 把数据准备好，然后「喂」给 GPU。说白了，就是 cudaMemcpy(d_gpu, h_cpu, size, cudaMemcpyHostToDevice)。
我在项目中遇到过一个问题：主机内存是 pageable（可分页）的，拷贝速度特别慢。后来改用 cudaHostAlloc 分配 page-locked 内存，速度直接翻倍。这个后面会细讲，先记住结论——page-locked 内存对 HostToDevice 拷贝有奇效。
关键点：HostToDevice 拷贝会阻塞 CPU 线程，直到数据传输完成。这意味着 CPU 在拷贝期间啥也干不了。
2.2.2 DeviceToHost（设备到主机）
GPU 算完了，结果得拿回 CPU 做后续处理。这就是 cudaMemcpy(h_cpu, d_gpu, size, cudaMemcpyDeviceToHost)。
这个方向有个坑：它也是阻塞的。而且因为 PCIe 带宽的限制，从 GPU 往 CPU 拷数据往往比反方向慢一些。我记得有一次做图像处理，每帧都要把结果拷回 CPU，结果帧率直接腰斩。后来改成异步拷贝才解决。
避坑指南：我曾经在 DeviceToHost 拷贝后直接访问主机内存，以为数据已经到位了。结果发现 GPU 内核还没执行完——因为 cudaMemcpy 虽然是同步的，但它只保证拷贝本身完成，不保证之前提交的内核已经执行完毕。所以记得在拷贝前加 cudaDeviceSynchronize()。
2.2.3 DeviceToDevice（设备到设备）
这个模式有意思：数据在 GPU 内部搬来搬去，不经过 PCIe。比如 cudaMemcpy(d_dst, d_src, size, cudaMemcpyDeviceToDevice)。
你可能会问：直接在 GPU 上操作指针不就行了，为啥还要拷贝？嗯，场景其实不少——比如做数据重排、多 GPU 间的数据交换（虽然这个场景更推荐用 cudaMemcpyPeer），或者某些库函数要求输入输出地址不同。
DeviceToDevice 拷贝的速度非常快，因为它走的是 GPU 内部的高速总线。我建议能用这个模式就别绕到 CPU 再拷回来，白白浪费带宽。
2.3 阻塞特性分析
为什么说 cudaMemcpy 是同步的？因为它会阻塞调用它的 CPU 线程，直到数据传输完成。说白了，就是 CPU 得「等」着。
来看个例子：
cudaMemcpy(d_gpu, h_cpu, size, cudaMemcpyHostToDevice);
// 这行代码执行时，CPU 被阻塞在这里
// 直到数据全部拷到 GPU 显存里，才会继续往下走
kernel<<<grid, block>>>(d_gpu);
// 内核启动是异步的，但数据拷贝已经完成了这种阻塞特性有好有坏：
- 好处：编程简单，逻辑清晰。你不需要操心数据什么时候到位，拷贝完就是到位了。
- 坏处：浪费 CPU 时间。如果数据量很大（比如几百 MB），CPU 就干等着，啥也干不了。
我刚开始学 CUDA 时，总觉得阻塞拷贝是「天经地义」的。后来发现，很多高性能场景下，CPU 和 GPU 应该「各干各的」——CPU 准备下一批数据，GPU 算当前这批。这时候就需要异步拷贝了，那是后面章节的内容。
2.4 知识体系图
下面这张图总结了 cudaMemcpy 的核心逻辑，我建议你多看几遍：
2.5 性能对比表
最后给个表格，方便你快速对比：
| 模式 | 方向 | 是否阻塞 CPU | 典型带宽 | 适用场景 |
|---|---|---|---|---|
| HostToDevice | CPU → GPU | 是 | ~12 GB/s (PCIe 3.0 x16) | 输入数据准备 |
| DeviceToHost | GPU → CPU | 是 | ~12 GB/s (PCIe 3.0 x16) | 结果回传 |
| DeviceToDevice | GPU 内部 | 是 | ~900 GB/s (HBM2e) | 显存内数据搬移 |
小提示：如果你发现 DeviceToDevice 拷贝很慢，先检查一下两个地址是不是在同一个 GPU 上。跨 GPU 拷贝要用 cudaMemcpyPeer，带宽会差很多。
好了，同步拷贝的基础就这些。记住一句话：同步拷贝简单可靠，但会浪费 CPU 时间。后面我们会讲怎么用异步拷贝来解决这个问题。