CUDA Graphs 101

Sally Stevenson, Senior System SW Engineer | GTC 2023
March 22, 2023

目录 (Table of Contents)

课程计划 (Lesson Plan)

本演讲的课程计划包括：
* 什么是 CUDA Graphs？
* 编程模型概述
* 性能提示与技巧
* CUDA Graphs 的新特性
* CUDA 设备图启动

1. 什么是 CUDA Graphs？

1.1 加速工作启动与执行 (Speeding Up Work Launch And Execution)

CUDA Graphs 旨在通过将 CUDA 流中的依赖关系转换为图结构，从而加速工作流的启动和执行。

  • CUDA Streams 中的工作 (CUDA Work in Streams)：操作按顺序调度并立即执行，需要通过显式等待（Wait）来管理跨流依赖。
  • 依赖图 (Graph of Dependencies)：任何 CUDA 流都可以映射到一个图。一旦工作流完全定义，图中的操作将被统一调度。这种方式允许在已知工作流的情况下进行更优化的执行。
CUDA Work in Streams vs Graph of Dependencies
CUDA Work in Streams vs Graph of Dependencies

1.2 在预先知晓工作流时进行执行优化 (Execution Optimization When Workflow is Known Up-Front)

CUDA Graphs 适用于那些工作流在执行前已知且重复性高的场景，从而实现执行优化。典型的应用包括：

  • 循环与函数卸载 (Loop & Function offload)：将循环或函数操作卸载到 GPU。
  • 深度学习推理 (DL Inference)：优化神经网络的推理过程。
  • 线性代数 (Linear Algebra)：加速复杂的线性代数运算图。
  • 深度神经网络训练 (Deep Neural Network Training)：优化反向传播等训练过程。
  • 高性能计算模拟 (HPC Simulation)：例如分子动力学模拟等科学计算。
CUDA Graphs应用场景
CUDA Graphs应用场景

2. 性能影响

2.1 深度学习 (Deep Learning): PyTorch - 机器学习框架 (PyTorch - Machine Learning Framework)

在 PyTorch 基准测试中，CUDA Graphs 展现了显著的性能提升，尤其是在某些模型上。测试环境为 CUDA 12.0, DGX-H100, Ubuntu 20.04。

  • Retinanet：多节点图实现了约 2.5x 的加速，单节点图实现了约 1.7x 的加速。
  • Mask R-CNN：多节点图实现了约 2.1x 的加速，单节点图实现了约 1.0x 的加速。
  • BERT：多节点图实现了约 1.2x 的加速，单节点图实现了约 1.0x 的加速。
  • RNN-T：未显示出显著加速。
PyTorch Benchmark Performance
PyTorch Benchmark Performance

2.2 高性能计算 (HPC): GROMACS - 分子动力学模拟 (GROMACS - Molecular Dynamics Simulation)

在 GROMACS 分子动力学模拟中，CUDA Graphs 也带来了性能提升。测试环境为 CUDA 11.8, DGX-A100, Ubuntu 20.04。

单 GPU (Single-GPU)

在单 GPU 配置下，对于 GROMACS Water Box 模拟，CUDA Graphs 在不同的系统规模（以千原子数计）下，相比 Streams 均实现了略微的加速，通常在 1.01x 到 1.05x 之间。

多 GPU (Multi-GPU)

在多 GPU 配置下，GROMACS Water Box 模拟显示：
* Graphs Today：当前的 CUDA Graphs 实现也提供了类似的轻微加速。
* Upcoming Release (Projected)：预计在未来的版本中，CUDA Graphs 将在更大的系统规模下提供更显著的加速，例如在 768k 原子数时可达约 1.15x 的加速。

GROMACS Performance
GROMACS Performance

2.3 通信 (Comms): Aerial - 网络与通信框架 (Aerial - Networking & Communications Framework)

在 Aerial 网络与通信框架中，CUDA Graphs 有助于降低延迟和减少抖动。测试环境为 CUDA 11.7, DGX-A100, 1xA100, Ubuntu 20.04。

平均延迟加速 (Average Latency Speedup)

  • PDSCH：实现了约 1.7x 的加速。
  • PRACH：实现了约 1.3x 的加速。
  • SSB：实现了约 1.2x 的加速。
  • CSI-RS：实现了约 1.1x 的加速。
  • PUSCH：实现了约 1.1x 的加速。
  • PDCCH 和 PUCCH 的加速接近 1x。

