【CUDA】（六）流和并发

（六）流和并发

本篇笔记参考如下：

https://blog.csdn.net/DevPath/article/details/155607318

在本章前，我们的关注点可能仅限于内核级的并发，在此级别的并发中，单一的任务或内核被GPU的多个线程并行执行。本章将研究网格级的并发。在网格级并发中，多个内核在同一设备上同时执行，这往往会让设备利用率更好。

6.1 流和事件概述

CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行。流能封装这些操作，保持操作的顺序，允许操作在流中排队，并使它们在先前的所有操作之后执行，并且可以查询排队操作的状态。

因为所有在CUDA流中排队的操作都是异步的，所以在主机与设备系统中可以重叠执行其他操作。在同一时间内将流中排队的操作与其他有用的操作一起执行，可以隐藏执行那些操作的开销。

CUDA编程的一个典型模式是以下形式：

1.将输入数据从主机移到设备上。

2.在设备上执行一个内核。

3.将结果从设备移回主机中。

在这些情况下，可

以完全隐藏CPU和GPU之间的通信延迟。通过将内核执行和数据传输调度到不同的流中，这些操作可以重叠，程序的总运行时间将被缩短。流在CUDA的API调用粒度上可实现流水线或双缓冲技术。

CUDA的API函数一般可以分为同步或异步。具有同步行为的函数会阻塞主机端线程，直到它们完成。具有异步行为的函数被调用后，会立即将控制权归还给主机。

6.1.1 CUDA流

所有的CUDA操作（包括内核和数据传输）都在一个流中显式或隐式地运行。流分为两种类型：

·隐式声明的流（空流）

·显式声明的流（非空流）

如果没有显式地指定一个流，那么内核启动和数据传输将默认使用空流。前面章节所使用的例子都是空流或默认流。

非空流可以被显式地创建和管理。如果想要重叠不同的CUDA操作，必须使用非空流。

基于流的异步的内核启动和数据传输支持以下类型的粗粒度并发：

• 重叠主机计算和设备计算
• 重叠主机计算和主机与设备间的数据传输
• 重叠主机与设备间的数据传输和设备计算
• 并发设备计算

默认流的代码来说

