CUDA C编程权威指南-第六章：流和并发

系列导航：导读 | 上一篇：第5章 共享内存和常量内存 | 下一篇：第7章 调整指令级原语

系列第 6 篇。前面几章主要围绕单次内核、单次传输；这一章讲流（Stream）与并发：用 CUDA 流把主机与设备、传输与计算、多内核并发组织起来，用事件做同步和精确计时，顺带提一下流回调、流优先级和可视化分析。

---

前言：从内核级并发到网格级并发

到目前为止，我们关注的一直是内核级并发：一个核函数由大量线程在 GPU 上并行执行，并从编程模型、执行模型和内存模型多个角度做了优化。这一章会讲到：除了这种「单任务、多线程」的并发，CUDA 还支持网格级并发——在单个设备上同时执行多个核函数启动，从而进一步提高设备利用率。

本章将系统学习如何用 CUDA 流 表达和实现网格级并发，并解决以下问题：

• 流与事件的本质：什么是流、默认流与阻塞行为，以及事件在流中的作用。
• 重叠执行：主机与设备并发、数据传输与核函数执行的重叠，以及多流并发。
• 同步机制：流同步、事件同步、避免不必要的同步，以及用事件精确计时。
• 流回调与流优先级：在流中注册主机端回调、调整流优先级（原书若涉及可查阅）。
• 可视化与多 GPU：用 NVIDIA Visual Profiler 查看执行时间线，以及多 GPU 编程入门（原书若涉及可查阅）。

第 2 章曾提到：若计算时间大于数据传输时间，可通过流来重叠计算与通信，隐藏传输延迟。本章将把这一思想具象化为可用的 API 与完整示例。

---

一、本章在全书中的位置与学习目标

1.1 为什么要学「流和并发」

前几章我们一直在做内核级并发：单次核函数内大量线程在 GPU 上并行，并从编程模型、执行模型和内存模型做了优化。这里区分指出：除了这种「单任务、多线程」的并发，CUDA 还支持网格级并发——在单个设备上同时执行多个核函数启动，从而进一步提高设备利用率。使用 CUDA 流是实现网格级并发的主要手段。此外，通过流可以重叠：(1) 主机与设备的执行；(2) 数据传输（H2D/D2H）与核函数执行。当计算时间足以掩盖部分传输时间时（第 2 章已提及），重叠能显著缩短端到端时间、提高有效带宽。因此，理解流的本质、默认流的阻塞行为、事件在同步与计时中的作用，以及如何用多流实现「传输与计算」的流水线，是写出高性能、高吞吐 CUDA 程序的关键一步；本章与第 2 章（性能模型）、第 4 章（带宽）紧密衔接。

1.2 学完本章，你应该能回答

学习目标	检验方式
理解两种并发级别（内核级 vs 网格级）及流在网格级并发中的作用	能说明「单核多线程」与「多核/多传输同时或交错执行」的区别
理解流与事件的本质：流是顺序操作队列，事件是流内标记点	能解释流内顺序执行、流间可并发的含义；事件的同步与计时用途
掌握流的创建与销毁、默认流与阻塞流的语义	能写出 cudaStreamCreate/Destroy；说明为何混用默认流会破坏多流并发
掌握在流中启动核函数与 cudaMemcpyAsync，以及固定内存的必要性	能写出 kernel<<<..., stream>>> 与 cudaMemcpyAsync(..., stream)；说明为何 H2D/D2H 需 pinned memory
理解重叠执行的三种形式：主机/设备、传输/计算、多流	能描述双流「H2D→kernel→D2H」流水线如何重叠
掌握流同步与 CUDA 事件：Synchronize、Record、ElapsedTime、StreamWaitEvent	能正确使用事件做跨流依赖与核函数计时；避免不必要的同步
了解流回调与流优先级（书中对应节）	能说出回调的限制（回调内不可调用 CUDA API）；优先级为提示性
会用 Visual Profiler 查看执行时间线以验证重叠	能根据时间线判断是否实现预期重叠、是否存在隐式同步

1.3 博客阅读导图（本章架构）

第 6 章  流和并发
│
├── 二、流与事件的本质（书 6.1 节）
│   ├── 内核级并发与网格级并发
│   ├── 流与事件的本质、异步执行与同步
│   └── 隐式同步与默认流的影响
│
├── 三、CUDA 流（书 6.2 节）
│   ├── 流的创建与销毁
│   ├── 默认流与阻塞行为 ★ 易错点
│   └── 在流中启动核函数与异步传输（固定内存）
│
├── 四、流的并发执行（书 6.3 节）★ 重难点
│   ├── 主机与设备并发
│   ├── 传输与计算的重叠
│   └── 多流并发与流水线模式
│
├── 五、流同步与 CUDA 事件（书 6.4～6.5 节）
│   ├── 流同步、避免不必要的同步
│   ├── CUDA 事件 API 与跨流依赖
│   └── 用事件精确计时
│
├── 六、流回调与流优先级（书对应节）
│   ├── 注册设备回调函数（cudaLaunchHostFunc）
│   └── 调整流优先级
│
├── 七、可视化与多 GPU 入门
│   ├── NVIDIA Visual Profiler 执行时间线
│   └── 多 GPU 与本章的衔接（简述）
│
└── 八、本章小结与重难点回顾