平均抖动降低 (Average Jitter Reduction)

  • PRACH：实现了约 2.2x 的显著抖动降低。
  • PDSCH：实现了约 1.8x 的抖动降低。
  • CSI-RS：实现了约 1.3x 的抖动降低。
  • PUSCH：实现了约 1.4x 的抖动降低。
Aerial Performance
Aerial Performance

3. 性能来源与重复使用

3.1 减少短运行内核的系统开销 (Reducing System Overheads Around Short-Running Kernels)

性能提升主要来源于减少围绕短运行内核的系统开销。

执行期间的时间花费分解 (Breakdown of time spent during execution)

初始情况下，每次内核启动都包含 Grid Initialization 和内核执行（例如 2µs Kernel）。在这种模式下，Grid Initialization 占用了大量时间，导致总开销高达 64%。

Execution Time Breakdown - Initial
Execution Time Breakdown - Initial

CPU 侧启动开销减少 (CPU-side launch overhead reduction)

通过优化 CPU 侧的启动过程，减少了每次内核启动所需的 CPU 开销。这使得总开销从 64% 降低到 45%。

Execution Time Breakdown - CPU-side reduction
Execution Time Breakdown - CPU-side reduction

设备侧执行开销减少 (Device-side execution overhead reduction)

进一步地，通过优化设备侧的执行开销，例如将 Grid Initialization 流程与内核执行更紧密地集成或预处理，使得总开销进一步降低到 33%。
对于三个 2µs 的内核，总时间缩短了 29%。

Execution Time Breakdown - Device-side reduction
Execution Time Breakdown - Device-side reduction

3.2 两个独立的优化：分离测量启动成本和网格初始化 (Measuring Launch Cost Separately From Grid Initialization)

CUDA Graphs 实现了两个独立的优化，分别针对 Host 和 Device 上的启动成本。测试环境为 CUDA 12.1, RTX A5000, Intel i7-7800X, Ubuntu 18.04。

每个节点的 Host 启动时间 (Host Launch Time Per Node)

  • Stream Launch (流启动)：对于 100 节点直线图和 2 分叉 100 节点直线图，启动时间约为 1.8-1.9 微秒。
  • First Graph Launch (首次图启动)：首次启动图的成本显著降低，约为 0.4-0.5 微秒。
  • Graph Relaunch (图重新启动)：图的重新启动成本最低，约为 0.3-0.4 微秒，显示了重复使用的优势。

每个节点的 Device 启动时间 (Device Launch Time Per Node)

  • Stream Launch (流启动)：对于两种图类型，设备启动时间约为 1.4 微秒。
  • First Graph Launch (首次图启动)：首次设备图启动时间约为 0.8 微秒。
  • Graph Relaunch (图重新启动)：设备图重新启动时间最低，约为 0.7 微秒。

这些数据显示了 CUDA Graphs 通过减少启动和重新启动的开销，尤其是在重复执行时，能够带来显著的性能提升。

Launch Time Per Node
Launch Time Per Node

3.3 Graphs 受益于重复使用：CUDA Graphs 应用于可重复的工作负载 (CUDA Graphs Should Be Used For Repeatable Workloads)

CUDA Graphs 的主要优势在于对可重复工作负载的优化。通过多次重新启动已实例化的图，可以摊销首次创建和实例化的开销。测试环境为 Straight-Line Graph, CUDA 12.1, RTX A5000, Intel i7-7800X, Ubuntu 18.04。

  • Graphs vs Streams, 最佳情况盈亏平衡点 (Best-Case Break Even Point)：
    • 对于图长度为 1 或 2 个节点，需要至少 9 次图启动才能比流更快。
    • 对于图长度为 4 个节点，需要至少 6 次图启动才能实现盈亏平衡。
    • 对于图长度为 8 或 16 个节点，需要至少 4 次图启动。
    • 对于图长度为 32、64、128 或 256 个节点，需要至少 3 次图启动。

结论：大多数图需要至少 3-4 次启动才能比流更快，这强调了 CUDA Graphs 在重复性工作负载中的价值。

Graphs vs Streams Break Even Point
Graphs vs Streams Break Even Point

4. 编程模型概述

4.1 三阶段执行模型 (THREE-STAGE EXECUTION MODEL Minimizes Execution Overheads - Pre-Initialize As Much As Possible)

CUDA Graphs 采用三阶段执行模型来最小化执行开销，尽可能地进行预初始化。

  1. 定义 (Define)：

    • 创建一个“单一图模板”。
    • 可以在宿主代码中创建，或从库中构建。此阶段定义了图的结构和操作。

  2. 实例化 (Instantiate)：

    • 从模板创建“多个可执行图”。
    • 这是模板的快照。此阶段设置并初始化 GPU 执行结构（创建一次，运行多次），为实际执行做好准备。

  3. 执行 (Execute)：

    • 可执行图在 CUDA 流中运行。
    • 图中的并发性不受单个 CUDA 流的限制，可以跨多个流或在单个流中实现更细粒度的并发。