cudaMemcpy(..., cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(...);
cudaMemcpy(..., cudaMemcpyDeviceToHost);

1.设备视角：严格的“排队机制”

在 GPU（设备）看来，这三行代码就像是排队买票。

• 默认流（Default Stream）：所有的任务都被丢进了同一个队列。
• 顺序执行：哪怕内核启动（kernel<<<...>>>）很快就发出了指令，GPU 也会严格按照“拷贝数据进 GPU -> 运行内核 -> 拷贝数据回 CPU”的顺序执行。GPU 不会“跳级”，也不会管 CPU 现在在忙什么。

2. 主机视角：阻塞与非阻塞

这是理解 CUDA 性能优化的关键点。CPU 在执行这三行代码时，态度是完全不同的：

操作类型	行为模式	对 CPU 的影响
cudaMemcpy	同步 (Blocking)	阻塞。CPU 会停在这里，盯着进度条看，直到数据传完。这期间 CPU 什么都干不了（空闲）。
kernel<<<…>>>	异步 (Non-blocking)	瞬发。CPU 只是“下达命令”给 GPU，下完命令立刻执行下一行代码，不会等待 GPU 算完。

数据传输也可以被异步发布，但是必须显式地设置一个CUDA流来装载它们。

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

请注意附加的流标识符作为第五个参数。默认情况下，流标识符被设置为默认流。这个函数与主机是异步的，所以调用发布后，控制权将立即返回到主机。

cudaError_t cudaStreamCreate(cudaStream_t* pStream);

cudaStreamCreate创建了一个可以显式管理的非空流。之后，返回到pStream中的流就可以被当作流参数供cudaMemcpyAsync和其他异步CUDA的API函数来使用。

举例：

cudaStream_t stream;
cudaStreamCreate(&stream);

// 使用异步传输
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

// 在指定流中启动内核
kernel<<<grid, block, 0, stream>>>(...);

cudaStreamDestroy(stream);

执行异步数据传输时，必须使用固定（非分页）的主机内存。可以使用 cudaMallocHost 或 cudaHostAlloc 来分配这种特殊内存。

原理与风险：

• 确保物理位置不变：固定内存确保数据在 CPU 内存中的物理地址在程序生命周期内保持不变。
• 避免未定义行为：如果使用普通的“可分页”内存进行异步转移，操作系统可能会在传输过程中移动数据块的物理位置，而 CUDA 硬件此时仍在尝试访问旧地址，这会导致程序崩溃或数据损坏（未定义行为）。

举例如下：

// 1. 分配固定内存 (Pinned Memory)
float *h_data;
cudaMallocHost(&h_data, size); //

// 2. 创建两个 CUDA 流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1); //
cudaStreamCreate(&stream2);

// 3. 流水线操作 (将数据分为两半)
int half = size / 2;

// --- 处理第一部分数据 (流1) ---
cudaMemcpyAsync(d_data, h_data, half, ..., stream1); 
kernel<<<grid, block, 0, stream1>>>(d_data, ...); // 在流1中启动

// --- 处理第二部分数据 (流2) ---
// 当流1在计算时，流2可以同时开始传输数据！
cudaMemcpyAsync(d_data + half, h_data + half, half, ..., stream2);
kernel<<<grid, block, 0, stream2>>>(d_data + half, ...);

// 4. 等待所有任务完成并销毁
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

普通模式（默认流 + 可分页内存）

CPU 会在 cudaMemcpy 处被阻塞，且 GPU 在同一时间只能做一件事：

1. 拷贝数据 1 (CPU 等待)
2. 计算数据 1 (CPU 恢复执行)
3. 拷贝数据 2 (CPU 再次等待)

流水线模式（异步流 + 固定内存）

因为固定内存确保了物理地址不变，操作系统不会在传输中途捣乱，所以 GPU 可以在计算第一组数据的同时，拷贝第二组数据。

下图展示了使用3个流的CUDA操作。数据传输和内核计算都是均匀分布在3个并发流中的。

6.1.2 流调度

从概念上讲，所有的流可以同时运行。但是，当将流映射到物理硬件时并不总是这样的。

6.1.2.1 虚假的依赖关系

在 Blackwell 架构中，CUDA 流（Streams）的执行不再仅仅依赖于传统的硬件队列，而是演变为一种高度动态的资源编排机制。

1. 消除“虚假依赖”：从 Hyper-Q 到分布式任务分布引擎 (DTE)

• Fermi 架构限制：采用单一硬件工作队列（Single Hardware Work Queue），所有流的任务在硬件入口处被“序列化”。即使是不同流中的非相关任务，也会因为队列顶部的阻塞操作而产生硬件层面的线性等待。
• Blackwell 架构优化：Blackwell 彻底抛弃了 Kepler 时代仅有的 32 个硬件工作队列（Hyper-Q）限制。它引入了更先进的分布式任务分布引擎 (Distributed Task Distribution Engine)，能够以极高的吞吐量将数以万计的并发网格分发至 GPC（图形处理集群）。这意味着在硬件底层，Blackwell 几乎实现了全透明的流间解耦，虚假依赖在微架构层面被物理消除。

6.1.2.2 Hyper-Q技术

Kepler GPU家族中的Hyper-Q技术，使用多个硬件工作队列，从而减少了虚假的依赖关系。Hyper-Q技术通过在主机和设备之间维持多个硬件管理上的连接，允许多个CPU线程或进程在单一GPU上同时启动工作。

在 Blackwell 中，异步计算（Async Compute）已经从“多个流排队”进化到了“硬件原生重叠”。

• 并发量级：Blackwell 支持极高数量的并发网格（Grids）执行，远超 Fermi 的 16 路。
• 动态切分：它能够动态地将 SM（流处理器）资源切分给不同的流。如果流 A 的任务只占用了 30% 的资源，剩下的 70% 会立刻被分配给流 B，而不会产生任何微小的“虚假等待”。

流调度不再局限于单块 GPU 内部。通过第五代 NVLink，多块 GPU 之间的流可以被看作是一个巨大的逻辑资源池。

• 跨卡并发：可以像在单卡上写流一样，在 72 块 GPU 组成的集群（GB200 NVL72）中调度流，硬件会自动处理卡间的依赖关系。

特性	Fermi (图 6-2)	Kepler (图 6-3)	Blackwell (最新)
工作队列	单一硬件队列 (单流水线)	32 个硬件工作队列 (Hyper-Q)	海量并行队列 + 动态资源切分
虚假依赖	严重：阻塞操作会卡住后续所有流	大幅减少：每个流有独立通道	基本消除：硬件原生支持乱序发射与抢占
并发能力	16 路网格并发	32 路流并发	数千个并发线程块，跨卡协同调度

6.1.3 流的优先级

在默认情况下，所有创建的流具有相同的优先级。但在支持流优先级的硬件上，CUDA 允许在创建流时指定其在硬件调度器中的相对权重。

• 硬件分发权重：优先级不代表“绝对抢占”，而是代表“调度权重”。当 GPU 资源（SM）空闲时，硬件任务分布引擎（DTE）会优先从高优先级队列中抽取线程块（Thread Blocks）进行分发。
• 并发与交错：高优先级流并不会完全停止低优先级流，而是通过在 SM 指令级调度中获得更高频次的发射机会，从而缩短任务的整体周转时间（Turnaround Time）。

由于不同架构（从 Fermi、Kepler 到 Blackwell）支持的优先级阶数不同，必须先通过 API 查询当前设备的有效范围：

int priority_low, priority_high;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// 通常结果：priority_low = 0 (最低), priority_high = -1 或 -2 (最高)

创建指定优先级的流

在创建非默认流时，使用 cudaStreamCreateWithPriority 代替 cudaStreamCreate：

cudaStream_t high_prio_stream;
// 创建一个最高优先级的流
cudaStreamCreateWithPriority(&high_prio_stream, cudaStreamNonBlocking, priority_high)

在 Blackwell 架构中，流优先级与硬件调度器的结合变得更加紧密和智能：

• 微秒级抢占 (Microsecond Preemption)：Blackwell 支持更细粒度的硬件抢占。如果一个高优先级流的内核启动，硬件调度器可以在更短的时间切片内暂停低优先级的线程块，切换上下文，以确保高优先级任务的确定性延迟。
• 异步拷贝引擎（ACE）优先级：在使用异步传输时，固定内存的传输请求也会根据流优先级进行排队。Blackwell 的多个异步拷贝引擎可以根据流的优先级标识，优先处理关键流的数据搬运，从而在通信层面也消除延迟。
• 计算图 (CUDA Graphs) 的集成：在 Blackwell 上，开发者可以将一整套带有优先级的流操作捕获为计算图。硬件调度器在执行图任务时，会自动优化高优先级路径上的资源分配，避免在复杂的任务依赖网络中出现“优先级反转”现象。

我们可以通过以下代码得到设备的优先级范围

int priority_low, priority_high;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);

