CUDA高性能计算系列06：流 (Stream) 与并发执行-编程实验室

CUDA高性能计算系列06：流 (Stream) 与并发执行

摘要：在之前的文章中，我们的视角主要集中在 GPU 内部（Kernel 优化）。但在宏观层面，CPU 和 GPU 是两个独立的处理器，GPU 内部也有拷贝引擎（Copy Engine）和计算引擎（Compute Engine）。本篇将引入CUDA Stream (流)的概念，教你如何像“流水线工厂”一样编排任务，实现数据传输与计算的完美重叠 (Overlap)，从而隐藏掉昂贵的 PCIe 传输延迟。

1. 默认的同步行为：串行执行

在没有显式使用 Stream 的情况下，所有的 CUDA 命令（Kernel 启动，cudaMemcpy）都会被提交给一个默认流 (Default Stream)。默认流保证了命令的顺序执行。

让我们回顾一下典型的深度学习训练循环：

for(inti=0;i<n_batches;++i){cudaMemcpy(d_in,h_in,size,H2D);// Copy H2D (Host to Device)myKernel<<<grid,block>>>(d_in,d_out);// Kernel ComputecudaMemcpy(h_out,d_out,size,D2H);// Copy D2H (Device to Host)}

时间线视图 (Timeline)：

问题：

当 GPU 在计算 Kernel 时，PCIe 总线是闲置的。
当数据在 PCIe 上传输时，GPU 核心是闲置的。
资源利用率低下。

2. CUDA Stream：异步并发的钥匙

CUDA Stream是一个 GPU 操作队列。

流内顺序：同一个流中的操作严格按顺序执行。
流间乱序：不同流中的操作可以并发执行（只要硬件资源允许）。

2.1 硬件引擎支持

现代 NVIDIA GPU 拥有独立的引擎：

Compute Engine：负责执行 Kernel。
Copy Engine (H2D)：负责 Host 到 Device 传输。
Copy Engine (D2H)：负责 Device 到 Host 传输。

这意味着，理论上我们可以同时做这三件事：向 GPU 拷入数据、GPU 计算、从 GPU 拷出数据。

3. 异步编程实战

要实现重叠，我们需要做两件事：

创建多个 Stream。
使用异步 API。

3.1 关键 API

// 1. 创建流cudaStream_t stream[n_streams];for(inti=0;i<n_streams;++i)cudaStreamCreate(&stream[i]);// 2. 异步内存拷贝 (必须使用 Pinned Memory!)// 注意：最后一个参数传入 stream[i]cudaMemcpyAsync(d_dst,h_src,size,kind,stream[i]);// 3. 异步 Kernel 启动// 第 4 个参数传入 stream[i] (第 3 个参数是 shared memory size，通常为 0)myKernel<<<grid,block,0,stream[i]>>>(...);// 4. 同步流 (等待流内所有操作完成)cudaStreamSynchronize(stream[i]);// 5. 销毁流cudaStreamDestroy(stream[i]);

3.2 锁页内存 (Pinned Memory)

非常重要：要使用cudaMemcpyAsync实现真正的异步传输，Host 端的内存必须是Pinned Memory (锁页内存)，不能是操作系统默认的分页内存。

// 申请float*h_ptr;cudaMallocHost((void**)&h_ptr,size);// 代替 malloc// 释放cudaFreeHost(h_ptr);// 代替 free

3.3 流水线设计 (Pipelining)

我们将大任务切分成NNN个小块（Chunks），分别放入NNN个 Stream 中处理。

深度优先 (Depth-First) 调度（推荐）：
Loop over streams:

Async Copy H2D (Streamiii)
Async Kernel (Streamiii)
Async Copy D2H (Streamiii)

时间线视图 (Pipeline)：

可以看到，Stream 2 的 Copy H2D和Stream 1 的 Kernel重叠了；Stream 3 的 Copy H2D、Stream 2 的 Kernel和Stream 1 的 Copy D2H甚至可以三者重叠！

4. 完整代码示例

#include<stdio.h>#include<cuda_runtime.h>#defineN30000000// 总数据量#defineN_STREAMS4// 流的数量constintSTREAM_SIZE=N/N_STREAMS;constintSTREAM_BYTES=STREAM_SIZE*sizeof(float);__global__voidsimpleKernel(float*a,float*b,float*c,intoffset){inti=offset+blockIdx.x*blockDim.x+threadIdx.x;if(i<N){// 增加计算负载，让 Kernel 耗时比 Copy 长，更容易观察到重叠for(intk=0;k<100;k++)c[i]=sinf(a[i])+cosf(b[i]);}}intmain(){float*h_a,*h_b,*h_c;float*d_a,*d_b,*d_c;// 1. 分配 Pinned Host MemorycudaMallocHost((void**)&h_a,N*sizeof(float));cudaMallocHost((void**)&h_b,N*sizeof(float));cudaMallocHost((void**)&h_c,N*sizeof(float));// 初始化数据...for(inti=0;i<N;i++){h_a[i]=i;h_b[i]=i;}// 2. 分配 Device MemorycudaMalloc((void**)&d_a,N*sizeof(float));cudaMalloc((void**)&d_b,N*sizeof(float));cudaMalloc((void**)&d_c,N*sizeof(float));// 3. 创建 StreamcudaStream_t stream[N_STREAMS];for(inti=0;i<N_STREAMS;++i)cudaStreamCreate(&stream[i]);// 4. 流水线循环for(inti=0;i<N_STREAMS;++i){intoffset=i*STREAM_SIZE;// Async Copy H2DcudaMemcpyAsync(&d_a[offset],&h_a[offset],STREAM_BYTES,cudaMemcpyHostToDevice,stream[i]);cudaMemcpyAsync(&d_b[offset],&h_b[offset],STREAM_BYTES,cudaMemcpyHostToDevice,stream[i]);// Async KernelsimpleKernel<<<STREAM_SIZE/256,256,0,stream[i]>>>(d_a,d_b,d_c,offset);// Async Copy D2HcudaMemcpyAsync(&h_c[offset],&d_c[offset],STREAM_BYTES,cudaMemcpyDeviceToHost,stream[i]);}// 5. 同步所有流 (等待所有任务完成)cudaDeviceSynchronize();// 6. 清理资源for(inti=0;i<N_STREAMS;++i)cudaStreamDestroy(stream[i]);cudaFreeHost(h_a);cudaFreeHost(h_b);cudaFreeHost(h_c);cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return0;}

5. 性能验证：Nsight Systems

要真正确认“重叠”发生了，最好的工具是Nsight Systems (nsys)。它能生成可视化的时间线。

运行命令：

nsys profile -o my_timeline ./my_stream_app

然后用 GUI 打开my_timeline.qdrep。

预期效果：
你会看到 Copy 栏和 Compute 栏在时间轴上是垂直对齐的（重叠），而不是阶梯状分布。如果 Copy H2D 比较快，Kernel 比较慢，你应该能看到 Kernel 紧密地排在一起，几乎没有空隙。

6. 总结与下篇预告

通过 CUDA Stream，我们打破了“传数据 -> 算数据 -> 传回数据”的串行枷锁，实现了Host 与 Device、Copy 与 Compute的全面并发。这对于处理视频流、实时推理等延迟敏感型任务至关重要。

现在，我们已经榨干了内存带宽（第3-5篇）和流水线并发（第6篇）。但是，我们的 Kernel 代码内部还在用着最基础的if-else吗？你知道if语句在 GPU 上可能导致巨大的性能陷阱吗？

下一篇[CUDA系列07_Warp Divergence与指令优化](./CUDA系列07_Warp Divergence与指令优化.md)，我们将深入微观指令层面，探讨Warp Divergence (线程束分化)现象，并学习如何写出对 SIMT 架构友好的分支代码。

参考文献

NVIDIA Corporation.CUDA C++ Programming Guide - Asynchronous Concurrent Execution. 2024.
Khronos Group.Vulkan & CUDA Interop / Concurrency.

CUDA高性能计算系列06：流 (Stream) 与并发执行