Three-Stage Execution Model
Three-Stage Execution Model

4.2 图节点能进行哪些操作？ (WHAT OPERATIONS CAN A GRAPH NODE DO? Everything You Would Expect)

一个 CUDA Graph 是由节点组成的，这些节点之间通过依赖关系连接，代表一系列操作。

  • 图中的节点可以跨多个设备执行。
  • 图节点的操作类型包括：
    • Kernel Launch (内核启动)：在 GPU 上运行 CUDA 内核。
    • CPU Function Call (CPU 函数调用)：作为回调函数在 CPU 上执行。
    • Memcopy/Memset (内存拷贝/设置)：GPU 数据管理操作。
    • Memory Alloc/Free (内存分配/释放)：内存管理操作。
    • External Dependency (外部依赖)：通过外部信号量/事件管理依赖关系。
    • Child Graph (子图)：图可以是分层的，一个节点可以是一个独立的子图。
Graph Node Operations
Graph Node Operations

4.3 直接创建图 (CREATE GRAPHS DIRECTLY Map Graph-Based Workflows Directly Into CUDA)

CUDA 提供了直接的 API 来创建和管理图，允许开发者将基于图的工作流直接映射到 CUDA。

以下代码片段展示了直接创建图的典型流程：

// 定义工作和依赖关系的图
cudaGraphCreate(&graph);

cudaGraphAddNode(graph, kernel_a, {...}); // 添加内核a节点
cudaGraphAddNode(graph, kernel_b, {kernel_a}, ...); // 添加内核b节点，依赖于a
cudaGraphAddNode(graph, kernel_c, {kernel_a}, ...); // 添加内核c节点，依赖于a
cudaGraphAddNode(graph, kernel_d, {kernel_b, kernel_c}, ...); // 添加内核d节点，依赖于b和c

// 实例化图并应用优化
cudaGraphInstantiate(&graphExec, graph);

// 启动可执行图 100 次
for (int i=0; i<100; i++)
    cudaGraphLaunch(graphExec, stream);

这段代码首先通过 cudaGraphCreate 创建一个空的图，然后使用 cudaGraphAddNode 添加一系列内核节点及其依赖关系。接着，cudaGraphInstantiate 将定义的图编译并实例化为可执行图。最后，通过循环调用 cudaGraphLaunch 多次启动同一个可执行图，从而利用了图的重复使用优势。

Direct Graph Creation Code Example
Direct Graph Creation Code Example

4.4 流捕获 (Stream Capture)

流捕获概述

流捕获允许在不重写现有 CUDA 代码的情况下，将一系列流操作转换为一个 CUDA 图，从而利用图的性能优势。

流捕获示例及图转换 (Page 16)
流捕获示例及图转换 (Page 16)

核心步骤：

  1. 开始捕获： 使用 cudaStreamBeginCapture(&stream1); 启动对指定流的捕获。
  2. 构建流工作： 正常编写 CUDA 内核启动、事件记录和等待等操作。
  3. 结束捕获： 使用 cudaStreamEndCapture(stream1, &graph); 将捕获到的操作转换为一个 CUDA 图。
流捕获原理 - 遵循流间依赖关系 (Page 17)
流捕获原理 - 遵循流间依赖关系 (Page 17)

捕获过程会遵循流间的依赖关系（例如通过 cudaStreamWaitEvent 建立的依赖），以创建图中的分支和连接，从而形成一个有向无环图（DAG）。

结合图与流工作

可以将库调用捕获为子图，并将其添加到现有图中。

结合图与流工作，捕获库调用作为子图 (Page 18)
结合图与流工作，捕获库调用作为子图 (Page 18)

示例步骤：
1. 创建根节点： 通过显式 API (如 cudaGraphAddNode) 为主图创建根节点。
2. 捕获库调用： 使用 cudaStreamBeginCapture 和 cudaStreamEndCapture 将一个 libraryCall 捕获成一个独立的子图 library_graph。
3. 插入子图： 使用 cudaGraphAddSubGraphNode 将捕获到的子图作为节点插入到主图中。
4. 继续构建主图： 继续通过显式 API 构建主图的其余部分。

实践中的流捕获

仅可捕获完全异步的序列。

实践中的流捕获 - 原始代码与捕获代码对比 (Page 19)
实践中的流捕获 - 原始代码与捕获代码对比 (Page 19)