采用cudaStreamCreateWithPriority创建非默认流

cudaStream_t high_prio_stream;
// 创建一个最高优先级的流
cudaStreamCreateWithPriority(&high_prio_stream, cudaStreamNonBlocking, priority_high);

完整代码如下：

#include <stdio.h>
#include <cuda_runtime.h>

// 一个简单的耗时核函数，用于模拟繁重的计算任务
__global__ void heavy_compute_kernel(int iterations) {
    double val = 0.0;
    for (int i = 0; i < iterations; i++) {
        val = sin(val) + cos(val);
    }
}

int main() {
    int priority_low, priority_high;
    
    // 1. 获取当前设备的优先级范围
    // 数值越小，优先级越高
    cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
    printf("设备支持的优先级范围: 最低=%d, 最高=%d\n", priority_low, priority_high);

    // 2. 创建不同优先级的非默认流
    cudaStream_t st_high, st_low;
    cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
    cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);

    // 创建计时事件
    cudaEvent_t start_h, stop_h, start_l, stop_l;
    cudaEventCreate(&start_h); cudaEventCreate(&stop_h);
    cudaEventCreate(&start_l); cudaEventCreate(&stop_l);

    printf("启动核函数并行竞争...\n");

    // 3. 在低优先级流中启动核函数 (先启动)
    // 根据描述，由于核函数启动是异步的，主机会立即返回并执行下一行
    cudaEventRecord(start_l, st_low);
    for(int i=0; i<10; i++) {
        heavy_compute_kernel<<<100, 1024, 0, st_low>>>(1000000); //
    }
    cudaEventRecord(stop_l, st_low);

    // 4. 在高优先级流中启动核函数 (后启动，尝试插队)
    cudaEventRecord(start_h, st_high);
    for(int i=0; i<10; i++) {
        heavy_compute_kernel<<<100, 1024, 0, st_high>>>(1000000); //
    }
    cudaEventRecord(stop_h, st_high);

    // 5. 等待所有流完成同步
    cudaStreamSynchronize(st_high);
    cudaStreamSynchronize(st_low);

    // 计算执行时间
    float ms_h, ms_l;
    cudaEventElapsedTime(&ms_h, start_h, stop_h);
    cudaEventElapsedTime(&ms_l, start_l, stop_l);

    printf("高优先级流耗时: %.2f ms\n", ms_h);
    printf("低优先级流耗时: %.2f ms\n", ms_l);

    if (ms_h < ms_l) {
        printf("验证成功：高优先级任务在硬件层面获得了优先调度。\n");
    }

    // 6. 销毁资源
    cudaStreamDestroy(st_high);
    cudaStreamDestroy(st_low);
    cudaEventDestroy(start_h); cudaEventDestroy(stop_h);
    cudaEventDestroy(start_l); cudaEventDestroy(stop_l);

    return 0;
}

编译并执行代码stream_priority_test.cu

nvcc -O3 stream_priority_test.cu -o stream_priority_test
./stream_priority_test

得到以下结果

设备支持的优先级范围: 最低=0, 最高=-5
启动核函数并行竞争...
高优先级流耗时: 0.03 ms
低优先级流耗时: 218.44 ms
验证成功：高优先级任务在硬件层面获得了优先调度。

6.1.4 CUDA事件

CUDA 事件 (Events) 是轻量级的同步和计时对象。CUDA 事件主要服务于以下三个目的：

• 性能测量（精确定时）：事件可以记录在 GPU 流中执行到特定点的时间戳。这比 CPU 计时器更准确，因为 CUDA 内核启动是异步的，CPU 计时器往往只能测到指令发射的时间，而非执行完成的时间。
• 流间同步（跨流协调）：事件可以作为不同流之间的“信号灯”。例如，流 A 可以等待流 B 中的某个特定事件完成后再继续执行，从而实现复杂的任务依赖关系。
• 进度跟踪：主机（CPU）可以通过查询事件状态来判断 GPU 是否已经运行到了程序中的某个特定点。

