3. 异步传输

好，咱们进入第三个核心话题。前面讲了同步传输和页锁定内存，说白了都是在为“怎么让数据搬得更快”做准备。但真正让GPU发挥威力的，是异步传输。

我刚开始接触CUDA时，总觉得GPU计算就是“把数据拷过去，算，再拷回来”。后来发现，这其实是在浪费生命。你想想看，GPU在算的时候，CPU在干嘛？大部分时间在傻等。异步传输，就是让CPU和GPU各干各的，谁也不耽误谁。

3.1 流（Stream）—— 异步传输的基石

流是什么？你可以把它想象成一条任务流水线。每个流里，任务按顺序执行。但不同流之间，任务可以并行。

默认情况下，所有操作都在同一个流（默认流）里。这就像只有一个工人，干完一件才能干下一件。而创建多个流，就像雇了好几个工人，可以同时干活。

但要注意，cudaMemcpyAsync 必须搭配页锁定内存使用。我之前讲过，页锁定内存是异步传输的前提。如果你传的是可分页内存，cudaMemcpyAsync 会退化成同步版本，那就白忙活了。

创建和使用流

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);

// 同步等待所有流完成
cudaDeviceSynchronize();

// 清理
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

你看，代码里每个操作都指定了流。这样，stream1的数据传输和stream2的数据传输可以同时进行。如果硬件支持，甚至计算和传输也能重叠。

3.2 双缓冲（Double Buffering）—— 让流水线永不空转

双缓冲，说白了就是准备两份数据缓冲区。一份在计算，另一份在传输。这样，计算和传输就能像流水线一样，无缝衔接。

我举个例子。假设你要处理100个数据块。没有双缓冲时，流程是这样的：

1. 传输数据块1到GPU
2. GPU计算数据块1
3. 传输结果回CPU
4. 传输数据块2到GPU
5. GPU计算数据块2
6. ...

每一步都是串行的，GPU经常闲着等数据。

用了双缓冲后，流程变成：

1. 传输数据块1到缓冲区A
2. GPU开始计算缓冲区A的数据
3. 同时，传输数据块2到缓冲区B
4. GPU计算完A，立刻切换到缓冲区B
5. 同时，传输数据块3到缓冲区A（覆盖旧数据）
6. ...

这样，GPU几乎一直在算，数据也一直在传。效率高多了。

双缓冲代码示例

// 假设有N个数据块
for (int i = 0; i < N; i++) {
    // 当前使用的缓冲区索引
    int buffer_idx = i % 2;

    // 异步传输数据到当前缓冲区
    cudaMemcpyAsync(d_buffer[buffer_idx], h_data[i], size,
                    cudaMemcpyHostToDevice, stream_transfer);

    // 等待传输完成（用事件同步）
    cudaEventRecord(event_transfer_done, stream_transfer);
    cudaStreamWaitEvent(stream_compute, event_transfer_done);

    // 在计算流中启动内核
    kernel<<<grid, block, 0, stream_compute>>>(d_buffer[buffer_idx]);

    // 记录计算完成事件
    cudaEventRecord(event_compute_done, stream_compute);
    cudaStreamWaitEvent(stream_transfer, event_compute_done);
}

嗯，这里要注意。事件同步是必须的。我曾经犯过一个错误：忘了加 cudaStreamWaitEvent，结果计算流在数据还没传完时就启动了内核，算出来的全是垃圾数据。排查了好久才发现。

3.3 计算与传输重叠 —— 真正的性能提升

计算与传输重叠，是异步传输的终极目标。说白了，就是让GPU一边算，一边传数据。这样，总时间就变成了 max(计算时间, 传输时间)，而不是 计算时间 + 传输时间。

但有个前提：你的GPU必须支持并发执行。大部分现代GPU都支持，但老一点的型号可能不行。可以用 cudaDeviceGetAttribute 来检查：

int concurrent_kernels;
cudaDeviceGetAttribute(&concurrent_kernels,
                       cudaDevAttrConcurrentKernels, 0);
if (concurrent_kernels) {
    printf("支持并发内核执行\n");
}

另外，数据传输方向也有讲究。Host到Device的传输，和Device到Host的传输，可以同时进行。但同方向的传输，一般只能串行。所以，设计双缓冲时，最好让传输和计算的方向错开。

3.4 知识体系图

下面这张图，展示了异步传输、流、双缓冲之间的关系。你可以把它当作一个参考框架。

3.5 实战中的注意事项

最后，我总结几个实战中容易踩的坑：

• 内存分配：双缓冲需要两份设备内存。别算错了，否则会越界。
• 事件同步：一定要用 cudaEvent 来同步流。别用 cudaDeviceSynchronize，那会阻塞所有流，得不偿失。
• 数据依赖：如果下一个数据块的计算依赖上一个结果，那双缓冲就没用了。这种情况，老老实实串行吧。
• 性能调优：我习惯用 nvprof 或 Nsight Systems 来查看时间线。如果发现传输和计算没有重叠，说明流没用好。

好了，异步传输这块就讲到这里。记住，流是工具，双缓冲是方法，重叠是目标。三者结合，才能让GPU真正跑起来。