对于许多应用程序，流捕获需要进行一些调整，因为并非所有代码都能“直接工作”于捕获模式。例如，像 cudaDeviceSynchronize() 这样的同步调用，在捕获中会带来复杂性。

流捕获的局限性

某些操作无法被捕获。

  1. 默认（“空”）流无法被捕获。
    流捕获局限性 - 默认流无法捕获 (Page 21)
    libraryCall(cudaStreamDefault) 必须替换为指定流的 libraryCall(stream)。

  2. 同步调用无法被捕获。
    流捕获局限性 - 同步调用无法捕获 (Page 22)
    例如，cudaMalloc(...)、cudaMemcpy(...) 和 cudaFree(...) 必须替换为它们的异步版本，如 cudaMallocAsync(..., stream)、cudaMemcpyAsync(..., stream) 和 cudaFreeAsync(..., stream)。

  3. 没有异步对应物的调用必须发生在捕获之外。
    流捕获局限性 - 无异步对应物调用 (Page 23)
    例如，cudaMallocHost(...) 和 cudaFreeHost(...) 等主机内存操作，由于没有异步版本，需要在 cudaStreamBeginCapture 之前或 cudaStreamEndCapture 之后执行。

  4. 流捕获不能同步。
    流捕获局限性 - 不能同步 (Page 24)
    cudaDeviceSynchronize() 虽然可以存在于代码中，但其行为在捕获的图上下文中会有所不同，捕获本身不提供同步机制。

图的实践

仅仅使代码可捕获就完成了吗？
一段可捕获的代码可能包含如下结构：

cudaMallocHost(...);
cudaStreamBeginCapture(stream, ...);
cudaMallocAsync(..., stream);
cudaMemcpyAsync(..., stream);
hostLogic(...);
libraryCall(stream);
cudaFreeAsync(..., stream);
cudaStreamEndCapture(stream, &graph);
cudaFreeHost(...);
// Instantiate & launch the graph

那么，这段代码现在就图就绪了吗？不一定，还需要进一步优化。

4.5 图采用技巧 - 避免常见陷阱

图采用技巧 - 内存管理 (Page 26)
图采用技巧 - 内存管理 (Page 26)

技巧 #1： 通过 cudaMalloc/FreeAsync 将内存管理放入捕获中。
将 cudaMallocAsync 和 cudaFreeAsync 调用放置在 cudaStreamBeginCapture 和 cudaStreamEndCapture 之间，使其成为图的一部分。

图采用技巧 - 其他分配 (Page 27)
图采用技巧 - 其他分配 (Page 27)

技巧 #1.5： 对于其他分配（如 cudaFreeHost(...)）—— 保持内存存在，否则在启动时无法访问。这类操作通常应发生在图实例化和启动之外。

图采用技巧 - CPU 回调 (Page 28)
图采用技巧 - CPU 回调 (Page 28)

技巧 #2： 将任何重要逻辑放入 CPU 回调。
hostLogic(...) 等 CPU 端逻辑，应通过 cudaLaunchHostFunc(stream, hostLogic, ...); 封装，使其能在 GPU 图中异步执行。

图采用技巧 - 线程局部捕获模式 (Page 29)
图采用技巧 - 线程局部捕获模式 (Page 29)

技巧 #3： 如果您的应用程序是多线程的且线程独立运行，请考虑使用线程局部捕获模式。
使用 cudaStreamBeginCapture(stream, threadLocal); 来启动线程局部的捕获。

5. 图实例化与工作启动方式

5.1 图实例化 (Page 31)

图实例化等同于图的编译步骤：
* 准备并优化图以供执行。
* 实例化后，可执行图结构被锁定。
* 结构性更改需要重新实例化。
与代码编译一样，实例化不是一个简单的步骤，需要额外的时间。而且，像任何编译步骤一样，实例化不会为您处理所有事情。

图实例化
图实例化

5.2 无自动放置与结构更改限制

无自动放置 (Page 32)

用户必须为每个节点定义执行位置。
如果图中的一个分支可以在2个GPU上运行，我们如何选择在何处运行什么？
最佳选择可能取决于数据局部性——在执行层是未知的。

无自动放置
无自动放置

无结构更改 (Page 33)

执行层不具备执行此操作所需的信息。
* 不能拆分图节点。
* 不能合并图节点。
* 不能重新分配节点的执行位置。
元素操作可以很容易地融合，但仅当操作语义已知时。执行层只看到二进制代码，因此无法执行此合并。

无结构更改
无结构更改

