Toriskia 's Blog

13 篇文章 · 12 个标签 · 6 个友链

← 返回文章列表

2026.04.28

CUDA 笔记（四）：异步执行

CUDA MLSys

一、什么是 Asynchronous Concurrent Execution#

CUDA 允许以下操作同时进行：

  • 在 host 上的计算
  • 在 device 上的计算
  • host 到 device 的内存传输
  • device 到 host 的内存传输
  • device 内部的内存传输
  • device 之间的内存传输

这些并发通过异步接口实现：派发函数或 kernel launch 立即返回（可能早于实际操作开始）。应用可以在等待操作完成的同时执行其他任务。当需要操作的最终结果时，必须通过某种形式的同步来确保操作已经完成。

典型的并发模式是让内存传输和计算重叠，从而减少甚至消除传输开销。

CUDA Streams

异步接口提供三种主要的同步方式：

  • 阻塞（blocking）：调用一个阻塞函数，等待操作完成
  • 轮询（polling）：调用一个立即返回的函数，提供操作状态信息
  • 回调（callback）：操作完成后执行预先注册的函数

CUDA 中异步执行的核心组件是 CUDA Stream 和 CUDA Event。

二、CUDA Stream#

Stream 是一种抽象，允许程序表达一个操作序列。Stream 像一个工作队列，程序可以向其中添加操作（内存复制、kernel launch 等），操作按入队顺序依次执行。

应用可以同时使用多个 stream。此时 runtime 会根据 GPU 资源状态从有任务就绪的 stream 中选择任务执行。

Stream 中的 API 调用和 kernel launch 对于 host 线程都是异步的。应用可以等待某个 stream 变空，也可以在 device 级别同步。

CUDA 有一个默认 stream，所有未显式指定 stream 的操作都会被排入其中。默认 stream 有特殊语义（见第五节）。

1. 创建和销毁 Stream#

cpp
1cudaStream_t stream; 2cudaStreamCreate(&stream); // 创建 3 4// 在 stream 中执行操作... 5 6cudaStreamDestroy(stream); // 销毁（会等待 stream 中所有工作完成后才销毁）

2. 在 Stream 中启动 Kernel#

在 <<<>>> 的第四个参数指定 stream：

cpp
1kernel<<<grid, block, shared_mem_size, stream>>>(...);

kernel launch 是异步的，调用立即返回。kernel 将在 stream 中执行，CPU 在此期间可以执行其他任务。

3. 在 Stream 中启动内存传输#

使用 cudaMemcpyAsync()，它在普通 cudaMemcpy() 的基础上增加了 stream 参数：

cpp
1// 将 size 字节从 src 复制到 dst，在 stream 中异步执行 2cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);

调用立即返回，而 cudaMemcpy() 会阻塞直到传输完成。

为了让涉及 CPU 内存的复制真正异步执行，host buffer 必须是页锁定（page-locked）内存。如果使用普通可分页内存，cudaMemcpyAsync() 会退化为同步行为。推荐使用 cudaMallocHost() 分配与 GPU 通信的 host buffer。

4. Stream 同步#

阻塞方式 cudaStreamSynchronize()：阻塞直到 stream 中所有工作完成。

cpp
1cudaStreamSynchronize(stream); 2// 此处 stream 已空，可以安全访问结果

轮询方式 cudaStreamQuery()：立即返回 stream 状态。

