3. CUDA流基础：CUDA流的概念、创建与同步机制

核心原理：什么是CUDA流

CUDA流（Stream）是GPU上按顺序执行的一系列操作的队列。从程序员视角看，流是GPU任务的“传送带”——同一流中的内核启动、内存拷贝、事件记录等操作严格按照提交顺序执行；不同流之间则相互独立，可以并发执行。从硬件视角看，流对应GPU上可并行调度的任务序列，现代GPU（如Ampere、Hopper架构）拥有多个硬件工作队列，能够同时处理来自不同流的任务。

理解流的关键在于：流本身不提供计算能力，它提供的是“调度边界”。默认情况下，所有CUDA操作都在默认流（default stream，流0）中执行，这意味着它们被串行化。当我们需要让内存拷贝与内核计算重叠、让多个内核在不同SM（流式多处理器）上并行时，就必须创建多个流。

CUDA流的创建与生命周期

创建流使用 cudaStreamCreate()，释放使用 cudaStreamDestroy()。CUDA 11.0 之后推荐使用 cudaStreamCreateWithPriority() 来指定流的优先级。流的创建是轻量级操作，但频繁创建/销毁会带来不必要的API调用开销，因此建议在初始化阶段创建好所需流并复用。

// 创建默认优先级流
cudaStream_t stream1;
cudaStreamCreate(&stream1);

// 创建带优先级的流（-2为最高优先级，0为默认，2为最低）
int priority_low, priority_high;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
cudaStream_t stream_high, stream_low;
cudaStreamCreateWithPriority(&stream_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&stream_low, cudaStreamNonBlocking, priority_low);

// 使用流启动内核
kernelA<<<grid, block, 0, stream_high>>>(args);
kernelB<<<grid, block, 0, stream_low>>>(args);

// 释放流
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream_high);
cudaStreamDestroy(stream_low);

注意 cudaStreamNonBlocking 标志：它使该流不隐式同步默认流。默认情况下，默认流会等待所有非默认流完成，反之亦然——这种隐式同步会破坏并发性。使用 cudaStreamNonBlocking 可以解除这种约束，让流真正独立运行。

同步机制：事件与流同步

CUDA提供了多种同步原语，理解它们的语义是避免死锁和性能陷阱的关键：

同步API	作用	阻塞范围
cudaDeviceSynchronize()	等待所有流中的所有操作完成	整个设备，最重
cudaStreamSynchronize(stream)	等待指定流中的所有操作完成	仅该流
cudaStreamWaitEvent(stream, event)	让流等待某个事件发生	流内后续操作
cudaEventRecord(event, stream)	在流中记录一个事件标记	无阻塞，异步
cudaEventSynchronize(event)	等待事件被记录完成	主机线程

事件（Event）是流间依赖的“信号灯”。典型用法是：在流A中记录事件，然后让流B等待该事件，从而实现跨流依赖而不阻塞主机。

cudaEvent_t event;
cudaEventCreate(&event);

// 流1：计算中间结果
kernel_preprocess<<<grid, block, 0, stream1>>>(input, temp);
cudaEventRecord(event, stream1);  // 在流1中记录事件

// 流2：等待流1完成后再执行
cudaStreamWaitEvent(stream2, event, 0);
kernel_main<<<grid, block, 0, stream2>>>(temp, output);

// 主机等待流2完成
cudaStreamSynchronize(stream2);
cudaEventDestroy(event);

案例：流水线重叠优化

一个典型的深度学习推理场景：需要将输入数据从主机拷贝到设备（H2D），执行推理内核，再将结果拷回主机（D2H）。如果使用默认流，整个过程串行执行；使用多流流水线后，第i+1批数据的H2D拷贝可以与第i批的内核计算重叠，显著提升吞吐。

for (int i = 0; i < N; i++) {
    int stream_idx = i % num_streams;
    cudaStream_t s = streams[stream_idx];
    
    // 异步拷贝输入到设备
    cudaMemcpyAsync(d_inputs[i], h_inputs[i], bytes, 
                    cudaMemcpyHostToDevice, s);
    // 在流中执行推理内核
    inference_kernel<<<grid, block, 0, s>>>(d_inputs[i], d_outputs[i]);
    // 异步拷贝结果回主机
    cudaMemcpyAsync(h_outputs[i], d_outputs[i], bytes, 
                    cudaMemcpyDeviceToHost, s);
}
// 等待所有流完成
for (int i = 0; i < num_streams; i++) {
    cudaStreamSynchronize(streams[i]);
}

误区与复盘：流使用的常见陷阱