5.3 图执行语义 (Page 35)

将图工作与其他非图CUDA工作进行排序。
如果可以在CUDA流中放置，就可以与图一起运行。

launchWork(cudaGraphExec_t i1, cudaGraphExec_t i2, CPU_Func cpu, cudaStream_t stream) {
    A <<< 256, 256, 0, stream >>>();   // Kernel launch
    cudaGraphLaunch(i1, stream);       // Graph launch
    cudaStreamAddCallback(stream, cpu); // CPU callback
    cudaGraphLaunch(i2, stream);       // Graph launch

    cudaStreamSynchronize(stream);
}
图执行语义
图执行语义

5.4 图忽略流序列化规则 (Page 36)

启动流仅用于与其他工作排序。
即使图被启动到一个流中，图中的分支仍然会并发执行。

图忽略流序列化规则
图忽略流序列化规则

5.5 三种工作启动方式

性能提升伴随着限制的增加。

图重新实例化 (Graph Re-instantiation)

  • 每次迭代重建工作：构建图 -> 实例化图 -> 启动图 -> 销毁图。
  • 参数：可能更改。
  • 拓扑：可能更改。
  • 迭代 10 次。

图更新 (Graph Update)

  • 每次迭代更新图：定义图 -> 更新图 -> 启动图。
  • 参数：可能更改。
  • 拓扑：可能不更改。
  • 迭代 10 次。
    如果只有参数发生了变化怎么办？重新实例化是唯一的选择吗？图更新是解决此问题的方法。

图重新启动 (Graph Re-Launch)

  • 每次启动相同的图：定义图 -> 启动图。
  • 参数：可能不更改。
  • 拓扑：可能不更改。
  • 迭代 10 次。
3 种工作启动方式 (Page 39)
3 种工作启动方式 (Page 39)

5.6 案例研究

自动驾驶车辆 (Page 40)

不更新的重新启动。
在实时系统中，图必须保持静态以实现可靠的启动时间。

案例研究：自动驾驶车辆
案例研究：自动驾驶车辆

金融应用 (Page 41)

带更新的重新启动。
图的孤立加速不是全貌。整个应用20-30%可图化，因此整体加速较小。

案例研究：金融应用
案例研究：金融应用

5.7 将现有代码适配到图 (Page 42)

根据程序结构选择方法。

方法 描述 速度提升
图重新实例化 每次迭代重建工作 不比流快
图更新 每次迭代更新图 比流快达 1.2x
图重新启动 每次启动相同的图 比流快达 2.5x 
// Graph Re-instantiation
for(i=0; i<N; i++) {
    cudaStreamBeginCapture(stream);
    A<<< ..., stream >>>(data);
    B<<< ..., stream >>>(data);
    ...
    Z<<< ..., stream >>>(data);
    cudaStreamEndCapture(stream, &g);
    cudaGraphInstantiate(&graphExec, g);
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);
}

// Graph Update
for(i=0; i<N; i++) {
    cudaStreamBeginCapture(stream);
    A<<< ..., stream >>>(data[i]);
    B<<< ..., stream >>>(data[i]);
    ...
    Z<<< ..., stream >>>(data);
    cudaStreamEndCapture(stream, &g);
    cudaGraphExecUpdate(graphExec, g);
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);
}

// Graph Re-Launch
cudaStreamBeginCapture(stream);
A<<< ..., stream >>>(data);
B<<< ..., stream >>>(data);
...
Z<<< ..., stream >>>(data);
cudaStreamEndCapture(stream, &g);
cudaGraphInstantiate(&graphExec, g);
for(i=0; i<N; i++) {
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);
}
将现有代码适配到图
将现有代码适配到图

5.8 细粒度更新：单节点更新与启用/禁用

单节点更新 (Page 43)

一种更细粒度的参数更新方法。
如果您了解您的工作流程，您可以单独更新节点。

// Define graph
cudaGraphCreate(&graph);
cudaGraphAddNode(graph, kernel_a, {...});
...

// Instantiate graph
cudaGraphInstantiate(&graphExec, graph);

// Iterate 100 times
for (int i=0; i<100; i++) {
    generateNewParams(&newParams);
    // Update the parameters for A between launches
    cudaGraphExecKernelNodeSetParams(graphExec, kernel_a, newParams);
    cudaGraphLaunch(graphExec, stream);
}
单节点更新
单节点更新

单节点启用/禁用 (Page 44)

避免在次要拓扑更改时重新实例化。
节点也可以完全启用/禁用。

// Define graph
cudaGraphCreate(&graph);
cudaGraphAddNode(graph, kernel_a, {...});
...