---

二、流与事件的本质（书 6.1 节）

本节对应书中「Understanding the nature of streams and events」，为全章奠定概念基础。

2.1 内核级并发与网格级并发

这里区分区分了 CUDA C 编程中的两个并发级别：

并发级别	含义	此前章节侧重
内核级并发	单个核函数由大量线程在 GPU 上并行执行（网格 → 块 → 线程）	第 2～5 章
网格级并发	多个核函数启动（或多个传输任务）在同一个设备上同时或交错执行	本章

在网格级并发中，多个内核（或内核与内存传输）可以同时占用设备资源。例如：当一个内核在部分 SM 上执行时，其他 SM 可以执行另一个内核；或者，当 DMA 引擎在执行主机到设备的数据传输时，已就绪的核函数可以在设备上执行。这一章会讲到：使用 CUDA 流是实现网格级并发的主要手段。

2.2 流与事件的本质

CUDA 流可以理解为一个按顺序执行的操作队列。程序把一系列操作（如内存拷贝、核函数启动）按顺序加入某个流，流中的操作将按入队顺序依次执行。不同流之间的操作则可以由运行时交错或并发执行（取决于硬件与驱动），从而实现主机与设备、传输与计算、多内核之间的重叠。

CUDA 事件是插入到流中的标记点，用于标记流中某一时刻的「完成」状态。事件可用来：(1) 同步：让某个流等待另一流中的事件再继续；(2) 计时：记录两个事件之间的时间差，从而精确测量某段核函数或传输的耗时。需要留意的是：理解流与事件的本质，是正确使用异步 API 和避免隐蔽同步的前提。

多流重叠执行时，传输与核函数可交错进行，典型时序如下（流 A 与流 B 并发，H2D 表示主机到设备传输，Kernel 表示核函数，D2H 表示设备到主机传输）：

sequenceDiagram
    participant Host
    participant StreamA as 流 A
    participant StreamB as 流 B
    participant GPU as 设备
    Note over StreamA,StreamB: 流内顺序执行，流间可重叠
    StreamA->>GPU: H2D A
    StreamB->>GPU: H2D B
    StreamA->>GPU: Kernel A
    StreamB->>GPU: Kernel B
    StreamA->>Host: D2H A
    StreamB->>Host: D2H B

2.3 异步执行与同步

CUDA 的许多 API 是异步的：调用返回时，操作可能尚未开始或尚未完成。例如：

• 核函数启动：kernel<<<grid, block>>>(...) 会立即返回，内核在设备上异步执行。
• 异步内存拷贝：cudaMemcpyAsync() 在指定流中排队传输，调用立即返回。

若要安全地使用尚未完成的操作的结果（例如读回设备数据），必须进行显式同步，例如 cudaDeviceSynchronize()、cudaStreamSynchronize(stream) 或等待某个事件。这一章会讲到：隐式同步（如调用同步版 cudaMemcpy()）会阻塞所有流的进展，在追求重叠执行时应尽量避免。

理解与体会：异步是「发起即返回」，同步是「等到完成再用结果」。流和事件把「谁在等谁」表达清楚：流内顺序由入队顺序保证，流间依赖可用事件（cudaStreamWaitEvent）表达；滥用全局同步或默认流会破坏这种表达，把本可重叠的执行串行化。隐式同步是本章最容易踩的坑——一旦在多个流之间插入了默认流上的操作或同步版 cudaMemcpy，所有流都会与该操作「对齐」，失去并发；后面「默认流与阻塞」「避免不必要的同步」会反复呼应这一点。

---

三、CUDA 流（书 6.2 节）

本节对应书中对流的创建与销毁、默认流、以及在流中启动核函数与异步传输的讲解。

3.1 流的创建与销毁

流由类型为 cudaStream_t 的句柄表示。创建与销毁流的 API 如下（与书中及 CUDA 编程指南一致）：

// 创建流
cudaError_t cudaStreamCreate(cudaStream_t *pStream);

// 销毁流（若流中仍有未完成的工作，会等待其完成后再销毁）
cudaError_t cudaStreamDestroy(cudaStream_t stream);

典型用法示例（书中风格）：

cudaStream_t stream;
CHECK(cudaStreamCreate(&stream));

// 在该流中发起核函数、cudaMemcpyAsync 等操作
// ...

CHECK(cudaStreamDestroy(stream));

注意：若在调用 cudaStreamDestroy(stream) 时设备仍在该流中执行任务，销毁会阻塞直到该流中所有工作完成。因此，若希望尽早释放资源，应在不再向该流提交新任务后、在适当时机再调用 cudaStreamDestroy。

3.2 默认流与阻塞行为 ★ 易错点

当不指定流时，核函数启动和同步版内存拷贝（如 cudaMemcpy）会使用 默认流（也称 NULL 流或 stream 0）。书中及 CUDA 编程指南指出：默认流具有阻塞语义——