事件的主要API：

步骤	API 函数	说明
创建	cudaEventCreate(&event)	初始化一个事件对象。
记录	cudaEventRecord(event, stream)	在指定的流中插入一个标记。当 GPU 执行到此处时，事件状态变为“已发生”。
等待	cudaEventSynchronize(event)	阻塞主机：CPU 会停在此处，直到 GPU 完成该事件。
计算	cudaEventElapsedTime(&ms, start, stop)	计算两个事件之间的时间差（单位：毫秒），分辨率约为 0.5 微秒。
销毁	cudaEventDestroy(event)	释放事件占用的资源。

在 Blackwell 架构中，事件的调度与同步受益于硬件层面的多项改进：

• 超低延迟同步：Blackwell 进一步优化了硬件任务分布引擎，使得事件的记录和跨流等待的硬件开销更低，能够支持更高频率的任务切换。
• 计算图加速：在 Blackwell 上，建议将事件同步逻辑封装进 CUDA Graphs。硬件可以直接解析图中的事件依赖，减少了由于主机端参与同步而产生的指令延迟。
• 内存一致性优化：Blackwell 架构在处理事件同步时，能更高效地确保内存写入对其他流或主机可见，从而减少了为保证数据一致性而产生的等待周期。

编写代码如下cuda_event_test.cu

#include <stdio.h>
#include <cuda_runtime.h>

// 模拟计算任务
__global__ void simple_kernel(int iterations) {
    double val = threadIdx.x;
    for (int i = 0; i < iterations; i++) {
        val = sqrt(val + 5.0) * tan(0.1);
    }
}

int main() {
    // 1. 定义流与事件句柄
    cudaStream_t stream_A, stream_B;
    cudaEvent_t start_event, stop_event, sync_event;

    // 2. 初始化流
    cudaStreamCreate(&stream_A);
    cudaStreamCreate(&stream_B);

    // 3. 初始化事件
    // 使用 cudaEventDefault 或用 cudaEventDisableTiming 优化非计时事件
    cudaEventCreate(&start_event);
    cudaEventCreate(&stop_event);
    cudaEventCreateWithFlags(&sync_event, cudaEventDisableTiming); // 仅用于同步，降低开销

    printf("开始执行流间同步与计时实验...\n");

    // 4. 记录起始时间戳
    cudaEventRecord(start_event, stream_A);

    // 在流 A 中启动预处理内核
    simple_kernel<<<128, 512, 0, stream_A>>>(100000); 

    // 5. 设置同步点：在流 A 中记录同步事件
    // 只有当流 A 的任务执行到这里，sync_event 才会变为“完成”状态
    cudaEventRecord(sync_event, stream_A);

    // 6. 流间等待：让流 B 等待流 A 的信号
    // 这是硬件级的等待，不阻塞主机 CPU
    cudaStreamWaitEvent(stream_B, sync_event, 0); 

    // 流 B 在接收到 sync_event 信号后才会启动后续内核
    simple_kernel<<<128, 512, 0, stream_B>>>(200000);

    // 7. 记录结束时间戳
    // 注意：记录在流 B 中，以测算整个链条的完成时间
    cudaEventRecord(stop_event, stream_B);

    // 8. 主机同步：等待 GPU 完成所有操作
    cudaEventSynchronize(stop_event);

    // 9. 计算并输出耗时
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start_event, stop_event); //

    printf("全流程 GPU 硬件耗时: %.4f ms\n", milliseconds);

    // 10. 清理资源
    cudaEventDestroy(start_event);
    cudaEventDestroy(stop_event);
    cudaEventDestroy(sync_event);
    cudaStreamDestroy(stream_A);
    cudaStreamDestroy(stream_B);

    return 0;
}

代码实现的流程如下:

步骤	流 A (stream_A)	流 B (stream_B)	主机 (CPU)
1	记录 Start 事件	(空闲)	发射指令后立即返回
2	执行计算 (10w 次迭代)	(空闲)	执行后续代码
3	记录 Sync 事件	(等待信号)	执行后续代码
4	(执行完毕或继续其他)	接收信号并启动计算	执行后续代码
5	(空闲)	记录 Stop 事件	cudaEventSynchronize 等待结束

6.1.5 流同步

在非默认流中，所有的操作对于主机线程都是非阻塞的，因此会遇到需要在同一个流中运行主机和运算操作同步的情况。

从主机的角度来说，CUDA操作可以分为两大类：

·内存相关操作

·内核启动

在 CUDA 中，所有的非空流默认都是 阻塞流。这意味着它们虽然在 CPU 上是异步的，但在 GPU 内部的执行顺序会受到默认流的影响。