// Instantiate graph
cudaGraphInstantiate(&graphExec, graph);

// Iterate 100 times
for (int i=0; i<100; i++) {
    checkIfShouldEnable(&enableNode);
    // Toggle A on/off between launches
    cudaGraphNodeSetEnabled(graphExec, kernel_a, enableNode);
    cudaGraphLaunch(graphExec, stream);
}
单节点启用/禁用
单节点启用/禁用

6. 性能优化与新特性

6.1 四阶段执行模型

上传步骤可以从启动中分离，以实现更好的流水线操作。

四阶段执行模型
四阶段执行模型

  1. 定义 (Define):

    • 单个图“模板”。
    • 在主机代码中创建或从库中构建。

  2. 实例化 (Instantiate):

    • 多个“可执行图”。
    • 模板的快照，设置并初始化GPU执行结构（创建一次，运行多次）。

  3. 上传 (Upload) (3a):

    • 预上传启动资源。
    • 通常作为首次启动的一部分隐式执行，也可显式请求。

  4. 执行 (Execute) (3b):

    • 可执行图在CUDA流中运行。
    • 图中的并发性不受流限制。

6.2 图上传

旨在降低首次启动成本。

图上传以降低首次启动成本
图上传以降低首次启动成本

原始工作流程 (Original Workflow):

  • 构建图 -> 额外设置 -> 上传并启动图 -> 启动图 -> 同步。
  • 资源初始化在“上传并启动图”阶段进行。

带上传的工作流程 (With Upload):
* 构建图 -> 上传图 -> 额外设置 -> 启动图 -> 启动图 -> 同步。
* 资源初始化在单独的“上传图”阶段进行，从而可以减少首次启动的延迟。

6.3 性能来源回顾

通过减少短运行内核的系统开销来提升性能。

性能来源回顾
性能来源回顾
  • 启动 (Launch): 64% 的开销。
  • CPU端启动开销减少 (CPU-side launch overhead reduction): 45% 的开销。
  • 设备端执行开销减少 (Device-side execution overhead reduction): 33% 的开销。
  • 总结: 三个2µS内核的总时间缩短了29%。

6.4 性能技巧与窍门

优化GPU运行时

通过关注内核来优化GPU运行时。

通过关注内核优化GPU运行时
通过关注内核优化GPU运行时
  • 不推荐 (左图): 回调函数、内存设置 (Memset) 和内存拷贝 (Memcpy) 分散在内核执行之间。
  • 推荐 (右图): 尽可能将所有操作转换为内核，形成内核密集型工作流。
  • 核心信息: 内核密集型工作流将获得更多的GPU加速。

改进依赖解析

改进依赖解析
改进依赖解析
  • 依赖关系解析: 内核到内核的依赖解析比内核到其他操作（如Memset/Memcpy）更快。
  • 开销: 在内核之间直接连接的开销通常较小，而涉及Memset或Memcpy的连接开销较大。

Ampere带来新的硬件能力

Ampere架构为CUDA Graphs带来了新的硬件能力。

Ampere带来的新硬件能力
Ampere带来的新硬件能力

  • 主机启动延迟 (Host Launch Latencies) (长度 = 100, CUDA 12.1, Intel I7-7800X, Ubuntu 18.04):

    • RTX A5000（Ampere架构）在“Repeated Fork Join”和“Single-Entry Parallel Straight-Line”模式下，相对于TITAN V显示出显著更高的加速比。

  • 设备启动延迟 (Device Launch Latencies):

    • RTX A5000在设备启动延迟方面也表现出类似的优势。

  • 建议: 升级到Ampere或更高版本，以受益于新的硬件功能。

性能评估提示

  • 在实际运行之前，创建并实例化一个最大尺寸的图:

    • 这会预热驱动程序资源，使得后续图的实例化更快。

  • 将首次启动与第二次启动分开进行基准测试:

    • 首次启动包含一个在第二次启动中不存在的上传步骤。

  • 为工作选择正确的工具:

    • 性能分析器 (Profiler) 更适合分析单个节点。
    • CUDA事件更适合整个图的时间测量。

6.5 CUDA 11以来的新功能