• 在默认流中排入一个操作时，它会等待所有其他（阻塞）流中已排入的操作完成后才开始执行。
• 默认流中的一个操作未完成时，其他流中的操作不能开始。

因此，若混用默认流和自定义流，且自定义流是用 cudaStreamCreate() 创建的阻塞流，则默认流会与这些流相互「卡住」，无法实现多流并发。若要让自定义流与默认流并发，可以：(1) 使用 非阻塞流（cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking)）；或 (2) 在 CUDA 7 及以后，使用每线程默认流（编译选项 --default-stream per-thread 或宏 CUDA_API_PER_THREAD_DEFAULT_STREAM），使每个主机线程拥有独立的默认流。

下表归纳书中及官方文档中关于流的类型与默认流行为（便于对照原书）：

流类型	创建方式	与默认流的关系
默认流 (NULL)	不指定流即使用	与所有阻塞流互相同步，不利于多流并发
阻塞流	cudaStreamCreate(&stream)	与默认流互相同步
非阻塞流	cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking)	与默认流可并发

易错点：混用默认流与自定义流是失去多流并发的常见原因。若代码中既有 kernel<<<grid, block>>>(...)（默认流）又有 kernel<<<grid, block, 0, myStream>>>(...)，默认流会与所有阻塞流互相同步，时间线上会出现「所有流等默认流、默认流再等所有流」的串行化。对策：要么全部使用非默认流（并显式指定流），要么使用每线程默认流（--default-stream per-thread），避免在重叠路径上使用默认流或同步拷贝。

3.3 在流中启动核函数与异步传输

在指定流中启动核函数：在执行配置中增加流参数，即第四个参数（第三参数可为共享内存大小或 0）：

kernel_name<<<gridDim, blockDim, 0, stream>>>(argument list);

例如：

dim3 block(256);
dim3 grid((n + block.x - 1) / block.x);
myKernel<<<grid, block, 0, stream>>>(d_A, d_B, d_C, n);

在流中发起异步内存传输：使用 cudaMemcpyAsync，并传入目标流：

cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t count,
                            cudaMemcpyKind kind, cudaStream_t stream);

需要留意的是：要使主机与设备之间的 cudaMemcpyAsync 真正异步并与其他操作重叠，主机端缓冲区必须是固定内存（Pinned Memory）。否则，运行时可能退化为同步传输，无法重叠。固定内存通常通过 cudaMallocHost() 分配（或 cudaHostAlloc()），并在用完后用 cudaFreeHost() 释放。

下面给出一个完整示例：在单个流中完成「主机→设备」拷贝、核函数执行、「设备→主机」拷贝（与书中「基于流的向量加法」风格一致，便于对照书中代码清单 6-x）：

#include <stdio.h>
#include <stdlib.h>

#define CHECK(call)                                                                 \
    do {                                                                           \
        cudaError_t err = (call);                                                   \
        if (err != cudaSuccess) {                                                  \
            printf("Error: %s:%d, code: %d, reason: %s\n", __FILE__, __LINE__, err,\
                   cudaGetErrorString(err));                                        \
            exit(1);                                                                \
        }                                                                          \
    } while (0)

__global__ void vecAdd(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

int main(int argc, char **argv) {
    int n = 1 << 20;
    size_t nBytes = n * sizeof(float);

    // 使用固定内存以便 cudaMemcpyAsync 真正异步
    float *h_A, *h_B, *h_C;
    CHECK(cudaMallocHost((void **)&h_A, nBytes));
    CHECK(cudaMallocHost((void **)&h_B, nBytes));
    CHECK(cudaMallocHost((void **)&h_C, nBytes));
    for (int i = 0; i < n; i++) {
        h_A[i] = 1.0f; h_B[i] = 2.0f;
    }

    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((void **)&d_A, nBytes));
    CHECK(cudaMalloc((void **)&d_B, nBytes));
    CHECK(cudaMalloc((void **)&d_C, nBytes));

    cudaStream_t stream;
    CHECK(cudaStreamCreate(&stream));

    // 流中顺序：H2D -> kernel -> D2H
    CHECK(cudaMemcpyAsync(d_A, h_A, nBytes, cudaMemcpyHostToDevice, stream));
    CHECK(cudaMemcpyAsync(d_B, h_B, nBytes, cudaMemcpyHostToDevice, stream));

    dim3 block(256);
    dim3 grid((n + block.x - 1) / block.x);
    vecAdd<<<grid, block, 0, stream>>>(d_A, d_B, d_C, n);

    CHECK(cudaMemcpyAsync(h_C, d_C, nBytes, cudaMemcpyDeviceToHost, stream));

    CHECK(cudaStreamSynchronize(stream));

    // 简单验证
    float maxError = 0.0f;
    for (int i = 0; i < n; i++) {
        float e = fabsf(h_C[i] - 3.0f);
        if (e > maxError) maxError = e;
    }
    printf("max error = %f\n", maxError);

    CHECK(cudaStreamDestroy(stream));
    CHECK(cudaFree(d_A)); CHECK(cudaFree(d_B)); CHECK(cudaFree(d_C));
    CHECK(cudaFreeHost(h_A)); CHECK(cudaFreeHost(h_B)); CHECK(cudaFreeHost(h_C));
    cudaDeviceReset();
    return 0;
}

