第三节：cudaMemcpyAsync与流——异步拷贝的核心机制

各位同学，今天我们来聊聊CUDA编程里一个非常关键的话题——异步拷贝。说白了，就是让数据传输和计算同时进行，别让GPU闲着等数据。

我记得刚接触CUDA那会儿，写了个简单的向量加法，数据拷贝花了大部分时间，计算反而一眨眼就完了。当时我就想：这效率也太低了吧？后来才明白，原来有cudaMemcpyAsync这个好东西。

3.1 为什么需要异步拷贝？

先看一个典型的CUDA程序流程：

cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);  // 拷贝到GPU
kernel<<<grid, block>>>(d_data);                          // 执行计算
cudaMemcpy(h_result, d_data, size, cudaMemcpyDeviceToHost); // 拷回结果

这段代码的问题在哪？每一步都是同步的。CPU在调用cudaMemcpy时，会一直等到数据传输完成才返回。你想想看，这期间CPU啥也干不了，GPU也在等数据到位才能开始算。

说白了，这就是串行执行——浪费了宝贵的并行能力。

3.2 CUDA流——异步操作的容器

要实现异步，我们需要一个东西来管理这些并发的操作。这就是CUDA流（Stream）。

你可以把流想象成一条流水线。每个流里的操作是顺序执行的，但不同流之间的操作可以并发执行。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 在stream1中异步拷贝
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
// 在stream2中异步拷贝
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);

// 在stream1中执行内核
kernel1<<<grid, block, 0, stream1>>>(d_data1);
// 在stream2中执行内核
kernel2<<<grid, block, 0, stream2>>>(d_data2);

这里有个细节要注意：cudaMemcpyAsync必须配合固定内存（pinned memory）使用。我在项目中遇到过，有人直接用普通malloc分配的内存传给cudaMemcpyAsync，结果程序直接报错。嗯，这个坑我踩过。

3.3 数据传输与计算的重叠模型

现在我们来画一张图，看看异步拷贝到底怎么工作的。

从这张图可以看得很清楚：同步模式下，拷贝和计算是串行的。异步模式下，不同流里的操作可以重叠执行，大大缩短了总时间。

3.4 实战：双缓冲技术

我个人最喜欢的一个技巧是双缓冲（Double Buffering）。说白了，就是准备两块缓冲区，一块在计算，另一块在传输数据，交替进行。

// 双缓冲示例
cudaStream_t stream[2];
cudaStreamCreate(&stream[0]);
cudaStreamCreate(&stream[1]);

float *d_data[2], *h_data[2];
// 分配固定内存和显存...

for (int i = 0; i < num_batches; i++) {
    int current = i % 2;
    int next = (i + 1) % 2;
    
    // 异步拷贝下一批数据
    if (i + 1 < num_batches) {
        cudaMemcpyAsync(d_data[next], h_data[next], 
                       batch_size, cudaMemcpyHostToDevice, stream[next]);
    }
    
    // 在当前流中执行内核
    kernel<<<grid, block, 0, stream[current]>>>(d_data[current]);
    
    // 异步拷回结果
    cudaMemcpyAsync(h_data[current], d_data[current],
                   batch_size, cudaMemcpyDeviceToHost, stream[current]);
}

我曾经在一个图像处理项目里用这个技巧，把处理时间从原来的5秒降到了2.8秒。效果立竿见影。

3.5 流的同步与事件

异步操作多了，怎么知道什么时候完成？这就需要事件（Event）来帮忙了。

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream1);
// ... 执行一些操作 ...
cudaEventRecord(stop, stream1);

// 等待事件完成
cudaEventSynchronize(stop);

float elapsed;
cudaEventElapsedTime(&elapsed, start, stop);
printf("耗时: %.2f ms\n", elapsed);

3.6 避坑指南

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

• 固定内存分配失败：cudaHostAlloc分配固定内存时，如果系统内存不足会失败。我曾经在内存紧张的服务器上吃过这个亏，后来改用cudaHostAllocPortable标志，效果好了很多。
• 流同步遗漏：在程序结束前，一定要用cudaStreamSynchronize等待所有流完成。否则程序退出时，异步操作可能还没执行完。
• 过度使用流：不是流越多越好。GPU的并发能力有限，一般2-4个流就够用了。我曾经试过创建16个流，结果性能反而下降了。
• 内存拷贝方向：cudaMemcpyAsync的最后一个参数是流，别漏了。我见过有人写cudaMemcpyAsync(d, h, size, cudaMemcpyHostToDevice)——少了流参数，编译能过，但实际是同步的。

好了，关于cudaMemcpyAsync和流的内容就讲到这里。记住，异步拷贝的核心就是让GPU别闲着，数据一边传，计算一边跑。这个思想在很多高性能计算场景里都适用。