CUDA 11以来的新功能概述
CUDA 11以来的新功能概述
  • CUDA 11.1: 事件节点 (Event Nodes), 外部依赖节点 (External Dependency Nodes), 空节点更新 (Empty Node Update), 子图更新 (Child Graph Update), 图上传 (Graph Upload)。
  • CUDA 11.2: 内核节点功能更新 (Kernel Node Function Update)。
  • CUDA 11.3: 图的用户对象 (User Objects For Graphs), 生成图的DOT文件用于调试 (Generation of a DOT File of a Graph For Debugging)。
  • CUDA 11.4: 内存分配与释放节点 (Memory Allocation & Free Nodes)。
  • CUDA 11.6: 单节点启用/禁用 (Single-Node Enable/Disable)。
  • CUDA 11.7: 每节点优先级 (Per-node Priorities)。
  • CUDA 12.0: 设备图启动 (Device Graph Launch)。

6.6 CUDA 11以来的性能改进

CUDA 11以来的性能改进
CUDA 11以来的性能改进
  • 直线图启动延迟 (Straight Line Graph Launch Latencies) (图长度 = 100, RTX A5000, Intel i7-7800X, Ubuntu 18.04):
    • 比较了CUDA 11.1和CUDA 12.0（主机延迟、设备延迟、总延迟）在首次启动和重复启动时的表现。
    • 结论: 自CUDA 11以来，性能也得到了提升！

7. CUDA 设备图启动

7.1 设备图启动概述

设备图启动允许图的动态控制流，将图的启动从主机卸载到设备。

设备图启动代码示例
设备图启动代码示例

  • CPU部分 (CPU portion):

    • 创建图G1和G2。
    • 实例化G1。
    • 实例化G2时指定DeviceLaunch。
    • 上传G2。
    • 启动G1。

  • GPU部分 (GPU portion):

    • 定义一个全局函数Y，在其内部启动图G2。

  • 说明: 实现了设备侧图启动，使得GPU能够自主地启动图，减少主机开销。

7.2 设备启动性能：与主机启动的比较

性能与节点数关系

设备启动性能与节点数关系
设备启动性能与节点数关系
  • 直线图启动延迟 (Launch Latencies for Straight-line Graphs) (CUDA 12.1, RTX A5000, Intel i7-7800X, Ubuntu 18.04)。
  • 结论:
    • 设备重新启动比主机启动更快，仅需11个节点。
    • “即发即忘 (fire-and-forget)”启动总是比主机启动至少快2倍！

性能与图结构影响

设备启动性能与图结构关系
设备启动性能与图结构关系
  • 主机和设备启动延迟 (Host & Device Launch Latencies) (图长度 = 100, CUDA 12.1, RTX A5000, Intel i7-7800X, Ubuntu 18.04)。
  • 比较了“直线图 (Straight line)”、“重复分支连接 (Repeated fork join)”和“单入口并行直线 (Single-entry parallel straight-line)”在主机启动、尾部启动和即发即忘启动下的延迟。
  • 结论: 设备启动受图结构的影响较小。

性能与图宽度影响

设备启动性能与图宽度关系
设备启动性能与图宽度关系
  • 并行直线图启动延迟 (Launch Latencies For Parallel Straight-line Graphs) (图长度 = 100, CUDA 12.1, RTX A5000, Intel i7-7800X, Ubuntu 18.04)。
  • 比较了主机启动、尾部启动和即发即忘启动随直线数量增加的延迟。
  • 结论: 设备启动在图宽度方面也具有更好的扩展性。

7.3 设备图的创建规则 (Rules For Device Graphs)

一个图只有在满足以下条件时才能从设备启动：

  1. 它只包含内核 (kernels)、内存拷贝 (memcopies) 和内存设置 (memsets)。
  2. 所有节点都驻留在单个设备上。
  3. 它被实例化用于设备启动 (例如通过 cudaGraphInstantiate(G2, DeviceLaunch))。
  4. 它已被明确上传到设备（如果不是从主机启动的话，例如通过 cudaGraphUpload(G2, ...)）。
  5. 它从另一个图启动（例如在 __global__ void Y(...) 中调用 cudaGraphLaunch(G2, ...)）。

代码示例:
device_launch.cu 文件中的代码展示了如何创建、实例化、上传和启动图：

// CPU端
void main() {
    cudaGraphCreate(&G1);
    // Build graph G1
    cudaGraphInstantiate(G1);

    cudaGraphCreate(&G2);
    // Build graph G2
    cudaGraphInstantiate(G2, DeviceLaunch); // 实例化用于设备启动
    cudaGraphUpload(G2, ...);               // 明确上传到设备

    cudaGraphLaunch(G1, ...);
}

// GPU端
__global__ void Y(cudaGraphExec_t G2) {
    cudaGraphLaunch(G2, ...); // 从另一个图启动
}

满足以上条件的图被称为设备图 (device graphs)。所有其他图都是主机图 (host graphs)。