编译时可使用：nvcc -o stream_vecadd stream_vecadd.cu。本例中虽然只使用了一个流，但展示了固定内存 + 异步传输 + 流内顺序的完整写法，为下一节多流重叠打下基础。

---

四、流的并发执行（书 6.3 节）★ 重难点

本节对应书中「Overlapping CPU and GPU execution」「Overlapping kernel execution and data transfer」以及「Exploiting grid-level concurrency」，是本章提升性能的核心。

4.1 主机与设备并发

由于核函数启动和 cudaMemcpyAsync 是异步的，主机线程在发起这些操作后可以立即继续执行，不必等待设备完成。因此，可以在同一段代码中交错执行「主机上的计算」与「设备上的计算或传输」，实现主机与设备并发。这一章会讲到：当设备在运行内核或进行 DMA 传输时，CPU 可以同时处理其他任务（如准备下一批数据、做 I/O 或后处理），从而提高整体吞吐。

实现要点：(1) 使用非默认流（或每线程默认流）和 cudaMemcpyAsync，避免在关键路径上使用会阻塞的 cudaMemcpy；(2) 仅在需要结果时再调用 cudaStreamSynchronize(stream) 或 cudaDeviceSynchronize()。

4.2 传输与计算的重叠

在支持并发拷贝与执行的 GPU 上（Fermi 及以后通常支持），可以在一段数据从主机传到设备的同时，对已经传好的另一段数据在设备上执行核函数；同理，可以在核函数写回的结果正在 D2H 传输的同时，对下一批数据执行 H2D 或另一内核。书中用「双缓冲」或「多流流水线」描述这类模式：把数据和任务分块，流 1 执行「块 1 的 H2D → 块 1 的 kernel → 块 1 的 D2H」，流 2 执行「块 2 的 H2D → 块 2 的 kernel → 块 2 的 D2H」，两个流交错执行，从而重叠传输与计算。

重叠效果受硬件限制：例如，同一时刻通常只能有一个「主机→设备」和一个「设备→主机」的传输在进行（双向拷贝可并发）；内核与传输能否并发取决于设备与驱动。这一章会讲到：通信计算比若满足「计算时间足以掩盖部分传输时间」，重叠能显著减少端到端时间（与第 2 章「若计算时间 > 数据传输时间，可通过流重叠计算与通信」一致）。有效带宽公式（书中第 2、4 章已出现）可写为：

[
\text{有效带宽} = \frac{(\text{读字节数} + \text{写字节数}) \times 10^{-9}}{\text{总时间（秒）}} \ \text{GB/s}
]

在重叠场景下，「总时间」为从第一批传输开始到最后一批结果取回的时间，通常小于「所有传输时间 + 所有计算时间」的简单相加，从而有效带宽会提高。

重难点理解：重叠能否带来收益取决于计算与通信的比值。若每块数据的计算时间远小于传输时间，则重叠后总时间仍由传输主导，收益有限；若计算时间足以「盖住」一部分传输（例如流 1 的 kernel 执行时，流 2 的 H2D 在进行），则总时间会明显缩短。书中与第 2 章一致：当计算时间大于或接近传输时间时，流重叠最有价值。实践中可通过 Visual Profiler 时间线观察是否真的重叠。

4.3 多流并发

多流并发即同时使用多个流，让多个「传输 + 内核」任务在不同流中排队，由运行时在设备上交错或并发执行。需要留意的是：要实现多流之间的真正并发，必须避免在中间插入默认流上的操作（如无参的 cudaMemcpy 或 kernel<<<grid, block>>>() 未指定流），否则会引发隐式同步，导致所有流等待默认流、默认流再等待所有流，失去并发性。

下面给出一个双流重叠传输与计算的完整示例（对应书中多流、分块处理的典型模式）。将向量加法按块划分到两个流中，每个流执行：H2D → kernel → D2H；两流交错提交，以实现传输与计算的重叠。

#include <stdio.h>
#include <stdlib.h>

#define CHECK(call)                                                                 \
    do {                                                                           \
        cudaError_t err = (call);                                                   \
        if (err != cudaSuccess) {                                                  \
            printf("Error: %s:%d, code: %d, reason: %s\n", __FILE__, __LINE__, err,\
                   cudaGetErrorString(err));                                        \
            exit(1);                                                                \
        }                                                                          \
    } while (0)

