2. cudaMemcpy基础：同步拷贝与异步拷贝的API详解，默认流与多流场景下的行为差异
说到CUDA编程，内存拷贝是绕不开的基本功。我刚开始接触CUDA时，觉得cudaMemcpy不就是把数据从CPU搬到GPU嘛，有啥好研究的？结果第一次做大规模数据处理，性能直接拉胯，排查了半天才发现是拷贝方式没选对。今天咱们就把这块掰开揉碎了讲清楚。
2.1 同步拷贝：最直接的cudaMemcpy
先看最基础的API：
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count,
enum cudaMemcpyKind kind);这个函数大家肯定都用过。它的行为是阻塞式的——调用返回时，数据一定已经拷贝完成了。说白了，CPU会一直等在那里，直到GPU把活干完。
关键点：cudaMemcpy是同步的，它会阻塞当前CPU线程。
我有个习惯，在写原型代码时先用cudaMemcpy，因为它简单、不容易出错。但生产环境里，我几乎不用它做大批量数据传输——太慢了，CPU白白等着。
2.2 异步拷贝：cudaMemcpyAsync的威力
真正干活的是这个：
cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count,
enum cudaMemcpyKind kind, cudaStream_t stream);注意最后一个参数——stream。这是异步拷贝的灵魂。调用cudaMemcpyAsync后，函数会立即返回，CPU不用等。拷贝操作被提交到指定的流中，由GPU在后台执行。
我的经验：异步拷贝配合流水线操作，能把数据传输和计算重叠起来。我曾经在一个视频处理项目中，用异步拷贝把整体延迟降低了40%。
但有个坑——异步拷贝要求源地址和目的地址必须是页锁定内存（pinned memory）。如果你传了普通malloc分配的内存，cudaMemcpyAsync会默默退化成同步拷贝，而且不会报错。嗯，这里要注意。
2.3 默认流 vs 多流：行为天差地别
为什么要把流单独拎出来讲？因为流决定了操作之间的执行顺序。
默认流（NULL stream）
如果你不显式创建流，所有操作都在默认流中执行。默认流有个特殊性质——它会对所有其他流产生隐式同步。什么意思？
- 默认流中的操作必须等所有其他流中的操作完成才能开始
- 其他流中的操作也必须等默认流中的操作完成才能开始
说白了，默认流就是个"霸道总裁"，谁都得让着它。这会导致你明明用了多流，却得不到并行效果。
避坑指南：我曾经在项目里同时用了默认流和自定义流，以为能并行执行，结果性能纹丝不动。排查了半天才发现是默认流在作祟。从那以后，我只要用多流，就坚决不用默认流。
多流场景
创建多个非默认流，才能真正实现操作并行：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 异步拷贝到不同流
cudaMemcpyAsync(dst1, src1, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(dst2, src2, size, cudaMemcpyHostToDevice, stream2);
// 不同流中的kernel可以并行执行
kernel1<<<grid, block, 0, stream1>>>(dst1);
kernel2<<<grid, block, 0, stream2>>>(dst2);同一个流中的操作是串行的，不同流中的操作可以并行。你想想看，如果数据拷贝和计算能同时进行，那效率能提升多少？
2.4 知识体系总览
下面这张图把核心逻辑串起来了：
2.5 实际场景对比
为了让你更直观地理解差异，我整理了一个对比表：
| 特性 | cudaMemcpy（同步） | cudaMemcpyAsync（异步） |
|---|---|---|
| CPU行为 | 阻塞等待 | 立即返回 |
| 内存要求 | 任意内存 | 必须页锁定内存 |
| 流支持 | 默认流（隐式） | 可指定任意流 |
| 并行能力 | 无（串行执行） | 可与计算重叠 |
| 适用场景 | 小数据量、原型验证 | 大数据量、性能优化 |
个人建议：刚开始学CUDA时，先用cudaMemcpy把功能跑通。等代码稳定了，再逐步替换成cudaMemcpyAsync做性能优化。一步到位反而容易踩坑。
好了，这一章的内容就到这里。记住核心：同步拷贝简单但低效，异步拷贝灵活但需要小心内存类型和流管理。下一章我们会深入页锁定内存的分配和管理，那是异步拷贝的基石。