3. cudaMemcpy基础：三种传输模式详解
说到CUDA编程，内存拷贝是绕不开的基本功。我刚开始接触CUDA时，觉得不就是把数据从CPU搬到GPU嘛，有什么难的？结果第一次写大规模数据处理程序，就因为没搞清楚传输模式，性能直接崩了。今天咱们就把cudaMemcpy的三种模式彻底讲透。
核心要点：cudaMemcpy是CUDA中最基础的数据传输函数，支持三种方向：主机→设备、设备→主机、设备→设备。每种模式都有其适用场景和性能特征。
3.1 函数原型与参数
先看函数签名，这个我建议你背下来，因为太常用了：
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind);四个参数，说白了就是：
- dst：目标地址指针
- src：源地址指针
- count：拷贝的字节数
- kind：传输方向类型
嗯，这里要注意，count的单位是字节，不是元素个数。我见过不少新手在这里栽跟头——想拷贝100个float，结果传了100而不是400。你想想看，这能对吗？
3.2 三种传输模式详解
cudaMemcpyKind枚举定义了四种类型，但常用的就三种。我整理成表格，一目了然：
| 枚举值 | 方向 | 典型场景 |
|---|---|---|
cudaMemcpyHostToDevice |
主机 → 设备 | 将输入数据从CPU内存拷贝到GPU显存 |
cudaMemcpyDeviceToHost |
设备 → 主机 | 将计算结果从GPU显存拷贝回CPU内存 |
cudaMemcpyDeviceToDevice |
设备 → 设备 | GPU显存内部的数据搬移（如中间结果整理） |
小提示：还有第四种cudaMemcpyHostToHost，但说实话，我几乎没用过。CPU内存之间拷贝用memcpy就行了，何必绕CUDA？
3.3 主机到设备（HostToDevice）
这是最常用的模式。把数据从CPU内存传到GPU显存，让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);为什么这个模式重要？因为GPU不能直接访问CPU内存。你得先把数据搬到显存里，GPU才能干活。我在项目中遇到过一个问题：数据量太大，一次性拷贝导致程序卡死。后来改用分块拷贝，才解决了。
避坑指南：我曾经在调试一个图像处理程序时，发现结果全是错的。排查了半天，原来是cudaMemcpy的返回值没检查。记住：一定要检查返回值！用cudaGetLastError()或直接检查cudaError_t。
3.4 设备到主机（DeviceToHost）
GPU算完了，结果得拿回来。这就是设备到主机的拷贝：
float* h_result = (float*)malloc(N * sizeof(float));
// 设备到主机拷贝
cudaMemcpy(h_result, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);你想想看，这个操作是同步的。也就是说，cudaMemcpy返回时，数据已经安全地躺在CPU内存里了。但这也意味着它会阻塞CPU线程，直到拷贝完成。
我个人习惯在做完设备到主机拷贝后，立刻用cudaDeviceSynchronize()确保所有操作完成。虽然cudaMemcpy本身是同步的，但加一道保险总没错。
3.5 设备到设备（DeviceToDevice）
这个模式很多人容易忽略。其实在GPU内部，有时候也需要搬数据。比如你算完一个中间结果，想把它整理到另一个缓冲区：
float* d_temp1, * d_temp2;
cudaMalloc(&d_temp1, N * sizeof(float));
cudaMalloc(&d_temp2, N * sizeof(float));
// 在GPU内部拷贝
cudaMemcpy(d_temp2, d_temp1, N * sizeof(float), cudaMemcpyDeviceToDevice);设备到设备的拷贝速度非常快，因为它走的是GPU内部总线，不经过PCIe。我记得有一次做多GPU协同计算，需要在不同GPU之间搬数据，当时就用了cudaMemcpyPeer（设备到设备的进阶版）。
性能对比：设备到设备拷贝的带宽通常是主机到设备的10倍以上。所以，能留在GPU里的数据，尽量别来回倒腾。
3.6 知识体系图
下面这张图展示了cudaMemcpy三种模式在整个数据传输流程中的位置：
3.7 实战经验总结
最后，分享几个我这些年积累的经验：
- 批量拷贝优于单次拷贝：能一次拷完的数据，别分多次。每次cudaMemcpy都有开销。
- 异步拷贝是性能利器：配合
cudaMemcpyAsync和流(stream)，可以实现计算和传输重叠。这个后面章节会细讲。 - 注意内存对齐：虽然cudaMemcpy不要求严格对齐，但对齐的地址传输效率更高。
- 别在性能热点里做主机到设备拷贝：PCIe带宽有限，频繁传输会成为瓶颈。
我的习惯：每次写完cudaMemcpy，我都会加一行printf("cudaMemcpy status: %s\n", cudaGetErrorString(cudaGetLastError()));来调试。等程序稳定了再删掉。这招帮我省了不少排查时间。
好了，三种传输模式就讲到这里。记住：主机到设备是送数据进去，设备到主机是拿结果出来，设备到设备是内部整理。搞清楚了这三个方向，你的CUDA程序就成功了一半。