__global__ void vecAdd(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

int main(int argc, char **argv) {
    int n = 1 << 20;
    size_t nBytes = n * sizeof(float);
    const int nStreams = 2;
    int segment = n / nStreams;
    size_t segBytes = segment * sizeof(float);

    float *h_A, *h_B, *h_C;
    CHECK(cudaMallocHost((void **)&h_A, nBytes));
    CHECK(cudaMallocHost((void **)&h_B, nBytes));
    CHECK(cudaMallocHost((void **)&h_C, nBytes));
    for (int i = 0; i < n; i++) {
        h_A[i] = 1.0f; h_B[i] = 2.0f;
    }

    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((void **)&d_A, nBytes));
    CHECK(cudaMalloc((void **)&d_B, nBytes));
    CHECK(cudaMalloc((void **)&d_C, nBytes));

    cudaStream_t stream[nStreams];
    for (int i = 0; i < nStreams; i++)
        CHECK(cudaStreamCreate(&stream[i]));

    for (int i = 0; i < nStreams; i++) {
        int offset = i * segment;
        CHECK(cudaMemcpyAsync(d_A + offset, h_A + offset, segBytes,
                              cudaMemcpyHostToDevice, stream[i]));
        CHECK(cudaMemcpyAsync(d_B + offset, h_B + offset, segBytes,
                              cudaMemcpyHostToDevice, stream[i]));

        dim3 block(256);
        dim3 grid((segment + block.x - 1) / block.x);
        vecAdd<<<grid, block, 0, stream[i]>>>(d_A + offset, d_B + offset,
                                              d_C + offset, segment);

        CHECK(cudaMemcpyAsync(h_C + offset, d_C + offset, segBytes,
                              cudaMemcpyDeviceToHost, stream[i]));
    }

    for (int i = 0; i < nStreams; i++)
        CHECK(cudaStreamSynchronize(stream[i]));

    float maxError = 0.0f;
    for (int i = 0; i < n; i++) {
        float e = fabsf(h_C[i] - 3.0f);
        if (e > maxError) maxError = e;
    }
    printf("max error = %f\n", maxError);

    for (int i = 0; i < nStreams; i++)
        CHECK(cudaStreamDestroy(stream[i]));
    CHECK(cudaFree(d_A)); CHECK(cudaFree(d_B)); CHECK(cudaFree(d_C));
    CHECK(cudaFreeHost(h_A)); CHECK(cudaFreeHost(h_B)); CHECK(cudaFreeHost(h_C));
    cudaDeviceReset();
    return 0;
}

书中通常会对比「单流顺序执行」与「多流重叠」的 nvprof 或 Visual Profiler 时间线：多流下 H2D、kernel、D2H 在时间轴上会出现重叠，总耗时往往明显缩短。

学习体会：多流流水线的典型模式是「按数据分块、每块绑定一个流、每流内顺序为 H2D → kernel → D2H」。流数量并非越多越好：过多流会增加调度与内存占用（每流可能需独立缓冲区或分段），建议根据设备与问题规模选择流数量（如 2～8 个流）。固定内存是前提，否则 cudaMemcpyAsync 可能退化为同步，重叠失效。

下表归纳本节涉及的三种并发模式（与书中描述一致）：

并发模式	做法简述	注意点
主机与设备	主机在发起 kernel/Async 拷贝后继续做 CPU 工作	避免在关键路径用同步 API
传输与计算	多流或双缓冲，使 H2D/D2H 与 kernel 在时间上重叠	主机缓冲需固定内存；设备需支持
多流网格级	多流中各提交独立的内核或「传输+内核」序列	避免默认流插入导致隐式同步

---

五、流同步与 CUDA 事件（书 6.4～6.5 节）

本节对应书中「Understanding synchronization mechanisms」「Avoiding unwanted synchronization」、CUDA 事件的使用以及用事件计时的内容。

5.1 流同步与避免不必要的同步

与流相关的常用同步方式包括：

• cudaDeviceSynchronize()：阻塞主机，直到所有流中此前排入的所有工作完成。这一章会讲到：这是最「重」的同步，会等待整个设备空闲，在需要最大并发时应避免在中间频繁调用。
• cudaStreamSynchronize(stream)：阻塞主机，直到指定流中此前排入的所有工作完成。其他流可继续执行，因此比 cudaDeviceSynchronize() 更细粒度。
• cudaStreamQuery(stream)：非阻塞地查询流是否已空（所有排入的操作已完成），返回 cudaSuccess 表示已空，cudaErrorNotReady 表示未空。可用于轮询或条件触发后续逻辑。

需要留意的是：避免不必要的同步（Avoiding unwanted synchronization）。例如，若在多个流之间没有数据依赖，就不要在每启动一个流后就调用 cudaStreamSynchronize 或 cudaDeviceSynchronize，否则会人为串行化本可并发的执行。应只在必须使用某流的结果时再同步该流（或等待相应事件）。原则：同步粒度越细、同步点越少，越有利于重叠；cudaStreamSynchronize(stream) 只等该流，优于 cudaDeviceSynchronize() 等全设备。

5.2 CUDA 事件

CUDA 事件是流中的逻辑标记点，类型为 cudaEvent_t。常用 API 如下（与书中及编程指南一致）：