流类型	主机端（CPU）表现	设备端（GPU）表现
非空流（自定义流）	非阻塞：函数启动后立即返回 CPU。	受限并行：其操作可能被空流中的任务阻塞。
空流（默认流）	阻塞：通常会导致同步行为。	霸道执行：它不仅阻塞自己，还会拦截并阻塞所有的“阻塞型”非空流。

1.显式同步

显式同步是在代码中明确插入的指令，用于强制程序运行到某一特定点。

• 设备级同步：调用 cudaDeviceSynchronize()。这会阻塞主机 CPU，直到 GPU 上所有流的任务全部完成。
• 流级同步：调用 cudaStreamSynchronize(stream)。CPU 仅等待指定的流，不干扰其他流。
• 事件级同步：调用 cudaEventSynchronize(event)。CPU 等待某个特定的标记点被 GPU 执行到。
• 设备端流同步：调用 cudaStreamWaitEvent(stream, event)。这是一种“非阻塞”同步，CPU 不会等待，而是由 GPU 硬件调度器在内部让一个流等待另一个流。

2.隐式同步

隐式同步通常是某些操作的副作用，由 CUDA 运行时自动触发。即使没有写同步函数，这些操作也会强制让 GPU 和 CPU 进行“握手”。

• 默认流（空流）的干扰：
  • 在默认流中启动任何操作都会阻塞所有非空流（阻塞型流）的执行。
  • 默认流会强制序列化设备上的所有操作。
• 内存分配与设置：调用 cudaMalloc 或 cudaMemset 等函数时，通常会触发隐式同步。
• 数据传输：在默认流中使用 cudaMemcpy 传输数据是同步的，主机会被强制闲置直到传输完成。
• L1/共享内存配置修改：在切换核函数的共享内存配置时，可能会触发同步。

6.2 并发内核执行

接下来用几个例子来演示一下。第一个示例演示了如何使用多个流并发运行多个核函数。这个简单的例子将介绍并发内核执行的几个基本问题，包括以下几个方面：

·使用深度优先或广度优先方法的调度工作

·调整硬件工作队列

·在Kepler设备和Fermi设备上避免虚假的依赖关系

·检查默认流的阻塞行为

·在非默认流之间添加依赖关系

·检查资源使用是如何影响并发的

6.2.1 非空流中的并发内核

我们采用simpleHyperqDepth.cu代码

核心代码：

// 深度优先提交 (Depth-First)
for (int i = 0; i < n_streams; i++)
{
    kernel_1<<<grid, block, 0, streams[i]>>>();
    kernel_2<<<grid, block, 0, streams[i]>>>();
    kernel_3<<<grid, block, 0, streams[i]>>>();
    kernel_4<<<grid, block, 0, streams[i]>>>();
}

流程如下（假设 n_streams = 2）：

1. 循环 i = 0 (处理第一个流)：
   • 将 kernel_1 放入 streams[0]。
   • 将 kernel_2 放入 streams[0]。
   • 将 kernel_3 放入 streams[0]。
   • 将 kernel_4 放入 streams[0]。
   • 注意：在 streams[0] 内部，这四个任务必须按顺序执行。
2. 循环 i = 1 (处理第二个流)：
   • 将 kernel_1 放入 streams[1]。
   • 将 kernel_2 放入 streams[1]。
   • … 依此类推。

---

设备端（GPU）：硬件调度细节

这是最容易产生理解偏差的地方。虽然你使用了多个流（意图是并行），但由于提交顺序是“深度优先”，GPU 的执行表现往往如下：

A. 内部顺序性

在同一个流（比如 streams[0]）里，GPU 严格遵守 FIFO（先入先出）。也就是说，kernel_2 绝对不会在 kernel_1 完成之前开始。

B. 虚假依赖 (False Dependency) 风险

在早期的 CUDA 架构或未开启 Hyper-Q 的情况下，GPU 只有一个硬件工作队列。当你按深度优先提交时：

• 硬件队列的前端被 streams[0] 的四个任务占满。
• 即使 streams[1] 的 kernel_1 已经发送到了 GPU，但它在队列里排在 streams[0] 的 kernel_4 后面。
• 结果：原本可以在不同流之间并行的任务，被迫变成了串行执行。

C. Hyper-Q 的介入

由于代码开头设置了 CUDA_DEVICE_MAX_CONNECTIONS = 32，这会启用 Hyper-Q。

• 流程变化： Hyper-Q 允许 GPU 拥有多个硬件管理队列。
• 理想效果： 当 streams[0] 的 kernel_1 占用了部分计算资源（SM）后，如果还有空闲的计算资源，Hyper-Q 能够跳过 streams[0] 中排队的 kernel_2，直接从另一个硬件队列抓取 streams[1] 的 kernel_1 出来运行。

我们编译并利用nsys分析simpleHyperqDepth.cu

nvcc -arch=sm_120 simpleHyperqDepth.cu -o simpleHyperqDepth
nsys profile --stats=true ./simpleHyperqDepth

6.2.2 使用OpenMP的调度操作

