📚 课程目录第2章内存拷贝的同步与异步

2. 内存拷贝的同步与异步

做吞吐压测，最怕什么？

我最怕的就是GPU在那闲着等数据。你想想看，一张A100跑着几百T的算力，结果因为内存拷贝没处理好，计算单元空转，那画面太美我不敢看。

这一节，咱们就聊聊内存拷贝的同步与异步。说白了，就是怎么让GPU干活的时候，数据也能在后台悄悄传完。

2.1 同步拷贝：cudaMemcpy 的"笨办法"

cudaMemcpy 是大家最熟悉的API。它干一件事：把数据从A搬到B，搬完才返回。

我刚开始做压测时，就吃过它的亏。代码写得很漂亮，结果一跑profiling，发现GPU利用率不到30%。为什么？因为每次 cudaMemcpy 都在阻塞整个设备。

核心特点：

默认在默认流（default stream）中执行
调用返回时，数据拷贝已经完成
会阻塞后续所有CUDA操作（包括kernel launch）

// 同步拷贝示例
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// 这行代码返回时，数据已经拷完了
// 但GPU可能还在忙别的？不，它得等你拷完
myKernel<<<grid, block>>>(d_data);

嗯，这里要注意：默认流在CUDA 7.0之后有个"坑"。如果你没显式创建流，所有操作都挤在默认流里，它会跟其他流隐式同步。我当年排查一个性能问题，查了两天才发现是默认流在作祟。

2.2 异步拷贝：cudaMemcpyAsync 的"骚操作"

那怎么让拷贝不阻塞？用 cudaMemcpyAsync。

这个API调用后立刻返回，数据拷贝在后台进行。但有个前提——你得用 固定内存（pinned memory）。

避坑指南：

我曾经在项目里直接用 malloc 分配的内存传给 cudaMemcpyAsync，结果程序直接崩了。后来才想起来，异步拷贝需要页锁定内存，否则DMA控制器没法直接访问。

// 正确做法：先分配固定内存
cudaHostAlloc(&h_data, size, cudaHostAllocDefault);
// 然后异步拷贝
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);
// 这行代码立刻返回，拷贝在后台进行
myKernel<<<grid, block, 0, stream>>>(d_data);

为什么一定要固定内存？你想想看，GPU的DMA引擎要直接读写主机内存，如果内存被操作系统换页了，那数据不就乱套了？固定内存就是告诉OS："这块内存你别动，GPU要用。"

2.3 流(Stream)：让GPU变成"多线程"

流是什么？我习惯把它理解成GPU上的"独立执行队列"。

每个流里的操作是顺序执行的，但不同流之间的操作可以并行。这就好比一条流水线上，你可以同时处理多个订单。

我的经验：

在吞吐压测中，我一般会创建2-4个流。不是越多越好，因为GPU的硬件队列有限。我记得有一次为了炫技，创建了16个流，结果性能反而下降了——调度开销太大了。

// 创建多个流
cudaStream_t streams[4];
for (int i = 0; i < 4; i++) {
    cudaStreamCreate(&streams[i]);
}

// 每个流独立执行拷贝+计算
for (int i = 0; i < 4; i++) {
    cudaMemcpyAsync(d_data + i * chunk, h_data + i * chunk, 
                    chunk, cudaMemcpyHostToDevice, streams[i]);
    myKernel<<<grid, block, 0, streams[i]>>>(d_data + i * chunk);
}

2.4 事件(Event)：精准的"计时器"与"同步点"

流之间怎么同步？用事件。

事件就像赛道上的计时器。你可以记录某个操作什么时候完成，也可以让一个流等另一个流的事件。

事件的两个核心用途：

计时：精确测量GPU操作的耗时
同步：让流B等待流A的某个操作完成

// 事件计时示例
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream);
myKernel<<<grid, block, 0, stream>>>(d_data);
cudaEventRecord(stop, stream);

cudaEventSynchronize(stop);  // 等待事件完成
float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel耗时: %f ms\n", ms);

事件同步更实用。比如流0在拷贝数据，流1想等它拷完再计算：

cudaEvent_t copyDone;
cudaEventCreate(&copyDone);

// 流0：拷贝数据
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream0);
cudaEventRecord(copyDone, stream0);  // 记录拷贝完成事件

// 流1：等待拷贝完成，然后计算
cudaStreamWaitEvent(stream1, copyDone, 0);
myKernel<<<grid, block, 0, stream1>>>(d_data);

2.5 协同机制：一张图看懂全流程

说了这么多，咱们画张图总结一下。这张图展示了我最常用的"双缓冲+双流"模式：

2.6 实战建议：压测中的最佳实践

说了这么多理论，来点干货。我在做吞吐压测时，一般遵循这几个原则：

固定内存是前提——没有固定内存，异步拷贝就是空谈。用 cudaHostAlloc 替代 malloc。
流数量要适中——2-4个流通常最优。太多反而增加调度开销。
事件要善用——别滥用 cudaDeviceSynchronize，那会杀死所有并行性。用事件做细粒度同步。
拷贝粒度要匹配——单次拷贝太小（<1MB）会浪费带宽，太大（>100MB）会阻塞太久。我一般控制在4-16MB。

一个小技巧：

用 cudaMemcpyAsync 时，可以配合 cudaStreamSynchronize 做流级别的等待，而不是全局同步。这样其他流还能继续干活。

好了，这一节的内容就到这。记住一句话：同步拷贝保平安，异步拷贝提性能，流和事件是让GPU忙起来的关键。