cudaEventCreate(cudaEvent_t *event);
cudaEventDestroy(cudaEvent_t event);
cudaEventRecord(cudaEvent_t event, cudaStream_t stream);  // 将事件插入流
cudaEventSynchronize(cudaEvent_t event);                    // 阻塞直到事件完成
cudaEventQuery(cudaEvent_t event);                          // 非阻塞查询事件是否完成
cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t stop);  // 两事件间耗时(ms)

cudaEventRecord(event, stream) 把事件插入到 stream 的当前末尾；当流执行到该点时，事件被「完成」。cudaEventSynchronize(event) 会阻塞主机直到该事件完成，因此可用来「等到流执行到某一位置」再继续，而不必等整个流结束。cudaStreamWaitEvent(stream, event, 0) 则让另一流在该事件完成之前不执行其后排入的操作，从而建立流与流之间的依赖。

5.3 用事件计时

这一章会讲到：用 CPU 计时器 给核函数计时时，必须配合 cudaDeviceSynchronize() 或 cudaStreamSynchronize()，否则测到的只是「启动开销」。CUDA 事件可以在设备上精确记录时间戳，适合对单个流内的核函数或传输做计时，且不需要阻塞所有流。

典型模式：在流中在核函数前后各插入一个事件，然后同步流（或至少同步「停止」事件），再用 cudaEventElapsedTime 得到两次记录之间的毫秒数。公式上，若记 ( t_{\mathrm{start}} )、( t_{\mathrm{stop}} ) 为两事件对应的 GPU 时间戳，则：

[
\text{耗时（ms）} = t_{\mathrm{stop}} - t_{\mathrm{start}}
]

cudaEventElapsedTime 返回的即为该差值（单位：毫秒）。注意：两个事件必须在同一流中，且 start 必须在 stop 之前被记录，否则行为未定义。

下面给出用事件对指定流中的核函数计时的完整示例（与书中「用事件计时」代码风格一致）：

#include <stdio.h>
#include <stdlib.h>

#define CHECK(call)                                                                 \
    do {                                                                           \
        cudaError_t err = (call);                                                   \
        if (err != cudaSuccess) {                                                  \
            printf("Error: %s:%d, code: %d, reason: %s\n", __FILE__, __LINE__, err,\
                   cudaGetErrorString(err));                                        \
            exit(1);                                                                \
        }                                                                          \
    } while (0)

__global__ void vecAdd(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

int main(int argc, char **argv) {
    int n = 1 << 20;
    size_t nBytes = n * sizeof(float);

    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((void **)&d_A, nBytes));
    CHECK(cudaMalloc((void **)&d_B, nBytes));
    CHECK(cudaMalloc((void **)&d_C, nBytes));

    cudaStream_t stream;
    cudaEvent_t start, stop;
    CHECK(cudaStreamCreate(&stream));
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));

    dim3 block(256);
    dim3 grid((n + block.x - 1) / block.x);

    CHECK(cudaEventRecord(start, stream));
    vecAdd<<<grid, block, 0, stream>>>(d_A, d_B, d_C, n);
    CHECK(cudaEventRecord(stop, stream));

    CHECK(cudaStreamSynchronize(stream));

    float elapsedMs = 0.0f;
    CHECK(cudaEventElapsedTime(&elapsedMs, start, stop));
    printf("Kernel elapsed time = %.3f ms\n", elapsedMs);

    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));
    CHECK(cudaStreamDestroy(stream));
    CHECK(cudaFree(d_A)); CHECK(cudaFree(d_B)); CHECK(cudaFree(d_C));
    cudaDeviceReset();
    return 0;
}

下表归纳本节涉及的同步与事件 API（便于与书中表 6-x 对照）：

API / 行为	作用
cudaDeviceSynchronize()	等待所有流上此前的工作完成
cudaStreamSynchronize(s)	等待指定流此前的工作完成
cudaStreamQuery(s)	非阻塞查询流是否已空
cudaEventRecord(e,s)	在流 s 中插入事件 e
cudaEventSynchronize(e)	阻塞直到事件 e 完成
cudaStreamWaitEvent(s,e,0)	流 s 在 e 完成前不执行其后操作
cudaEventElapsedTime(&t, e0, e1)	得到两事件间耗时（ms）

---

六、流回调与流优先级（书对应节）

本节对应书中「Registering device callback functions」与「Adjusting stream priorities」的内容。

6.1 流回调

CUDA 允许在流中插入主机端回调：当流执行到该点时，在主机上执行一个指定的函数。可用于在「某段设备工作完成后」自动触发主机逻辑（如解压下一批数据、更新 UI），而无需轮询或单独开线程。书中在「注册设备回调函数」一节中，若使用旧版 API 会提到 cudaStreamAddCallback；当前 CUDA 编程指南推荐使用 cudaLaunchHostFunc（cudaStreamAddCallback 已标记为弃用）：

cudaError_t cudaLaunchHostFunc(cudaStream_t stream, void (*fn)(void *), void *userData);