OpenMP是CPU的并行编程模型，它使用编译器指令来识别并行区域。支持OpenMP指令的编译器可以将它们用作如何并行化应用程序的提示。用很少的代码，在主机上就可以实现多核并行。

核心代码如下

omp_set_num_threads(n_streams);
#pragma omp parallel
{
    int i = omp_get_thread_num();
    kernel_1<<<grid, block, 0, streams[i]>>>();
    kernel_2<<<grid, block, 0, streams[i]>>>();
    kernel_3<<<grid, block, 0, streams[i]>>>();
    kernel_4<<<grid, block, 0, streams[i]>>>();
}

通过以下命令编译simpleHyperqOpenmp.cu

nvcc -O3 -Xcompiler -fopenmp simpleHyperqOpenmp.cu -o simpleHyperqOpenmp -lgomp
./simpleHyperqOpenmp

得到如下结果

CUDA_DEVICE_MAX_CONNECTIONS = 32
> Using Device 0: NVIDIA GeForce RTX 5060 Ti with num_streams=4
> Compute Capability 12.0 hardware with 36 multi-processors
> grid 1 block 1
Measured time for parallel execution = 0.215s

6.2.3 创建流间依赖关系

在理想情况下，流之间不应该有非计划之内的依赖关系（即虚假的依赖关系）。然而，在复杂的应用程序中，引入流间依赖关系是很有用的，它可以在一个流中阻塞操作直到另一个流中的操作完成。事件可以用来添加流间依赖关系。

假如我们想让一个流中的工作在其他所有流中的工作都完成后才开始执行，那么就可以使用事件来创建流之间的依赖关系。首先，将标志设置为cudaEventDisableTiming，创建同步事件

cudaEvent_t *kernelEvent;
kernelEvent = (cudaEvent_t *) malloc(n_streams * sizeof(cudaEvent_t));

for (int i = 0; i < n_streams; i++)
{
    CHECK(cudaEventCreateWithFlags(&(kernelEvent[i]),
                cudaEventDisableTiming));
}

使用cudaEventRecord函数，在每个流完成时记录不同的事件。然后，使用cudaStreamWaitEvent使最后一个流（即streams[n_streams-1]）等待其他所有流：

for (int i = 0; i < n_streams; i++)
{
    kernel_1<<<grid, block, 0, streams[i]>>>();
    kernel_2<<<grid, block, 0, streams[i]>>>();
    kernel_3<<<grid, block, 0, streams[i]>>>();
    kernel_4<<<grid, block, 0, streams[i]>>>();

    CHECK(cudaEventRecord(kernelEvent[i], streams[i]));
    CHECK(cudaStreamWaitEvent(streams[n_streams - 1], kernelEvent[i], 0));
}

6.3 重叠内核执行和数据传输

重叠内核执行与数据传输 是提升 GPU 性能的核心技术。它的本质是利用 GPU 内部的不同硬件引擎（计算引擎和复制引擎），让“算”和“搬”同时进行。

现代 GPU 拥有独立的硬件单元来处理不同的任务：

• 计算引擎 (Execution Engine)： 负责运行核函数（Kernel）。
• 复制引擎 (Copy Engine)： 负责主机（CPU）与设备（GPU）之间的数据传输（cudaMemcpy）。

由于这些硬件是独立的，GPU 可以在从内存读取下一组数据的同时，处理当前已经在显存里的数据。

6.3.1 使用深度优先调度重叠

实现向量加法的CUDA程序，其基本结构包含3个主要步骤：

·将两个输入向量从主机复制到设备中

·执行向量加法运算

·将单一的输出向量从设备返回主机中

为了在向量加法中实现重叠，需要将输入和输出数据集划分成子集，并将来自一个子集的通信与来自于其他子集的计算进行重叠。具体对向量加法来说，需要将两个长度为N的向量加法问题划分为长度为N/M的向量相加的M个子问题。因为这里的每个子问题都是独立的，所以每一个都可以被安排在不同的CUDA流中，这样它们的计算和通信就可以重叠了。

1）数据传输是通过同步复制函数实现的。要重叠数据传输和内核执行，必须使用异步复制函数。因为异步复制函数需要固定的主机内存，所以首先需要使用cudaHostAlloc函数，在固定主机内存中修改主机数组的分配：

float *h_A, *h_B, *hostRef, *gpuRef;
CHECK(cudaHostAlloc((void**)&h_A, nBytes, cudaHostAllocDefault));
CHECK(cudaHostAlloc((void**)&h_B, nBytes, cudaHostAllocDefault));
CHECK(cudaHostAlloc((void**)&gpuRef, nBytes, cudaHostAllocDefault));
CHECK(cudaHostAlloc((void**)&hostRef, nBytes, cudaHostAllocDefault));

2）在NSTREAM流中平均分配该问题的任务。每一个流要处理的元素数量使用以下代码定义：

int iElem = nElem / NSTREAM;
size_t iBytes = iElem * sizeof(float);

3）使用一个循环来为几个流同时调度iElem个元素的通信和计算

