3. CUDA流基础

核心原理:什么是CUDA流

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

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

流的核心价值:

  • 并发执行:不同流中的操作可以重叠,提高GPU利用率
  • 依赖控制:通过事件(Event)或流同步机制精确控制任务间的先后关系
  • 延迟隐藏:在数据拷贝的同时执行计算,掩盖PCIe传输延迟
  • 优先级调度:为不同流设置优先级,影响GPU调度器的任务选择顺序
CUDA流调度模型 默认流(流0)— 串行执行 H2D拷贝 内核A 内核B D2H拷贝 所有操作按提交顺序依次执行,GPU空闲等待 多流(流1/流2)— 并发执行 流1: H2D拷贝 流1: 内核A 流2: H2D拷贝 流2: 内核B 流1与流2的操作可重叠执行,提高吞吐量 时间 流1: 拷贝 流1: 内核A 流1: 拷贝 流2: 拷贝 流2: 内核B 流2: 拷贝 重叠区 cudaStreamSynchronize (流1) 优先级调度 • 高优先级流(cudaStreamPriorityHigh)优先获得SM资源 • 低优先级流(cudaStreamPriorityLow)在空闲时执行 • 优先级影响延迟敏感型任务的响应速度 内存拷贝 内核执行 同步点

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批的内核计算重叠,显著提升吞吐。

流水线实现要点:

  1. 将数据划分为N个批次,每个批次绑定一个流
  2. 每个流内按“H2D拷贝 → 内核执行 → D2H拷贝”的顺序提交
  3. 使用 cudaMemcpyAsync 配合流实现异步拷贝
  4. 主机侧使用 cudaStreamSynchronize 等待最终结果
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]);
}

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

陷阱1:默认流的隐式同步

默认流(流0)会与所有非默认流隐式同步。如果某个流使用 cudaMemcpy(同步版本)而非 cudaMemcpyAsync,它会阻塞默认流,导致并发失效。解决:统一使用 cudaMemcpyAsync 并创建流时加 cudaStreamNonBlocking 标志。

陷阱2:过度同步导致性能归零

在每次内核启动后都调用 cudaDeviceSynchronize() 会彻底破坏异步执行。正确做法是只在需要结果时才同步,或使用事件进行细粒度依赖控制。

陷阱3:忽略优先级调度的实际效果

优先级只影响尚未占用SM资源的任务调度顺序。如果低优先级任务已经占用SM,高优先级任务仍需等待其完成。因此,优先级调优应配合内核占用率控制(如限制block数量)才能发挥效果。

复盘建议:使用 nsys profile 或 ncu 工具查看流的时间线,确认操作是否真正重叠。如果发现流之间出现意外的间隙,优先检查是否存在隐式同步或同步版本API调用。