回调函数签名为 void hostFn(void *userData)。限制：回调函数内不得调用任何 CUDA API，否则可能死锁或未定义行为。需要留意的是：回调在主机线程上执行，流在该回调执行期间被视为「空闲」，可安全访问与该流关联的、已完成的设备结果。典型用法是：在某个流中排入「核函数 → 主机回调」，主机端无需轮询即可在该流中核函数完成后执行后续逻辑（如准备下一批数据或写回文件）。

6.2 流优先级

在支持流优先级的设备上（如 Pascal 及更新架构），可以用 cudaStreamCreateWithPriority 创建带优先级的流，运行时在调度时会优先执行高优先级流中的工作（作为提示，不保证严格顺序）。书中若涉及，会给出类似下面的用法：

int leastPriority, greatestPriority;
cudaDeviceGetStreamPriorityRange(&leastPriority, &greatestPriority);
// greatestPriority 为高优先级（数值上可能更小，依实现而定）
cudaStream_t streamHigh;
cudaStreamCreateWithPriority(&streamHigh, cudaStreamNonBlocking, greatestPriority);

流优先级适合将「延迟敏感」的内核与「吞吐型」的内核分开到不同流并赋予不同优先级，以在共享设备上获得更可预测的延迟（书中表 6-x 若有优先级范围说明，可在此引用）。书中「调整流优先级」一节会说明：并非所有设备都支持流优先级，可通过 cudaDeviceGetStreamPriorityRange 查询；若设备不支持，创建带优先级的流可能回退为普通非阻塞流。

---

七、可视化与多 GPU 入门（书对应节）

本节对应书中「Displaying application execution timelines with the NVIDIA Visual Profiler」；多 GPU 深入内容在原书中为后续章节（如第 9 章），此处仅作与本章知识的衔接简述。

7.1 使用 Visual Profiler 查看执行时间线

可使用 NVIDIA Visual Profiler（或 Nsight Systems）查看应用程序的执行时间线。在时间线视图中，可以直观看到不同流上的 H2D、D2H、核函数在时间轴上的排列与重叠情况，从而判断是否实现了预期的传输与计算重叠、是否存在不必要的同步或默认流阻塞。通过时间线可以快速定位「本可并发的流被默认流或同步调用串行化」的问题。建议在运行示例时使用 nvprof 或 nsys profile 采集时间线，再用 Visual Profiler 打开分析。

实践建议：优化流与并发时，先用 profiler 看时间线——若看到各流的 H2D/kernel/D2H 几乎完全错开、没有重叠，多半是默认流或某处同步导致；若重叠明显，再结合有效带宽公式评估是否达到预期。

7.2 多 GPU 与本章的衔接（简述）

原书多 GPU 编程在后续章节（如第 9 章）展开。与本章直接相关的是：每个设备有各自的默认流和显式创建的流；在多设备程序中，分配与启动前需通过 cudaSetDevice(deviceId) 指定当前设备，流和事件都是「每设备」的。理解「单设备上的多流」后，多 GPU 时即为「每设备一套流」，再叠加设备间数据传输与 P2P 等主题。

下表归纳本章涉及的主要 API 与概念（与书中描述一致，便于自检）：

类别	内容
流	cudaStreamCreate / cudaStreamDestroy；默认流、阻塞流、非阻塞流；kernel<<<..., stream>>>；cudaMemcpyAsync(..., stream)
同步	cudaDeviceSynchronize；cudaStreamSynchronize；cudaStreamQuery；隐式同步（慎用默认流与同步拷贝）
事件	cudaEventCreate / cudaEventDestroy；cudaEventRecord；cudaEventSynchronize / cudaEventQuery；cudaEventElapsedTime；cudaStreamWaitEvent
主机	固定内存 cudaMallocHost / cudaFreeHost；流回调 cudaLaunchHostFunc；流优先级 cudaStreamCreateWithPriority
并发	主机与设备重叠；传输与计算重叠；多流网格级并发；避免不必要的同步

---

八、本章小结与重难点回顾

8.1 知识小结（与书中对应）

• 两种并发级别：内核级并发（单核函数多线程）与网格级并发（多核函数/多传输同时或交错执行）；本章重点为网格级并发，通过 CUDA 流 实现。
• 流的本质：流是按顺序执行的操作队列；不同流间可并发。默认流会与阻塞流互相同步，要实现多流并发需使用非阻塞流或每线程默认流，并避免在中间插入默认流操作。
• 异步 API：核函数启动、cudaMemcpyAsync 为异步；主机端参与 H2D/D2H 的缓冲区需固定内存（cudaMallocHost），否则异步传输可能退化为同步。
• 重叠执行：主机与设备可并发；传输与计算可在多流中重叠；多流提交「H2D → kernel → D2H」的流水线可提高有效带宽、缩短总时间。
• 同步：cudaStreamSynchronize(stream) 比 cudaDeviceSynchronize() 粒度更细；应避免不必要的同步以保留并发性。
• 事件：用于在流内打点、跨流依赖（cudaStreamWaitEvent）以及精确计时（cudaEventRecord + cudaEventElapsedTime）。
• 流回调与流优先级：书中「Registering device callback functions」对应 cudaStreamAddCallback / cudaLaunchHostFunc（回调内不可调用 CUDA API）；「Adjusting stream priorities」对应 cudaStreamCreateWithPriority（设备支持时）。
• 可视化：用 NVIDIA Visual Profiler 查看执行时间线以验证重叠、定位隐式同步。多 GPU 时每设备独立流与 cudaSetDevice 是基础（深入见原书后续章节）。