for (int i = 0; i < NSTREAM; ++i)
{
    int ioffset = i * iElem;
    CHECK(cudaMemcpyAsync(&d_A[ioffset], &h_A[ioffset], iBytes,
                          cudaMemcpyHostToDevice, stream[i]));
    CHECK(cudaMemcpyAsync(&d_B[ioffset], &h_B[ioffset], iBytes,
                          cudaMemcpyHostToDevice, stream[i]));
    sumArrays<<<grid, block, 0, stream[i]>>>(&d_A[ioffset], &d_B[ioffset],
            &d_C[ioffset], iElem);
    CHECK(cudaMemcpyAsync(&gpuRef[ioffset], &d_C[ioffset], iBytes,
                          cudaMemcpyDeviceToHost, stream[i]));
}

运行得到结果如下

> ./simpleMultiAddDepth Starting...
> Using Device 0: NVIDIA GeForce RTX 5060 Ti
> Compute Capability 12.0 hardware with 36 multi-processors
> CUDA_DEVICE_MAX_CONNECTIONS = 1
> with streams = 4
> vector size = 262144
> grid (2048, 1) block (128, 1)

Measured timings (throughput):
 Memcpy host to device  : 1.086976 ms (0.964673 GB/s)
 Memcpy device to host  : 0.142944 ms (7.335572 GB/s)
 Kernel                 : 272.716797 ms (0.007690 GB/s)
 Total                  : 273.946716 ms (0.007655 GB/s)

Actual results from overlapped data transfers:
 overlap with 4 streams : 74.629761 ms (0.028101 GB/s)
 speedup                : 72.757561 
Arrays match.

代码将原本庞大的向量加法任务（262,144 个元素）平均拆分成了 4 份。

在传统的“顺序执行”模式下，GPU 的硬件引擎是闲置的：搬运数据时计算核心在等，计算时搬运引擎在等。而在“流并行”模式下，它利用了 GPU 拥有独立计算引擎和复制引擎的特性。

• 数据分区：每一块数据的大小为 iBytes。
• 异步操作：使用 cudaMemcpyAsync 和带流参数的内核启动。
• 重叠逻辑：
  • 当 Stream 0 正在执行计算（Kernel）时。
  • Stream 1 可以同时利用复制引擎将下一块数据从 CPU 搬运到 GPU。
  • 这就实现了“计算”与“传输”在时间轴上的重叠。

6.3.2 使用广度优先调度重叠

编译并执行

nvcc -arch=sm_120 simpleMultiAddBreadth.cu -o simpleMultiAddBreadth
./simpleMultiAddBreadth

得到以下结果

> ./simpleMultiAddBreadth Starting...
> Using Device 0: NVIDIA GeForce RTX 5060 Ti
> Compute Capability 12.0 hardware with 36 multi-processors
> CUDA_DEVICE_MAX_CONNECTIONS = 1
> with streams = 4
> vector size = 262144
> grid (2048, 1) block (128, 1)

Measured timings (throughput):
 Memcpy host to device  : 0.264352 ms (3.966590 GB/s)
 Memcpy device to host  : 0.157408 ms (6.661517 GB/s)
 Kernel                 : 283.888214 ms (0.007387 GB/s)
 Total                  : 284.309967 ms (0.007376 GB/s)

Actual results from overlapped data transfers:
 overlap with 4 streams : 71.616707 ms (0.029283 GB/s)
 speedup                : 74.810341 
Arrays match.

广度优先的核心在于**“分层提交”**。代码不再在一个循环里完成某个流的所有操作，而是分成了三个独立的循环：

1. 第一层循环（发射所有 H2D 拷贝）：连续向所有流派发“主机到设备”的传输任务。
2. 第二层循环（发射所有 Kernel）：连续向所有流派发计算任务。
3. 第三层循环（发射所有 D2H 拷贝）：连续向所有流派发“设备到主机”的回传任务。

在整体性能方面会发现，使用广度优先方法不如使用深度优先方法。由主机到设备复制队列上的争用导致的虚假依赖关系，在主机到设备间的传输完成前，将阻止所有的内核启动。

6.4 重叠GPU和CPU执行

相对而言，实现GPU和CPU执行重叠是比较简单的，因为所有的内核启动在默认情况下都是异步的。因此，只需简单地启动内核，并且立即在主机线程上实现有效操作，就能自动重叠GPU和CPU执行。

以下使用了3个CUDA操作（两个复制和一个内核启动）。记录一个停止事件，以标记所有CUDA操作的完成。

// asynchronously issue work to the GPU (all to stream 0)
CHECK(cudaMemcpyAsync(d_a, h_a, nbytes, cudaMemcpyHostToDevice));
kernel<<<grid, block>>>(d_a, value);
CHECK(cudaMemcpyAsync(h_a, d_a, nbytes, cudaMemcpyDeviceToHost));
CHECK(cudaEventRecord(stop));

