Sally Stevenson, Senior System SW Engineer | GTC 2023
March 22, 2023
本演讲的课程计划包括:
* 什么是 CUDA Graphs?
* 编程模型概述
* 性能提示与技巧
* CUDA Graphs 的新特性
* CUDA 设备图启动
CUDA Graphs 旨在通过将 CUDA 流中的依赖关系转换为图结构,从而加速工作流的启动和执行。
CUDA Graphs 适用于那些工作流在执行前已知且重复性高的场景,从而实现执行优化。典型的应用包括:
在 PyTorch 基准测试中,CUDA Graphs 展现了显著的性能提升,尤其是在某些模型上。测试环境为 CUDA 12.0, DGX-H100, Ubuntu 20.04。
在 GROMACS 分子动力学模拟中,CUDA Graphs 也带来了性能提升。测试环境为 CUDA 11.8, DGX-A100, Ubuntu 20.04。
在单 GPU 配置下,对于 GROMACS Water Box 模拟,CUDA Graphs 在不同的系统规模(以千原子数计)下,相比 Streams 均实现了略微的加速,通常在 1.01x 到 1.05x 之间。
在多 GPU 配置下,GROMACS Water Box 模拟显示:
* Graphs Today:当前的 CUDA Graphs 实现也提供了类似的轻微加速。
* Upcoming Release (Projected):预计在未来的版本中,CUDA Graphs 将在更大的系统规模下提供更显著的加速,例如在 768k 原子数时可达约 1.15x 的加速。
在 Aerial 网络与通信框架中,CUDA Graphs 有助于降低延迟和减少抖动。测试环境为 CUDA 11.7, DGX-A100, 1xA100, Ubuntu 20.04。
性能提升主要来源于减少围绕短运行内核的系统开销。
初始情况下,每次内核启动都包含 Grid Initialization 和内核执行(例如 2µs Kernel)。在这种模式下,Grid Initialization 占用了大量时间,导致总开销高达 64%。
通过优化 CPU 侧的启动过程,减少了每次内核启动所需的 CPU 开销。这使得总开销从 64% 降低到 45%。
进一步地,通过优化设备侧的执行开销,例如将 Grid Initialization 流程与内核执行更紧密地集成或预处理,使得总开销进一步降低到 33%。
对于三个 2µs 的内核,总时间缩短了 29%。
CUDA Graphs 实现了两个独立的优化,分别针对 Host 和 Device 上的启动成本。测试环境为 CUDA 12.1, RTX A5000, Intel i7-7800X, Ubuntu 18.04。
这些数据显示了 CUDA Graphs 通过减少启动和重新启动的开销,尤其是在重复执行时,能够带来显著的性能提升。
CUDA Graphs 的主要优势在于对可重复工作负载的优化。通过多次重新启动已实例化的图,可以摊销首次创建和实例化的开销。测试环境为 Straight-Line Graph, CUDA 12.1, RTX A5000, Intel i7-7800X, Ubuntu 18.04。
结论:大多数图需要至少 3-4 次启动才能比流更快,这强调了 CUDA Graphs 在重复性工作负载中的价值。
CUDA Graphs 采用三阶段执行模型来最小化执行开销,尽可能地进行预初始化。
定义 (Define):
实例化 (Instantiate):
执行 (Execute):
一个 CUDA Graph 是由节点组成的,这些节点之间通过依赖关系连接,代表一系列操作。
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 多次启动同一个可执行图,从而利用了图的重复使用优势。
流捕获允许在不重写现有 CUDA 代码的情况下,将一系列流操作转换为一个 CUDA 图,从而利用图的性能优势。
核心步骤:
cudaStreamBeginCapture(&stream1); 启动对指定流的捕获。cudaStreamEndCapture(stream1, &graph); 将捕获到的操作转换为一个 CUDA 图。捕获过程会遵循流间的依赖关系(例如通过 cudaStreamWaitEvent 建立的依赖),以创建图中的分支和连接,从而形成一个有向无环图(DAG)。
可以将库调用捕获为子图,并将其添加到现有图中。
示例步骤:
1. 创建根节点: 通过显式 API (如 cudaGraphAddNode) 为主图创建根节点。
2. 捕获库调用: 使用 cudaStreamBeginCapture 和 cudaStreamEndCapture 将一个 libraryCall 捕获成一个独立的子图 library_graph。
3. 插入子图: 使用 cudaGraphAddSubGraphNode 将捕获到的子图作为节点插入到主图中。
4. 继续构建主图: 继续通过显式 API 构建主图的其余部分。
仅可捕获完全异步的序列。
对于许多应用程序,流捕获需要进行一些调整,因为并非所有代码都能“直接工作”于捕获模式。例如,像 cudaDeviceSynchronize() 这样的同步调用,在捕获中会带来复杂性。
某些操作无法被捕获。
默认(“空”)流无法被捕获。
libraryCall(cudaStreamDefault) 必须替换为指定流的 libraryCall(stream)。
同步调用无法被捕获。
例如,cudaMalloc(...)、cudaMemcpy(...) 和 cudaFree(...) 必须替换为它们的异步版本,如 cudaMallocAsync(..., stream)、cudaMemcpyAsync(..., stream) 和 cudaFreeAsync(..., stream)。
没有异步对应物的调用必须发生在捕获之外。
例如,cudaMallocHost(...) 和 cudaFreeHost(...) 等主机内存操作,由于没有异步版本,需要在 cudaStreamBeginCapture 之前或 cudaStreamEndCapture 之后执行。
流捕获不能同步。
cudaDeviceSynchronize() 虽然可以存在于代码中,但其行为在捕获的图上下文中会有所不同,捕获本身不提供同步机制。
仅仅使代码可捕获就完成了吗?
一段可捕获的代码可能包含如下结构:
cudaMallocHost(...);
cudaStreamBeginCapture(stream, ...);
cudaMallocAsync(..., stream);
cudaMemcpyAsync(..., stream);
hostLogic(...);
libraryCall(stream);
cudaFreeAsync(..., stream);
cudaStreamEndCapture(stream, &graph);
cudaFreeHost(...);
// Instantiate & launch the graph
那么,这段代码现在就图就绪了吗?不一定,还需要进一步优化。
技巧 #1: 通过 cudaMalloc/FreeAsync 将内存管理放入捕获中。
将 cudaMallocAsync 和 cudaFreeAsync 调用放置在 cudaStreamBeginCapture 和 cudaStreamEndCapture 之间,使其成为图的一部分。
技巧 #1.5: 对于其他分配(如 cudaFreeHost(...))—— 保持内存存在,否则在启动时无法访问。这类操作通常应发生在图实例化和启动之外。
技巧 #2: 将任何重要逻辑放入 CPU 回调。
hostLogic(...) 等 CPU 端逻辑,应通过 cudaLaunchHostFunc(stream, hostLogic, ...); 封装,使其能在 GPU 图中异步执行。
技巧 #3: 如果您的应用程序是多线程的且线程独立运行,请考虑使用线程局部捕获模式。
使用 cudaStreamBeginCapture(stream, threadLocal); 来启动线程局部的捕获。
图实例化等同于图的编译步骤:
* 准备并优化图以供执行。
* 实例化后,可执行图结构被锁定。
* 结构性更改需要重新实例化。
与代码编译一样,实例化不是一个简单的步骤,需要额外的时间。而且,像任何编译步骤一样,实例化不会为您处理所有事情。
用户必须为每个节点定义执行位置。
如果图中的一个分支可以在2个GPU上运行,我们如何选择在何处运行什么?
最佳选择可能取决于数据局部性——在执行层是未知的。
执行层不具备执行此操作所需的信息。
* 不能拆分图节点。
* 不能合并图节点。
* 不能重新分配节点的执行位置。
元素操作可以很容易地融合,但仅当操作语义已知时。执行层只看到二进制代码,因此无法执行此合并。
将图工作与其他非图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);
}
启动流仅用于与其他工作排序。
即使图被启动到一个流中,图中的分支仍然会并发执行。
性能提升伴随着限制的增加。
不更新的重新启动。
在实时系统中,图必须保持静态以实现可靠的启动时间。
带更新的重新启动。
图的孤立加速不是全貌。整个应用20-30%可图化,因此整体加速较小。
根据程序结构选择方法。
| 方法 | 描述 | 速度提升 |
|---|---|---|
| 图重新实例化 | 每次迭代重建工作 | 不比流快 |
| 图更新 | 每次迭代更新图 | 比流快达 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);
}
一种更细粒度的参数更新方法。
如果您了解您的工作流程,您可以单独更新节点。
// 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);
}
避免在次要拓扑更改时重新实例化。
节点也可以完全启用/禁用。
// 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);
}
上传步骤可以从启动中分离,以实现更好的流水线操作。
定义 (Define):
实例化 (Instantiate):
上传 (Upload) (3a):
执行 (Execute) (3b):
旨在降低首次启动成本。
原始工作流程 (Original Workflow):
带上传的工作流程 (With Upload):
* 构建图 -> 上传图 -> 额外设置 -> 启动图 -> 启动图 -> 同步。
* 资源初始化在单独的“上传图”阶段进行,从而可以减少首次启动的延迟。
通过减少短运行内核的系统开销来提升性能。
通过关注内核来优化GPU运行时。
Ampere架构为CUDA Graphs带来了新的硬件能力。
主机启动延迟 (Host Launch Latencies) (长度 = 100, CUDA 12.1, Intel I7-7800X, Ubuntu 18.04):
设备启动延迟 (Device Launch Latencies):
建议: 升级到Ampere或更高版本,以受益于新的硬件功能。
在实际运行之前,创建并实例化一个最大尺寸的图:
将首次启动与第二次启动分开进行基准测试:
为工作选择正确的工具:
设备图启动允许图的动态控制流,将图的启动从主机卸载到设备。
CPU部分 (CPU portion):
DeviceLaunch。GPU部分 (GPU portion):
Y,在其内部启动图G2。说明: 实现了设备侧图启动,使得GPU能够自主地启动图,减少主机开销。
一个图只有在满足以下条件时才能从设备启动:
cudaGraphInstantiate(G2, DeviceLaunch))。cudaGraphUpload(G2, ...))。__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)。
依赖解析发生在整个图的粒度上。
图的封装边界是整个启动图。这个边界被称为执行环境 (execution environment)。
图的启动不能在父图内创建新的依赖。图内没有 fork/join 并行。
例如,在 Graph G1 中,如果节点 Y 启动了 Graph G2,那么 G2 成为 Kernel2 的一个依赖,而不是节点 Z 的依赖。
即发即弃模式会立即启动图。启动的图与父图并发运行。
后续工作会隐式地加入即发即弃启动。
即发即弃启动不能直接同步(例如,通过 cudaDeviceSynchronize())。
那么,如何插入工作依赖呢?
尾部启动在调用图完成后按顺序启动。
它提供了一种插入工作依赖的方式。
连接即发即弃启动。
即发即弃启动不能被直接同步(例如,通过 cudaDeviceSynchronize())。
只能通过尾部启动来强制排序。尾部启动会加入即发即弃的工作。
尾部启动队列。
尾部启动按照它们被入队的顺序执行。
在您的应用程序中实现主机无关的循环。
一个图可以重新启动自身,从而实现主机无关的循环。例如,一个 __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
}
}
克服父图封装,将依赖关系转移到上层。
兄弟启动在先行工作和后续工作之间插入子工作。
子工作与父工作并发启动。
子工作成为父图的父图的依赖,但不会阻塞调度图的重新启动。
调度图本身作为尾部启动重新启动。
在主机代码中,程序初始化期间创建多个图。
设备上的调度器内核根据数据包类型选择并启动相应的图来处理传入的数据包。调度器图本身可以作为尾部启动重新启动,以实现持续处理。
代码示例:
__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);
}
开始使用图 (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