8.2 重难点速查

重难点	要点
默认流与阻塞	默认流与所有阻塞流互相同步；混用默认流会串行化多流。对策：非默认流 + 非阻塞流，或 per-thread 默认流。
固定内存与异步传输	cudaMemcpyAsync 要真正异步、参与重叠，主机缓冲区必须是 pinned（cudaMallocHost）；否则可能退化为同步拷贝。
隐式同步	同步版 cudaMemcpy、默认流上的操作会触发隐式同步，阻塞所有流。追求重叠时避免在关键路径使用。
传输与计算重叠	多流流水线：每流 H2D → kernel → D2H；重叠效果取决于计算/通信比与设备能力；流数量需权衡（如 2～8）。
事件的两大用途	(1) 同步：cudaEventSynchronize、cudaStreamWaitEvent 建立流间依赖；(2) 计时：同流内两事件 + cudaEventElapsedTime 得 GPU 侧耗时。
避免不必要的同步	只在必须使用某流结果时再同步该流；细粒度同步优于全设备同步。

8.3 学习思考

• 与第 2 章的衔接：第 2 章提到「若计算时间大于数据传输时间，可通过流重叠计算与通信」；本章把这一思想落实为流、固定内存、异步拷贝与多流流水线，并用有效带宽公式衡量。理解「计算/通信比」有助于判断何时值得上多流、何时收益有限。
• 与第 4 章的关系：第 4 章关注单次内核的全局内存带宽与合并访问；本章关注「多次传输与多次内核」在时间轴上的重叠，从而从系统层面提高有效带宽。两者结合：单内核访存优化 + 流级重叠，才能把设备吞吐榨干。
• 实践建议：写多流代码时，(1) 先保证单流逻辑正确（固定内存 + 异步 API + 指定流）；(2) 再扩展到多流，注意不要插入默认流或同步拷贝；(3) 用 Visual Profiler 看时间线确认重叠；(4) 事件计时优先于 CPU 计时，且两事件须在同一流内。

下表归纳本章涉及的主要 API 与概念（与书中描述一致，便于自检）：

类别	内容
流	cudaStreamCreate / cudaStreamDestroy；默认流、阻塞流、非阻塞流；kernel<<<..., stream>>>；cudaMemcpyAsync(..., stream)
同步	cudaDeviceSynchronize；cudaStreamSynchronize；cudaStreamQuery；隐式同步（慎用默认流与同步拷贝）
事件	cudaEventCreate / cudaEventDestroy；cudaEventRecord；cudaEventSynchronize / cudaEventQuery；cudaEventElapsedTime；cudaStreamWaitEvent
主机	固定内存 cudaMallocHost / cudaFreeHost；流回调 cudaLaunchHostFunc；流优先级 cudaStreamCreateWithPriority
并发	主机与设备重叠；传输与计算重叠；多流网格级并发；避免不必要的同步

---

下一章预告

在下一篇博客中，我们将进入第 7 章：调整指令级原语：

• 线程束洗牌再探：在第 5 章入门基础上，深入 __shfl_*_sync 的各类变体与掩码语义
• 线程束表决函数：__ballot_sync、__all_sync、__any_sync
• 活跃掩码与 warp 同步：__activemask()、__syncwarp()
• 循环展开与快速数学：指令级优化与精度取舍

从「流与并发」到「单次内核内部的指令级调优」，是进一步榨取 GPU 性能的必经之路。

---

本章自测

1. 内核级并发与网格级并发的区别是什么？流主要实现哪一种？
2. 为何多流重叠执行时不能使用同步版 cudaMemcpy()？应改用何 API？
3. 用事件做核函数计时的典型步骤是什么（Record → ? → ElapsedTime）？

答案与解析

1. 内核级并发指单次核函数内大量线程并行；网格级并发指多次核函数启动或传输在设备上同时/交错执行。流主要实现网格级并发（多流内操作可重叠）。
2. 同步版 cudaMemcpy() 会阻塞主机直到传输完成，且会与所有流同步，导致多流无法重叠。应使用 cudaMemcpyAsync(..., stream) 并在流中排队，且主机端需使用锁页内存。
3. 在核函数前 cudaEventRecord(start, stream)，核函数后 cudaEventRecord(stop, stream)，然后 cudaEventSynchronize(stop)，最后 cudaEventElapsedTime(&ms, start, stop) 得到毫秒数。

---

系列导航：导读 | 上一篇：第5章 共享内存和常量内存 | 下一篇：第7章 调整指令级原语

---

本文为「CUDA C编程权威指南」系列博客第 6 篇，共 10 章。基于《Professional CUDA C Programming》by John Cheng, Max Grossman, Ty McKercher。