7.4 设备图启动封装 (Device Graph Launch Encapsulation)

依赖解析发生在整个图的粒度上。

图的封装边界是整个启动图。这个边界被称为执行环境 (execution environment)。
图的启动不能在父图内创建新的依赖。图内没有 fork/join 并行。

例如，在 Graph G1 中，如果节点 Y 启动了 Graph G2，那么 G2 成为 Kernel2 的一个依赖，而不是节点 Z 的依赖。

设备图启动封装
设备图启动封装

7.5 设备图启动模式 (DEVICE GRAPH LAUNCH MODES)

命名流: 即发即弃 (Fire-and-Forget)

即发即弃模式会立即启动图。启动的图与父图并发运行。

后续工作会隐式地加入即发即弃启动。
即发即弃启动不能直接同步（例如，通过 cudaDeviceSynchronize()）。

那么，如何插入工作依赖呢？

即发即弃模式
即发即弃模式

命名流: 尾部启动 (Tail Launch)

尾部启动在调用图完成后按顺序启动。
它提供了一种插入工作依赖的方式。

尾部启动模式
尾部启动模式

7.6 尾部启动封装与排序

尾部启动封装 (TAIL LAUNCH ENCAPSULATION)

连接即发即弃启动。

即发即弃启动不能被直接同步（例如，通过 cudaDeviceSynchronize()）。
只能通过尾部启动来强制排序。尾部启动会加入即发即弃的工作。

尾部启动封装
尾部启动封装

尾部启动排序 (TAIL LAUNCH ORDERING)

尾部启动队列。

尾部启动按照它们被入队的顺序执行。

尾部启动排序
尾部启动排序

7.7 允许自重启 (Self Relaunch is Permitted)

在您的应用程序中实现主机无关的循环。

一个图可以重新启动自身，从而实现主机无关的循环。例如，一个 __global__ 函数可以在满足条件时，通过尾部启动重新启动当前图。

代码示例:

// CPU端
void main() {
    cudaGraphCreate(&G);
    // Build graph G
    cudaGraphInstantiate(G, DeviceLaunch);
    cudaGraphLaunch(G, ...);
}

// GPU端
__global__ void Y(void) {
    if (condition) {
        cudaGraphExec_t G = cudaGetCurrentGraphExec();
        cudaGraphLaunch(G, tailLaunch); // 自重启图 G
    }
}
自重启示例
自重启示例

7.8 即将推出：兄弟启动 (COMING SOON: SIBLING LAUNCH)

克服父图封装，将依赖关系转移到上层。

兄弟启动在先行工作和后续工作之间插入子工作。
子工作与父工作并发启动。
子工作成为父图的父图的依赖，但不会阻塞调度图的重新启动。
调度图本身作为尾部启动重新启动。

兄弟启动
兄弟启动

7.9 示例用法: 运行时动态工作调度 (SAMPLE USAGE: Run-Time Dynamic Work Scheduling)

程序初始化阶段:

在主机代码中，程序初始化期间创建多个图。

示例用法：初始化
示例用法：初始化

设备上的调度器内核:

设备上的调度器内核根据数据包类型选择并启动相应的图来处理传入的数据包。调度器图本身可以作为尾部启动重新启动，以实现持续处理。

代码示例:

__global__ void scheduler(...) {
    Packet data = receivePacket(...);

    switch(data.type) {
        case 1:
            cudaGraphLaunch(G1, FireAndForget);
            break;
        case 2:
            cudaGraphLaunch(G2, FireAndForget);
            break;
        case 3:
            cudaGraphLaunch(G3, FireAndForget);
            break;
        case 4:
            cudaGraphLaunch(G4, FireAndForget);
            break;
        case 5:
            cudaGraphLaunch(G5, FireAndForget);
            break;
    }

    // Re-launch the scheduler to run after processing
    cudaGraphExec_t currentGraphExec = cudaGetCurrentGraphExec();
    cudaGraphLaunch(currentGraphExec, tailLaunch);
}
示例用法：调度器
示例用法：调度器

8. 附加信息 (Additional Info)

开始使用图 (Get Started With Graphs)。

请阅读编程指南的 CUDA graphs 部分。

查阅 CUDA 示例：
* simpleCudaGraphs
* jacobiCudaGraphs
* graphMemoryNodes
* graphMemoryFootprint

开发者博客：
* Getting Started With CUDA Graphs
* Employing CUDA Graphs in a Dynamic Environment
* Enabling Dynamic Control Flow With Device Graph Launch

GTC 演讲：
* Effortless CUDA Graphs

NVIDIA Logo
NVIDIA Logo