cpp
1cudaError_t status = cudaStreamQuery(stream); 2 3switch (status) { 4 case cudaSuccess: 5 // stream 已空 6 break; 7 case cudaErrorNotReady: 8 // stream 不为空 9 break; 10 default: 11 // 发生错误 12 break; 13}

三、CUDA Event#

Event 是一种在 stream 中插入标记的机制，类似于跟踪粒子，用于追踪 stream 中任务的进度。通过 event 可以：

  • 在第一个 kernel 完成后、第二个 kernel 完成前就开始依赖操作
  • 测量 kernel 或内存传输的耗时

1. 创建和销毁 Event#

cpp
1cudaEvent_t event; 2cudaEventCreate(&event); 3 4// 使用 event... 5 6cudaEventDestroy(event);

2. 在 Stream 中插入 Event#

cpp
1cudaEventRecord(event, stream);

Event 被插入到 stream 的当前队尾，当 stream 执行到此处时记录时间戳。

3. 计时#

cpp
1cudaStream_t stream; 2cudaStreamCreate(&stream); 3 4cudaEvent_t start, stop; 5cudaEventCreate(&start); 6cudaEventCreate(&stop); 7 8cudaEventRecord(start, stream); 9kernel<<<grid, block, 0, stream>>>(...); 10cudaEventRecord(stop, stream); 11 12cudaStreamSynchronize(stream); 13 14float elapsedTime; 15cudaEventElapsedTime(&elapsedTime, start, stop); 16 17cudaEventDestroy(start); 18cudaEventDestroy(stop); 19cudaStreamDestroy(stream);

4. 检查 Event 状态#

与 stream 类似，分阻塞和轮询两种方式：

阻塞 cudaEventSynchronize()：

cpp
1cudaEvent_t event; 2cudaStream_t stream; 3cudaStreamCreate(&stream); 4cudaEventCreate(&event); 5 6kernel1<<<grid, block, 0, stream>>>(...); 7cudaEventRecord(event, stream); // 在 kernel1 后插入 event 8kernel2<<<grid, block, 0, stream>>>(...); 9 10// 等待 event——此时 kernel1 一定完成了，但 kernel2 可能还在执行 11cudaEventSynchronize(event); 12dependentCPUtask(); // 安全地启动依赖 kernel1 结果的任务 13 14cudaStreamSynchronize(stream); // 此时 kernel2 也完成 15cudaEventDestroy(event); 16cudaStreamDestroy(stream);

轮询 cudaEventQuery()：

以下例子展示了如何让 CPU 工作与 GPU 执行及数据传回重叠：

cpp
1cudaStream_t stream1, stream2; 2cudaEvent_t event; 3bool copyStarted = false; 4 5cudaStreamCreate(&stream1); // 计算流 6cudaStreamCreate(&stream2); // 复制流 7cudaEventCreate(&event); 8 9kernel1<<<grid, block, 0, stream1>>>(d_data, size); 10cudaEventRecord(event, stream1); // kernel1 完成后触发 11kernel2<<<grid, block, 0, stream1>>>(); // kernel2 也可以继续跑 12 13while (!allCPUWorkDone() || !copyStarted) { 14 doNextChunkOfCPUWork(); // 同时做 CPU 工作 15 16 if (!copyStarted && cudaEventQuery(event) == cudaSuccess) { 17 // kernel1 已完成，开始在 stream2 中异步传回数据 18 cudaMemcpyAsync(h_data, d_data, size, cudaMemcpyDeviceToHost, stream2); 19 copyStarted = true; 20 } 21} 22 23cudaStreamSynchronize(stream1); 24cudaStreamSynchronize(stream2);

四、Callback 函数#

CUDA 允许在 stream 中插入 host 端回调函数。

cpp
1cudaError_t cudaLaunchHostFunc(cudaStream_t stream, void (*func)(void *), void *data);

回调函数签名：

cpp
1void hostFunction(void *data);

注意：

  • 回调函数中不能调用任何 CUDA API
  • Stream 在回调函数执行期间被视为空闲，后续所有被排入同一个流的 kernel 或内存传输都必须等这个回调函数返回后才会启动
  • 回调函数的开始等价于同步了一个在其之前插入的 event

五、Blocking / Non-blocking Stream 与默认 Stream#

1. 创建 Non-blocking Stream#

默认情况下 cudaStreamCreate() 创建的是 blocking stream。要创建 non-blocking stream：

cpp
1cudaStream_t stream; 2cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

2. Legacy Default Stream#

CUDA 提供 Legacy Default Stream（也叫 NULL stream 或 stream 0），当 kernel launch 或 cudaMemcpy() 不指定 stream 时使用。它是阻塞流，会与所有 blocking stream 同步——它必须等待所有 blocking stream 完成才能执行。

cpp
1cudaStream_t stream1, stream2; 2cudaStreamCreate(&stream1); // blocking 3cudaStreamCreate(&stream2); // blocking 4 5kernel1<<<grid, block, 0, stream1>>>(...); 6kernel2<<<grid, block>>>(...); // 默认流：等待 kernel1 完成 7kernel3<<<grid, block, 0, stream2>>>(...); // 等待 kernel2 完成

使用 non-blocking stream 可以避免这种隐式同步：

cpp
1cudaStream_t stream1, stream2; 2cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking); 3cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking); 4 5kernel1<<<grid, block, 0, stream1>>>(...); 6kernel2<<<grid, block>>>(...); // 不再等待 stream1 7kernel3<<<grid, block, 0, stream2>>>(...); 8// 三个 kernel 理论上可以并发执行，需要显式同步

