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

说到CUDA编程，内存拷贝是绕不开的基本功。我个人习惯把cudaMemcpy看作是CPU和GPU之间的「桥梁」——数据要过去，结果要回来，全靠这座桥。今天我们就来聊聊这座桥的三种通行模式。

2.1 同步拷贝的本质

先搞清楚一个概念：什么叫「同步拷贝」？

说白了，就是调用cudaMemcpy的时候，CPU会乖乖等着，直到数据全部搬完才继续往下走。你想想看，这就像你寄快递，站在快递站等包裹打包好才离开——虽然有点浪费时间，但胜在简单可靠。

我在项目中遇到过不少新手，一上来就追求异步拷贝，结果数据还没传完就开始计算，跑出来的结果全是错的。嗯，这里要提醒一句：先学会走，再学跑。同步拷贝虽然慢，但它是理解CUDA内存模型的基础。

核心要点：同步拷贝是阻塞式的，函数返回时数据一定已经到位。

2.2 三种传输方向

cudaMemcpy有三个方向参数，我刚开始学的时候也记混过。其实很简单，看名字就知道：

枚举值	方向	典型场景
cudaMemcpyHostToDevice	CPU → GPU	把输入数据传到显存
cudaMemcpyDeviceToHost	GPU → CPU	把计算结果取回来
cudaMemcpyDeviceToDevice	GPU → GPU	显存内部数据搬移

2.3 HostToDevice：把数据送进显存

这是最常用的模式。你在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);

这里有个坑，我曾经踩过：源地址和目的地址的顺序。cudaMemcpy的参数顺序是(dst, src, size, kind)，跟memcpy一样——先目的，后源。我见过有人写反了，结果数据不知道传到哪里去了。

注意：HostToDevice的带宽通常受限于PCIe总线。我实测过，PCIe 3.0 x16的理论带宽约16GB/s，实际能跑到12GB/s就算不错了。如果你的数据量很大，这个拷贝时间会成为瓶颈。

2.4 DeviceToHost：把结果拿回来

GPU算完了，你得把结果取回内存。代码几乎是对称的：

float *h_result = (float*)malloc(N * sizeof(float));
cudaMemcpy(h_result, d_data, N * sizeof(float),
           cudaMemcpyDeviceToHost);

我个人习惯在拷贝之前加一个cudaDeviceSynchronize()，确保GPU计算全部完成。虽然cudaMemcpy本身会等待，但加上同步调用能让代码意图更清晰——别人一看就知道「哦，这里在等GPU算完」。

小技巧：如果你只需要取回部分结果，可以用cudaMemcpy指定偏移量。比如只取前100个元素：

cudaMemcpy(h_result, d_data + offset, 100 * sizeof(float),
           cudaMemcpyDeviceToHost);

这样能省下不少拷贝时间。

2.5 DeviceToDevice：显存内部的搬运

这个模式很多人容易忽略。其实在显存内部拷贝数据，速度非常快——因为它不走PCIe，完全在GPU内部完成。

float *d_src, *d_dst;
cudaMalloc(&d_src, N * sizeof(float));
cudaMalloc(&d_dst, N * sizeof(float));

// 在显存内部拷贝
cudaMemcpy(d_dst, d_src, N * sizeof(float),
           cudaMemcpyDeviceToDevice);

我曾经在做一个图像处理项目时，需要在显存里维护多个中间缓冲区。一开始我用HostToDevice来回倒腾，性能惨不忍睹。后来改成DeviceToDevice直接拷贝，速度提升了将近10倍。你想想看，PCIe的带宽和显存带宽根本不是一个量级——现代GPU的显存带宽动辄几百GB/s，而PCIe只有十几GB/s。

性能对比：

DeviceToDevice：~500 GB/s（取决于GPU型号）
HostToDevice/DeviceToHost：~12 GB/s（PCIe 3.0 x16）

差了40多倍！所以能用DeviceToDevice解决的问题，千万别绕到CPU那边去。

2.6 三种模式的流程图

下面这张图展示了三种拷贝模式的数据流向，我画出来方便你理解：

2.7 避坑指南

最后分享几个我踩过的坑：

内存泄漏：cudaMalloc分配的内存一定要用cudaFree释放。我曾经在一个循环里忘了释放，跑了几个小时程序直接OOM崩溃。
指针混淆：CPU指针和GPU指针长得一模一样，但你不能在CPU上解引用GPU指针。我见过有人直接printf(d_data)，结果段错误。
大小不匹配：拷贝的字节数一定要算对。我习惯用sizeof(type) * count，而不是硬编码数字。
异步陷阱：虽然这章讲的是同步拷贝，但你要知道cudaMemcpy默认是同步的。如果你用了cudaMemcpyAsync，那就得自己处理同步了。

我的建议：刚开始学CUDA的时候，先用同步拷贝把流程跑通。等你对内存模型有了感觉，再考虑用异步拷贝来优化性能。别一上来就追求花哨，稳扎稳打才是正道。

好了，这一章就到这里。三种拷贝模式其实不难，关键是理解它们各自的应用场景和性能特征。下一章我们会聊聊异步拷贝和流的概念，到时候你会发现，同步拷贝只是冰山一角。

📚 课程目录

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

2.1 同步拷贝的本质

2.2 三种传输方向

2.3 HostToDevice：把数据送进显存

2.4 DeviceToHost：把结果拿回来

2.5 DeviceToDevice：显存内部的搬运

2.6 三种模式的流程图

2.7 避坑指南