特性	阻塞复制 (cudaMemcpy)	异步复制 (cudaMemcpyAsync)
CPU 行为	等待。CPU 会停在这一行，直到数据完全搬运完毕。	立即返回。CPU 只是把任务“交待”给 GPU，然后立刻执行下一行代码。
同步性	隐式同步。确保执行到下一行时，数据已经在目的地了。	非同步。执行下一行时，传输可能还没开始。
流（Stream）	通常在默认流中执行，会阻塞其他所有操作。	需要指定流（代码中默认是 stream 0），允许与其他流的任务并发。

虽然这些函数本身是异步的（CPU 不会等），但在 同一个流（Stream 0） 中，CUDA 保证任务是保序执行的：

1. kernel 必须等 H2D 拷贝完成才会开始。
2. D2H 拷贝必须等 kernel 运行完才会开始。

6.5 流回调

回调函数被放入 CUDA 流中，其执行顺序遵循流的 FIFO（先入先出） 原则。

• 加入流中：使用 cudaStreamAddCallback 将一个函数指针放入流中。
• 等待执行：该回调函数不会立即运行。它必须等待流中位于它之前的任务（如核函数或内存拷贝）全部完成。
• 阻塞后续任务：在回调函数执行完毕之前，该流中位于回调之后的任务都不会开始执行。

回调函数必须符合特定的签名。它是在 CPU 线程中执行的，通常由 CUDA 驱动程序内部的一个专用线程调用。

cudaError_t cudaStreamAddCallback(cudaStream_t stream, cudaStreamCallback_t callback, void *userData, unsigned int flags);

由于回调函数是在 CUDA 内部线程中运行的，使用时需要遵守以下禁令，否则会导致 死锁（Deadlock）：

• 禁止调用 CUDA API：回调函数内部绝对不能调用任何 CUDA API 函数（尤其是带有同步性质的），否则会导致系统死锁。
• 禁止进行同步操作：回调内部不能尝试获取可能会被其他正在等待 CUDA 任务的线程占用的互斥锁。

6.6 总结

一、 CUDA 流与异步架构

CUDA 流是实现网格级并发的基石，它是一系列按序在设备上执行的异步操作队列。

• 流的分类：分为显式声明的“非空流”和隐式默认的“空流”。若要实现不同操作的重叠执行，必须使用非空流。
• 异步本质：异步 API 调用（如内核启动和 cudaMemcpyAsync）会立即归还控制权给 CPU，从而允许 CPU 处理后续逻辑，而 GPU 在后台并行工作。
• 内存要求：执行异步数据传输必须使用由 cudaMallocHost 分配的“页锁定内存（Pinned Memory）”，以确保 DMA 传输的物理地址稳定性。

二、 流调度与硬件演进

GPU 如何处理多个流取决于其微架构，随着技术更迭，流间的干扰（虚假依赖）逐渐被消除。

• 架构对比：早期 Fermi 架构受限于单一工作队列，易产生串行化；Kepler 引入 Hyper-Q 提供了 32 个硬件连接；最新的 Blackwell 架构则通过“分布式任务分布引擎”实现了原生的乱序发射与 SM 资源的动态切分。
• 流优先级：开发者可以通过 API 设置流的优先级，高优先级流会在硬件调度中获得更高的发射权重，在 Blackwell 上甚至能实现微秒级的硬件抢占。

三、 事件机制与精细同步

CUDA 事件是轻量级的同步和计时对象，是流控管理的“信号灯”。

• 三大功能：精细的时间测量（优于 CPU 计时）、流间同步（跨流信号传递）以及主机端的进度跟踪。
• 同步策略：包括阻塞 CPU 的显式同步（如 cudaDeviceSynchronize）和由硬件处理、不阻塞 CPU 的流间同步（如 cudaStreamWaitEvent）。
• 隐式同步风险：某些副作用（如默认流的操作、内存分配或配置修改）会触发隐式同步，强制使 GPU 和 CPU 停下来“握手”，从而破坏并行性。

四、 并发执行的最佳实践

通过实验对比，文中揭示了任务提交顺序对性能的显著影响。

• 深度优先 (Depth-First)：按流提交任务（传-算-传），在计算密集型场景表现良好。
• 广度优先 (Breadth-First)：按任务类型分层提交（全传-全算-全回）。实验数据表明，广度优先能进一步压榨 RTX 5060 Ti 的性能，将加速比提升至约 74.81%。
• CPU 并行加速：利用 OpenMP 等模型可以在 CPU 端开启多线程，从而更快速地向 GPU 派发多个流的任务。

五、 计算与通信的重叠（Overlap）

这是全文的性能核心，旨在掩盖数据传输延迟。

• 流水线模式：通过将数据拆分为多个子集并分配至独立流，使 GPU 在计算当前块时，能利用空闲的复制引擎搬运下一块数据。
• 流回调 (Callbacks)：允许在 GPU 流完成特定任务后，在 CPU 端自动触发 C/C++ 函数，实现更高级的异步协作。但需严格禁止在回调中调用同步 CUDA API 以免死锁。