3. Per-thread Default Stream#

从 CUDA 7 开始，每个 host 线程可以拥有独立的默认流，而非共享 Legacy Default Stream。启用方式：

  • nvcc 编译选项 --default-stream per-thread
  • 预处理宏 CUDA_API_PER_THREAD_DEFAULT_STREAM

启用后，每个线程的默认流不会与其他 stream 同步。

六、显式/隐式同步#

函数行为
cudaDeviceSynchronize()等待所有 stream 的所有操作完成
cudaStreamSynchronize(stream)等待指定 stream 的所有操作完成
cudaStreamWaitEvent(stream, event)使 stream 之后的命令等待 event 触发后才执行
cudaStreamQuery(stream)非阻塞查询 stream 是否为空

通过 cudaStreamWaitEvent() 可以建立跨 stream 的依赖关系，从而构建有向无环图（DAG）。

如果两个操作之间在默认流（NULL stream）上提交了任何 CUDA 操作，则这两个来自不同 stream 的操作无法并发——除非它们是 non-blocking stream。

为获得更好的并发性能，应遵循：

  • 所有独立操作应在依赖操作之前提交
  • 各种形式的同步应尽量推迟

七、优先级#

可以给 stream 分配优先级。使用 cudaStreamCreateWithPriority() 创建，数字越小优先级越高：

cpp
1int minPriority, maxPriority; 2cudaDeviceGetStreamPriorityRange(&minPriority, &maxPriority); 3 4cudaStream_t stream1, stream2; 5cudaStreamCreateWithPriority(&stream1, cudaStreamDefault, minPriority); // 最低优先级 6cudaStreamCreateWithPriority(&stream2, cudaStreamDefault, maxPriority); // 最高优先级

优先级是对 runtime 的提示，不保证执行顺序，主要影响 kernel launch 的调度，不一定影响内存传输。

设置优先级不会抢占已经在执行的工作。

八、CUDA Graph 简介#

对于需要反复执行相同操作序列的应用，CUDA Graph 可以显著降低 CPU 开销。CUDA Graph 分三步工作：

  1. 捕获（Capture）：首次执行时捕获整个操作图（DAG）；也可以手动用 Graph API 构建
  2. 实例化（Instantiate）：创建运行时代结构，使后续 launch 尽可能快
  3. 执行（Execute）：反复 launch 预实例化的图，CPU 开销极小

CUDA 提供了两种创建 Graph 的方式：手动 API 构建和 Stream Capture。Stream Capture 是更简洁的方式，只需将一组已有 stream 操作用 cudaStreamBeginCapture / cudaStreamEndCapture 包裹即可自动捕获整个操作序列：

cpp
1bool graphCreated = false; 2cudaGraph_t graph; 3cudaGraphExec_t instance; 4 5for (int istep = 0; istep < NSTEP; istep++) { 6 if (!graphCreated) { 7 cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal); 8 9 for (int ikrnl = 0; ikrnl < NKERNEL; ikrnl++) 10 shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d); 11 12 cudaStreamEndCapture(stream, &graph); 13 cudaGraphInstantiate(&instance, graph, NULL, NULL, 0); 14 graphCreated = true; 15 } 16 17 cudaGraphLaunch(instance, stream); 18 cudaStreamSynchronize(stream); 19}
  • 捕获期间只能提交 operation 到同一 stream；跨 stream 依赖需用 cudaStreamWaitEvent 建立
  • 实例化后的 graph 可以直接被 launch，无需重新提交 API——适合反复执行的固定 DAG
  • cudaGraphLaunch 返回后操作即入队，需 cudaStreamSynchronize 等待完成

九、异步错误处理#

Stream 中的错误可能来自 kernel 或内存传输。这些错误不会立即返回，直到 stream 被同步时才会报告。有两种方式查询错误：

cpp
1// 同步后检查 2cudaStreamSynchronize(stream); 3 4// 查看错误但不清除 5cudaError_t err = cudaPeekAtLastError(); 6if (err != cudaSuccess) { 7 printf("Error: %s\n", cudaGetErrorString(err)); 8} 9 10// 查看并清除错误 11cudaError_t err2 = cudaGetLastError();

设置环境变量 CUDA_LAUNCH_BLOCKING=1 可以让每次 kernel launch 后自动同步，方便定位具体是哪个 kernel 出了问题（但会强制串行，显著降低性能）。