4. CUDA流与并发：让GPU真正忙起来

各位同学，今天我们来聊聊CUDA流。说实话，这是很多初学者容易忽略，但实际项目中又极其重要的一个概念。我刚开始做GPU加速时，就犯过一个低级错误——把所有操作都塞进默认流，结果GPU利用率惨不忍睹。后来才明白，CUDA流就是让GPU真正并行起来的关键。

4.1 什么是CUDA Stream？

简单来说，CUDA流就是一个操作队列。你把kernel启动、内存拷贝、事件记录等操作提交到流里，它们就会按顺序执行。但不同流之间的操作，是可以并发执行的。

你想想看，GPU里有那么多计算单元、内存控制器、拷贝引擎，如果只用一个流，那大部分硬件资源都在闲着。多流并发，才能把GPU的潜力榨干。

我习惯把流比作工厂里的流水线。一条流水线只能按顺序处理任务，但多条流水线同时开工，总产量就上去了。嗯，这个比喻虽然不完美，但意思到了。

4.2 Stream的创建与销毁

创建流很简单，API就两个：

cudaStream_t stream;
cudaError_t cudaStreamCreate(&stream);   // 创建
cudaError_t cudaStreamDestroy(stream);   // 销毁

我在项目中遇到过一个问题：频繁创建销毁流会导致性能抖动。后来我改用流池（Stream Pool），提前创建好一批流，按需分配。这样既避免了创建开销，又能控制并发度。

还有一个容易被忽略的点：默认流（NULL流）的行为。在CUDA 7.0之前，默认流会阻塞其他流。7.0之后引入了“每线程默认流”，每个host线程有自己的默认流，互不干扰。这个改动其实挺大的，我当年升级CUDA版本时就踩过这个坑。

4.3 异步操作：别让CPU闲着

CUDA的异步操作，说白了就是让CPU和GPU各干各的。你提交一个kernel启动命令，CPU不等它执行完，直接返回继续干别的。这才是真正的并行。

常见的异步操作包括：

• cudaMemcpyAsync：异步内存拷贝，需要配合流使用
• kernel启动：默认就是异步的（除非使用同步版本的API）
• cudaMemsetAsync：异步内存置零

举个例子：

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

// 异步拷贝到设备
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);

// 异步启动kernel
kernel1<<<grid, block, 0, stream1>>>(d_data1);
kernel2<<<grid, block, 0, stream2>>>(d_data2);

// 异步拷贝回主机
cudaMemcpyAsync(h_result1, d_data1, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(h_result2, d_data2, size, cudaMemcpyDeviceToHost, stream2);

cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

注意看，这里我用了cudaMemcpyAsync而不是cudaMemcpy。前者是异步的，后者是同步的。我曾经在代码里混用过这两个，结果发现性能还不如单流——因为同步拷贝把整个流水线都卡住了。

4.4 多Stream并发执行

多流并发，核心在于让不同流之间的操作重叠。比如：流1在计算，流2在拷贝，流3在等待。这样GPU的各个硬件单元都能同时工作。

我画了一张图，帮你理解多流并发的执行流程：

从图上可以看到，Stream1在拷贝时，Stream2在计算，Stream3在等待。这样GPU的拷贝引擎和计算单元就能同时工作。理想情况下，总执行时间接近最长的那个操作，而不是所有操作之和。

但要注意，并不是流越多越好。我记得有一次测试，把流数从4增加到16，性能反而下降了。原因有两个：

1. 硬件资源有限：GPU的拷贝引擎、计算单元数量是固定的。流太多会导致上下文切换开销。
2. 内存带宽竞争：多个流同时访问显存，会争抢带宽，反而降低效率。

4.5 流同步与事件

多流并发虽然好，但有时候我们需要让流之间协同工作。比如：Stream2的kernel必须等Stream1的数据拷贝完成才能启动。这时候就需要事件（Event）来同步。

cudaEvent_t event;
cudaEventCreate(&event);

// Stream1中记录事件
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream1);
cudaEventRecord(event, stream1);

// Stream2中等待事件
cudaStreamWaitEvent(stream2, event, 0);
kernel2<<<grid, block, 0, stream2>>>(d_data);

cudaEventDestroy(event);

这里cudaStreamWaitEvent是关键。它告诉Stream2：在event被记录之前，别执行后面的操作。这样就能实现流间的依赖关系。

我曾经在做一个多GPU流水线项目时，就是靠事件来协调不同GPU上的流。每个GPU处理完数据后，记录一个事件，下一个GPU的流等待这个事件。这样整个流水线就串起来了。

4.6 总结与建议

好了，关于CUDA流的内容就讲到这里。我最后给你几个实用建议：

• 默认流够用吗？ 对于简单任务，默认流就够了。但如果你想压榨GPU性能，多流是必须的。
• 先测再优化：不要盲目增加流数。用NVIDIA Nsight分析一下，看看GPU利用率是不是真的上去了。
• 注意内存固定：异步拷贝必须用固定内存，否则性能会大打折扣。
• 事件别滥用：事件同步会引入开销。能用流内顺序解决的，就别用事件。

嗯，这些经验都是我在实际项目中一点点踩坑踩出来的。希望你能少走弯路。

想要更多资料，联系可可