Performance Optimization Tutorial, Part 3 [S72686]: CUDA Techniques to Maximize Concurrency and System Utilization

Myrto Papadopoulou (NVIDIA DevTech Compute)
Igor Terentyev (NVIDIA DevTech Compute)
Guillaume Thomas-Collignon (NVIDIA DevTech Compute)

GPU Technology Conference | March 18th, 2025

*** With help of: Akshay Subramaniam, Allard Hendriksen, Athena Elafrou, Ben Pinzone, David Clark

GTC'25 性能优化系列教程

本演示是 GTC'25 性能优化系列教程的一部分，该系列包含以下内容：
- 最大化内存带宽和隐藏延迟的 CUDA 技术 [S72683]
- 最大化计算和指令吞吐量的 CUDA 技术 [S72685]
- 最大化并发性和系统利用率的 CUDA 技术 [S72686]
- 在 Grace-Hopper/Blackwell 上最大化应用程序性能的 CUDA 技术 [S72687]

议程 (Agenda)

• CUDA streams
• Programmatic Dependent Launch (程序化依赖启动)
• CUDA Graphs
• MIG, MPS, 和 Green Contexts
• Cluster Launch Control (集群启动控制)

术语 (Nomenclature)

• CTA (Cooperative Thread Array) == 线程块 (Thread Block)
• CGA (Cooperative Grid Array) == 线程块集群 (Thread Block Cluster)
• High-priority kernel == 与高优先级流关联的内核
• 代码片段 (Code snippets):

namespace cg = cooperative_groups;
using namespace cuda; // cuda::ptx::

CUDA Streams

GPU 上的异步执行

CUDA/GPU 任务（如内核、异步内存操作、主机回调等）与 CPU 是异步执行的。如下图所示，CPU 启动内核后可以继续执行其他代码，而不需要等待 GPU 内核完成。在默认情况下（未使用流），GPU 上的内核是按顺序执行的。

通过使用不同的 CUDA Streams，任务可以在 GPU 上彼此并行执行。这使得 GPU 能够同时处理来自不同流的任务，从而提高利用率。

CUDA Streams 简介

• CUDA Stream 类似于 CPU 线程在 GPU 上的对应物。
• 创建流时，它会与当前活动的 GPU 相关联（例如，通过 cudaSetDevice 设置）。
• 流中的任务按照它们被 CPU 提交的顺序执行。
• 下一个任务在前一个任务完成后开始执行（例外情况：通过 PDL 实现的内核重叠¹）。
• 不同的流可以乱序或并发地执行。

同步 (Synchronization):
- CPU 可以与一个流同步——等待该流中所有之前的任务完成。
- 流之间也可以相互同步——一个流中的下一个任务直到另一个流中的某个特定任务完成后才开始。

---

¹ Programmatic Dependent Launch - 将在本演示的后面部分介绍

默认流 (Default stream)

不同的流可以乱序或并发执行任务。

默认流 (0) 是特殊的:
- kernel<<<grid_size, block_size>>> 等价于 kernel<<<grid_size, block_size, 0, 0>>>。
- 它是在每个上下文中隐式创建的。
- 默认情况下，它不与使用默认标志创建的其他流中的操作重叠。
- 例如：以下内核将不会重叠执行。

kernel_A<<<grid_size, block_size, 0, stream_A>>>();
kernel_B<<<grid_size, block_size, 0, stream>>>(); // 这里的 stream 是默认流 0
kernel_C<<<grid_size, block_size, 0, stream_C>>>();

移除隐式同步:

• 为了实现异步的默认流行为，可以使用编译器选项：

  • nvcc --default-stream per-thread
  • nvcc -DCUDA_API_PER_THREAD_DEFAULT_STREAM=1 ...

• 使用 cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) 创建的流不与默认流同步。

建议:
- 避免使用默认流。
- 使用 cudaStreamNonBlocking 创建显式流。

流优先级 (Stream priorities)

流可以拥有优先级。
- 支持的范围: cudaDeviceGetStreamPriorityRange(...)
- 带优先级的流: cudaStreamCreateWithPriority(...)

优先级为调度任务提供了提示。
- 例如，较高优先级的 CTA 将在已经运行的较低优先级 CTA 完成后立即运行，而剩余的较低优先级 CTA 将在较高优先级的 CTA 完成后运行。

下图展示了将 kernel1_w 启动到低优先级流，随后将 kernel1_h 启动到高优先级流。尽管 kernel1_w 先启动，但 kernel1_h 会抢先执行。

未来的 Nsight Systems 版本将可以直接在 profiler 界面中查看到流的优先级（通过鼠标悬停在流上）。

GPU 同步

重度同步 (Heavy synchronization):
cudaDeviceSynchronize() 会阻塞 CPU，直到所有流上的所有 GPU 任务都完成。

这会导致 GPU 任务之间出现空闲，因为 CPU 无法提前提交新的任务，从而产生“启动延迟间隙 (launch latency gap)”。

一个常见的不良实践是在每次内核启动后调用 cudaDeviceSynchronize()（通常仅为了检查错误或计时）。这会严重影响性能，因为它完全消除了 CPU 和 GPU 之间的并发性。

GPU 同步行为

某些 API 调用是完全阻塞和同步的——如同被 cudaDeviceSynchronize() 包围：
cudaFree(...)

其他一些 API 调用也可以是完全阻塞和同步的——如同被 cudaDeviceSynchronize() 包围：
cudaMemcpy(...)，包括 cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost。
一些其他 API 函数，如 cudaDeviceSetCacheConfig(...) 也具有类似行为。

注：非参考标准，行为可能因硬件、页锁定/非页锁定内存、大小而异。

非阻塞和同步 - 默认流内核行为：

• cudaMemcpy(...) (例如 cudaMemcpyDeviceToDevice)
• cudaMemset(...)

下图展示了 API 调用（CPU 时间线上的红色方块）是非阻塞的，但它在默认流中是同步的，即它会等待流中先前的任务完成后才开始执行（GPU 时间线上的红色圆圈）。

阻塞和异步：

• cudaMalloc(...)

下图显示 API 调用在 CPU 上是阻塞的（红色圆圈），但 GPU 上的内核执行是异步的，可以继续进行。

注：非参考标准，行为可能有所不同。

GPU 同步技巧

• 使用 Nsight Systems 来验证阻塞/非阻塞和同步/异步行为！
• 仅在必要时（存在数据依赖时）使用流同步。

• 通过使用 Async（异步）操作来避免重度同步。

  • cudaMemcpyAsync, cudaMemsetAsync, cudaMallocAsync, cudaFreeAsync
  • 这些函数接受流作为参数，并遵循流语义。
  • 可以是同步的。例如，使用可分页主机内存的 cudaMemcpyAsync。

• 具有流语义的 CPU 任务：cudaLaunchHostFunc(stream, host_fn, data_ptr)。

GPU 错误检查与计时

检查内核错误

• 内核启动错误 (Kernel<<<...>>> launch error):

  • 由于无效的启动配置导致。
  • 如果内核在执行期间失败则不会被设置。
  • 使用 cudaGetLastError() [重置错误], cudaPeekAtLastError() [不重置错误] 来捕获。
  • 由返回执行错误的 API 设置（cudaDeviceSynchronize(), 等）。

• 执行错误 (Execution error):

  • 在内核失败后调用 cudaGetLastError() 和 cudaPeekAtLastError() 不会报告此错误。
  • 这种错误是“粘性”的 (Sticky)。
  • 由 cudaDeviceSynchronize(), cudaStreamSynchronize(...), cudaEventSynchronize(...) 报告。

下图展示了两种错误检查方式的区别：

• 左侧代码块：内核启动成功，但执行时发生错误。cudaGetLastError() 在此之后调用不会报告错误，因为它只检查启动错误。只有在 cudaDeviceSynchronize() 处才会捕获到执行错误。
• 右侧代码块：cudaDeviceSynchronize() 能够正确捕获到异步发生的执行错误，而随后的 cudaGetLastError() 不会再次报告该错误。

计时：

• 使用 Nsight Systems (Nsight Compute)。
• 通过 CUDA 事件。

良好实践：

• 在每次内核<<<...>>>调用后调用 cudaGetLastError() / cudaPeekAtLastError()。
• 检查每个 CUDA API 调用的返回值。

批量异步内存复制

• 常用的 cudaMemcpy/cudaMemcpyAsync：

  • 异步版本可能会同步，同步版本也可能表现为异步...
  • 许多小型复制操作的性能瓶颈主要在于启动延迟。
  • 在一致性（coherent）的 GH/GB 系统上，对于系统分配的内存，性能不是最佳的...

• 用于批量复制的新 API (CUDA 12.0)：

cudaMemcpyBatchAsync(void** dsts, void** srcs, size_t* sizes, size_t count, 
                     cudaMemcpyAttrributes* attrs, size_t* attrIdxs, size_t numAttrs, 
                     cudaStream_t stream)

• GH/GB（Grace Hopper / Grace Blackwell 超级芯片）与可分页内存并发运行。

流同步

cudaStreamSynchronize(stream) 会阻塞 CPU，直到先前提交到该流的所有任务完成。

事件

• 事件是流内的“检查点”。
• 当所有先前的任务在流中完成时，事件才完成。
• 可重用 - 新的 cudaEventRecord(...) 会重置完成状态。

• 事件同步：

  • cudaEventSynchronize(...) - 阻塞直到完成。
  • cudaEventQuery(...) - 非阻塞的完成状态检查。

• 可以获取完成时间戳：

  • cudaEventElapsedTime(start_event, stop_event) - 计算两个已完成事件之间的时间。

由于数据依赖性，同步是必需的。
最常见的情况：
流1中的任务消费由流2中的任务产生的数据（流1中的任务在流2中的任务完成后开始消费数据）。

cudaStreamWaitEvent(stream, event_from_other_stream):

• 流之间的同步“绕过”CPU。
• 非阻塞调用。
• 可以同步来自不同设备的流！

示例 - 小波变换（Wavelet Transform）

该示例展示了如何处理一个需要将一维CPU数组转换为二维CPU数组的任务，其中输出的每一列都是独立计算的。

• 输入: 一维 CPU 数组
• 输出: 二维 CPU 数组，每列独立计算

基础流程

处理流程的目标是，以一个一维CPU数组为输入，经过计算后，生成一个二维CPU数组，其中每一列都是独立计算的。

基础的串行处理流程如下：
1. 循环处理每一列：
1. 步骤 1: 将当前列的数据从CPU复制到GPU。
2. 步骤 2: 在GPU上执行计算。
3. 步骤 3: 将计算结果从GPU复制回CPU。

目标：重叠计算与复制

为了提升性能，核心目标是让计算和数据复制操作能够重叠（并行）执行。理想情况下，当GPU在计算第 i+1 列时，数据传输硬件可以同时在复制第 i 列的结果。

优化策略：使用固定内存（Pinned Memory）

为了实现异步的设备到主机（D2H）数据传输，通常需要使用固定（Pinned）内存作为中转缓冲区。

改进后的流程分为两个复制步骤：

• 步骤 3a: 将GPU上的计算结果异步复制到CPU端的固定内存中。
• 步骤 3b: 在CPU上，将数据从固定内存复制到最终的目标（非固定）内存中（Host-to-Host, H2H）。

注:
- 在X86架构上，需要通过固定缓冲区进行复制。
- 在Grace Hopper/Grace Blackwell (GH/GB) 架构上，如果使用 cudaMemcpyBatchAsync(...)，则可能不需要固定缓冲区。

优化策略：使用多流（奇偶列）

为了进一步提高并行度，可以使用多个CUDA流。例如，可以创建两个独立的流，一个处理所有奇数索引的列，另一个处理所有偶数索引的列，从而形成两条并行的处理流水线。

• 处理偶数列的流水线：
  

• 处理奇数列的流水线：
  

性能分析与实现方法

使用NVIDIA Nsight Systems工具可以观察到，通过流技术，内存复制（Memcpy DtoH, memcpy_h2h）和内核计算（void runna...）操作可以在时间上重叠执行。

异步CPU任务（如H2H复制）通常可以通过专用的CPU线程和同步原语（如互斥锁）来实现。但另一种更高效的方法是使用CUDA Streams结合 cudaLaunchHostFunc(...)，它允许将一个CPU函数调用插入到CUDA流中，由CUDA运行时在适当的时候回调执行。

三流流水线实现

一个更精细的实现是使用三个独立的流来管理流水线的不同阶段：
- stream_cpt：用于GPU计算。
- stream_d2h：用于设备到主机（固定内存）的异步复制。
- stream_h2h：用于主机端的回调函数，执行从固定内存到最终目标内存的复制。

代码实现框架：

在循环中，为每一列提交异步任务到各自的流中。使用 [col & 1] 的方式来实现双缓冲（ping-pong buffering），交替使用两个缓冲区。

for (int col = 0; col < ncol; ++col) {

    // 在计算流上启动GPU内核
    kernel<<<...>>>(d_out[col & 1], d_in, ..., stream_cpt);

    // 在D2H流上启动异步内存复制（GPU -> Pinned Memory）
    cudaMemcpyAsync(h_pin[col & 1], d_out[col & 1], ..., stream_d2h);

    // 在H2H流上启动一个主机函数，用于CPU端内存复制
    cudaLaunchHostFunc(stream_h2h, fn_h2h, &pars_h2h);
}

主机回调函数：

cudaLaunchHostFunc 调用的主机函数 fn_h2h 负责执行从固定内存到最终输出数组的 memcpy 操作。参数通过一个结构体 Pars_h2h 传递。

使用事件（Events）进行流间同步

为了确保流水线中各个阶段按正确的依赖关系执行（例如，必须在计算完成后才能开始复制结果），需要使用CUDA事件进行流间的显式同步。

• 依赖关系:
  1. D2H复制（stream_d2h）必须等待计算（stream_cpt）完成。
  2. H2H复制（stream_h2h）必须等待D2H复制（stream_d2h）完成。
  3. 为了防止下一轮计算覆盖正在被D2H复制操作读取的缓冲区，计算（stream_cpt）需要等待前一轮的D2H复制完成。

带事件同步的完整代码:

通过 cudaEventRecord 记录事件，并通过 cudaStreamWaitEvent 来让一个流等待另一个流中的事件。

CUDA Streams 使用技巧总结

• 避免重量级同步：例如，避免全局性的 cudaDeviceSynchronize()。
• 创建非阻塞流：使用 cudaStreamNonBlocking 标志创建流，并避免使用默认流（默认流具有同步行为）。
• 最小化同步：利用流的依赖关系进行数据同步。
• 使用异步API：例如，使用 cudaMemcpyBatchAsync。
• 创建高性能事件：如果不需要计时功能，使用 cudaEventDisableTiming 标志创建事件以获得更好的性能。
• 环境变量:
  • CUDA_DEVICE_MAX_CONNECTIONS：限制计算和复制引擎的并发连接数。
  • CUDA_DEVICE_MAX_COPY_CONNECTIONS：限制复制引擎的并发连接数。
  • CUDA_SCALE_LAUNCH_QUEUES：启动队列大小的缩放因子。

---

Programmatic Dependent Launch (PDL)

动机 (Motivation)

在标准的 CUDA 流执行模型中，流语义保证了内核的顺序执行。然而，内核之间的数据依赖关系通常是隐式的。如下图所示，Consumer kernel 3 依赖于 Producer 内核（primary 和 secondary）产生的数据。

Consumer kernel 的执行通常可以分为两个阶段：

1. 前序（Preamble）：这部分工作不依赖于 Producer 内核的输出。例如，共享内存初始化、指针运算、其他设置工作、从全局内存读取只读数据等。
2. 处理阶段：处理由 Producer 内核生成的数据。

在传统的流执行中，Consumer kernel 必须等待 Producer 完全结束后才能开始，即使其前序部分可以提前执行。

Programmatic Dependent Launch (PDL) 允许 consumer kernel 的前序部分与 producer kernel 的执行机会性地重叠，从而提高硬件利用率。

---

使用方法 (CUDA streams)

PDL 通过在生产者和消费者内核中插入特定的设备 API 来协调它们的执行。

关键时间点:
- 在 primary (生产者) 内核中，指示 secondary (消费者) 内核可以被触发的时间点。
- 在 secondary (消费者) 内核中，指示内核应该阻塞并等待 primary 内核完成的时间点。

设备 API (适用于 CC >= 9.0):
- cudaTriggerProgrammaticLaunchCompletion
- 位置：在 primary (生产者) 内核中调用。

• cudaGridDependencySynchronize
  • 位置：在 secondary (消费者) 内核中调用。

约束条件:

• secondary (消费者) 内核必须通过 cudaLaunchKernelEx 启动。
• 在同一个流上，primary 和 secondary 内核之间不能存在任何其他 GPU 工作（例如，不能有 CUDA 事件记录）。

---

使用方法 - 内核启动（CPU端）

下图展示了如何在主机端（CPU）代码中启动使用 PDL 的内核。

• 内核签名:

__global__ void primary_kernel(uint8_t* d_ptr);
__global__ void secondary_kernel(uint8_t* d_ptr);

• 启动代码:
  • primary_kernel 的启动方式保持不变。

primary_kernel<<<grid_dim, block_dim, 0, strm>>>(d_ptr);

- `secondary_kernel` 的启动需要使用 `cudaLaunchKernelEx`。
    - 首先，配置常规的启动参数（`blockDim`, `gridDim`, `dynamicSmemBytes`, `stream`）。
    - 然后，需要设置一个特殊的属性 `cudaLaunchAttributeProgrammaticStreamSerialization`，并将其值 `programmaticStreamSerializationAllowed` 设为 1。
    - 最后，使用 `cudaLaunchKernelEx` 启动内核。

---

使用方法 - PDL API（设备代码）

下图展示了如何在设备端（GPU）内核代码中使用 PDL API。

primary_kernel (生产者):

__global__ void primary_kernel(uint8_t* d_ptr) {
    work_A();
    cudaTriggerProgrammaticLaunchCompletion();
    work_B();
}

• secondary_kernel 可以在以下任一条件满足时被调度：

  • primary 内核中每个未退出的 CTA（Cooperative Thread Array）都至少调用了一次 cudaTriggerProgrammaticLaunchCompletion() API。
  • 如果 primary 内核中没有未退出的 CTA 调用该 API，则 secondary_kernel 可以在 primary 的所有 warp 完成后被调度。

• cudaTriggerProgrammaticLaunchCompletion() 不提供内存可见性保证。

secondary_kernel (消费者):

__global__ void secondary_kernel(uint8_t* d_ptr) {
    work_C();
    cudaGridDependencySynchronize();
    work_D();
}

• 一个线程在调用 cudaGridDependencySynchronize 时会阻塞，并等待 primary 内核完成（包括 work_B）。
• work_C 应该可以安全地与 work_B 并行执行。

• work_C 不应与 primary_kernel 的工作（work_A 或 work_B）有数据依赖。例如：

  • 不应访问由 work_A 或 work_B 修改的数据。
  • 不应修改由 work_A 或 work_B 访问的数据。
  • 读取也被 work_A 或 work_B 读取的只读数据是可以的。

• 典型的 work_C 包括：局部计算、共享内存初始化、从全局内存读取只读数据等。

• work_D 通常处理由 primary 内核产生的数据。

---

使用技巧 (Tips)

• 过早触发 cudaTriggerProgrammaticLaunchCompletion 的性能考量？

  • 如果 duration(work_B) > duration(work_C)，那么 secondary_kernel 可能会长时间等待 cudaGridDependencySynchronize()。
  • 这种情况在其他流上的内核本可以利用 secondary_kernel 占用的 SMs 时尤为重要。

• 如果在 work_C 之后调用的 cudaGridDependencySynchronize() 不正确地访问了由 work_B 修改的数据会怎样？

  • 这是无效的使用方式，并可能导致竞争条件（race condition）。
  • 工具无法捕获此类错误，因此需要谨慎使用。
  • 技巧: 在 secondary_kernel 中，对于常见用例，在任何全局内存访问之前调用 cudaGridDependencySynchronize。

• 如果 work_B 和 work_C 是空的会怎样？

  • 几乎没有重叠（纳秒级别的好处），但仍然可以使用。

• 兼容性:

  • 使用宏来保护 PDL API 的调用，以确保代码在不支持的架构上也能编译通过：
    #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)

---

3个内核的使用示例

在更复杂的依赖链中，一个内核可以同时作为消费者和生产者。

• primary_kernel 是生产者。
• secondary_kernel 是 primary_kernel 的消费者，同时也是 another_kernel 的生产者。
• another_kernel 是 secondary_kernel 的消费者。

---

GPU 时间线 (来自 Nsight Systems)

下图展示了不同情况下 PDL 的 GPU 执行时间线。

• 无 PDL: secondary_kernel 在 primary_kernel 完全结束后才开始执行。

• PDL (所有 CTA 调用 trigger): secondary_kernel 在 primary_kernel 执行期间就开始了，实现了重叠。

  • 如果 primary 是多波次（multi-wave）执行的，secondary 只有在最后一波 CTA 调用了 trigger 后才能被调度。

• PDL (隐式触发): 当 secondary 在末尾被隐式触发时（例如，没有 CTA 调用 trigger），也会有执行重叠。

代码示例可以在 https://github.com/NVIDIA/cuda-samples/tree/master/Samples/5_Domain_Specific/programmaticLaunch 找到。

---

进一步阅读

• CUDA C++ 编程指南中的 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programmatic-dependent-launch 部分，也包括了如何在 CUDA graphs 中使用 PDL。
• CUDA runtime API 参考: cudaTriggerProgrammaticLaunchCompletion 和 cudaGridDependencySynchronize。
• 相关的 PTX 指令 (如果在该级别使用): https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol。
• CUB 的归并排序可以使用 PDL。

---

CUDA Graphs

一种工作提交模型

• 图 (Graph) = 操作 (operations) & 操作间的依赖关系 (dependencies)。

• 操作类型:

  • 内核 (kernel)、内存拷贝 (memory copy)、内存设置 (memset)、条件节点 (conditional node)、子图 (child graph)、空节点 (empty node)、CPU 函数调用 (CPU function call)、异步内存分配 (async. memory allocation) 等。

• 依赖关系:

  • 在上图示例中，A, B, D, E 必须按此顺序执行。A, C, D, E 也必须按此顺序执行。
  • B 和 C 可以并行执行，或以任何顺序执行，但它们都必须在 A 之后和 D 之前执行。

• 性能优势:

  • 减少 CPU 启动提交开销（尤其是在重复启动同一个图时）。
  • 在 GPU/驱动端启用潜在的优化，因为工作是作为一个整体提交的。
  • 启用其他优化：条件执行（条件在运行时改变）等。

---

减少 CPU 启动提交开销

• CPU 时间线 (包括 CUDA API 调用):

  • Streams: 启动一系列内核需要多次调用 CUDA API，这会产生显著的 CPU 开销。
  • Graphs: 使用 cudaGraphLaunch 单次调用即可启动 CUDA graph 中封装的所有工作，从而节省了 CPU 时间。

• 让 CPU 领先于 GPU，以避免因 CPU 成为关键路径而导致的 GPU 空闲。对于短内核，CPU 开销的影响更为显著。

---

等效的流提交时间线

使用传统的流和事件来实现与上图相同的依赖关系会非常复杂。

• 依赖 A-C: 在 stream 1 上记录 event_end_A，然后让 stream 2 等待 event_end_A。
• 依赖 C-D: 在 stream 2 上记录 event_end_C，然后让 stream 1 等待 event_end_C。

CUDA Graphs 抽象了这种复杂的依赖管理。

---

减少 CPU 启动提交开销 (第二次启动)

下图比较了第二次启动相同工作时，基于流和基于图的 CPU 和 GPU 时间线。

• CPU 时间线:

  • Streams: 每次启动都需要完整的 API 调用序列，开销与第一次相同。
  • Graphs: 第二次启动只需一次 cudaGraphLaunch 调用，CPU 开销极低，节省了大量时间。

• GPU 时间线:

  • GPU 上的执行模式是相似的，但图的启动作为一个整体单元 (Graph-4 (GraphExec 1)) 出现。

注意: 时间线来自 Nsight Systems，使用默认的 --cuda-graph-trace graph 模式。节点级别的追踪可以通过 --cuda-graph-trace node 实现，但这可能会带来显著的开销。

---

关键步骤：定义、实例化、启动

1. 定义图 (Define graph)

   • 将操作和依赖关系封装在一个 cudaGraph_t 图中。
   • 两种方式：(a) 流捕获 (stream capture) 或 (b) 使用图 API。

2. 实例化图 (Instantiate graph)

   • 实例化图模板，生成一个可执行图 cudaGraphExec_t。
   • cudaGraphInstantiate(&graph_exec, graph)

3. 启动图 (Launch graph)

   • 可在 CUDA 流上启动可执行图：cudaGraphLaunch(graph_exec, stream)
   • 流仅用于依赖跟踪，不提供关于图节点在何处执行的信息。
   • 可选：在启动前可将图上传到流：cudaGraphUpload(graph_exec, stream)

4. 后续操作

   • 可以根据需要多次重新启动同一个可执行图。
   • 内核参数是否改变？可以使用 cudaGraphExecKernelNodeSetParams 从 CPU 端更新节点。
   • 是否需要不执行某些节点？如果已知，可以在图启动前使用 cudaGraphNodeSetEnabled() 禁用节点。
   • 发生较大变化？重新实例化图。

5. 销毁图 (Destroy graph)

   • 销毁可执行图和图模板。

---

CUDA 图创建（2 种方式）

1. 流捕获 (Stream Capture)

通过捕获在 CUDA 流上执行的工作来创建图。

代码示例：

• cudaStreamBeginCapture(strm1, cudaStreamCaptureModeGlobal): 开始在 strm1 上捕获操作，以构建图。
• kernel_A<<<1, 32, 0, strm1>>>(): 在 strm1 上启动核函数 A，此操作被捕获到图中。
• cudaEventRecord(event_end_A, strm1): 记录事件，同样被捕获。
• kernel_C<<<1, 32, 0, strm2>>>(): 在 strm2 上启动核函数 C。这里 strm2 是捕获的一部分，它会派生/加入依赖于 strm1。
• cudaStreamEndCapture(strm1, &graph): 结束在 strm1 上的捕获，并将捕获到的操作序列整合成一个图对象 graph。

要点:
- 在流捕获期间，GPU 不执行工作；工作仅被捕获到图中。
- 重要提示：在核函数启动后（包括捕获期间），不要跳过 cudaGetLastError() 调用。如果核函数启动包含无效参数（例如，不支持的网格大小、动态共享内存等），你可能会在图中遇到静默丢失的核函数。
- 免责声明：为简洁起见，幻灯片中的代码示例省略了错误检查代码。

---

2. 使用图 API (Use graph APIs)

手动、显式地创建图节点并定义它们之间的依赖关系。

代码示例：
1. 创建图和节点:

cudaGraph_t graph;
cudaGraphCreate(&graph, 0);
cudaGraphNode_t node_A, node_B, node_C, node_D, node_E;
cudaKernelNodeParams params[5] = {};
// <...> 填充核函数节点参数

1. 添加根节点 (Node A):

cudaGraphAddKernelNode(&node_A, graph, nullptr, 0, &params[0]);

- `&node_A`: 要添加的图节点。
- `graph`: 要添加到的图。
- `nullptr`, `0`: 该节点的依赖项数量（0 表示是根节点）。
- `&params[0]`: 参数。

1. 添加依赖节点 (Node B):

cudaGraphAddKernelNode(&node_B, graph, &node_A, 1, &params[1]);

- `&node_A`, `1`: `node_B` 依赖于 `node_A`。

代码示例（续）:

• node_C 同样依赖于 node_A。
• node_D 依赖于 node_B 和 node_C：

std::vector<cudaGraphNode_t> node_deps = {node_B, node_C};
cudaGraphAddKernelNode(&node_D, graph, node_deps.data(), node_deps.size(), &params[3]);

• node_E 依赖于 node_D。

---

CUDA 图创建方式的选择

这取决于具体情况，需要考虑一些权衡：

• 流捕获 (Stream Capture)

  • 优点:

    • 如果代码已经写好，这可能是利用图的最快方式。
    • 允许捕获库调用（例如，到一个子图中）。

  • 缺点:

    • 并非所有工作都是可捕获的（例如，可能需要将同步调用、CPU 逻辑转换为 cudaLaunchHostFunc）。
    • 更新节点参数需要更多工作（需要先获取你需要的节点）。

• 使用图 API 创建图（手动）

  • 优点:

    • 更容易表达依赖关系（无需流、事件等）。
    • 更容易探索不同的拓扑/扩展图，只需更新节点的依赖关系。
    • 更容易更新节点参数或启用/禁用节点（因为节点已经预先知道）。

  • 缺点:

    • 对于现有代码，需要维护一个单独的代码路径来同时支持流和图。

---

CUDA 图技巧与提醒

这是一个远未完整的列表。

• cudaGraphDebugDotPrint：用于可视化图。

  • 例如：cudaGraphDebugDotPrint(graph, "graph", 0 /*或 cudaGraphDebugDotFlagsVerbose */)
  • 可以将 dot 文件转换为 pdf；可以使用 c++filt 来 demangle 名称。

• cudaGraphGetNodes：获取图中节点的数量或列表。

  • 对于流捕获的图特别有用，可与 cudaGraphNodeGetType, cudaGraphKernelNodeGetParams 等一起使用。

• cudaGraphExecKernelNodeSetParams：更新图执行实例中核函数节点的参数。

• cudaGraphNodeSetEnabled：在图执行中启用/禁用一个节点。
• 如果可能，在你的关键路径之外创建和实例化你的图。在关键路径之外进行图的更新和启动。

---

启用条件执行

动机

处理过程可能依赖于在某些（GPU）工作处理之后才知道的运行时条件。

示例:
- 如果你的数据具有某些特征，则进行额外的处理。
- 如果你已经得到了一个足够好的答案，跳过后续的处理。
- 如果你的处理时间过长，提前退出。

---

传统方法的瓶颈

CPU 评估条件并决定接下来要启动什么。

如上图所示，当控制流返回到 CPU 进行条件判断时：
1. CPU 启动初始数据处理。
2. CPU 等待其完成。
3. CPU 评估条件 A，然后启动算法 1。
4. CPU 再次等待完成。
5. CPU 评估条件 B，然后启动后续工作。

潜在问题:
- CPU 无法远超前于 GPU。
- GPU 时间线上出现间隙 (Gap)，导致 GPU 空闲。
- 在关键路径上有启动开销。

---

在 GPU 上评估条件？

一个自然的想法是：如果我们可以在 GPU 上评估条件会怎么样？

---

朴素的 GPU 条件评估方法

无条件地启动所有 GPU 工作，并让 GPU 在每个核函数的序言（prologue）中评估条件。

• CPU 时间线：CPU 无条件地启动所有工作，不再需要等待。

• GPU 时间线：

  • 条件为真时，所有工作正常执行。
  • 条件为假时，在每个核函数的序言中评估条件，如果为假则提前退出。

• 问题: 如果处理不仅仅是核函数（例如，包含 memcpy 操作）怎么办？

潜在问题:
- 不可扩展：每个核函数都需要被修改。
- 不适用于非核函数工作。

---

使用图封装条件工作

将依赖于运行时条件的工作封装到一个条件节点的体图 (body-graph of a conditional node) 或一个设备启动的图 (device-launched graph) 中。

• CPU 时间线：CPU 评估条件 A，然后将算法 1 作为图启动。
• GPU 时间线：
  • 图中的一个核函数/节点评估条件 B。
  • 根据条件 B 的结果，有条件地启动/执行一个子图。
  • 如果条件为真，执行额外处理；如果为假，则跳过。
  • 这样避免了 GPU 空闲和返回 CPU 的开销。

---

启用条件执行总结

方法	优点	缺点
返回 CPU，评估条件并启动相应工作	无需修改 GPU 核函数。	CPU 等待 GPU 完成；无法远超前；GPU 时间线出现间隙；关键路径上有启动开销。
无条件启动所有 GPU 工作，并在 GPU 上评估条件	CPU 不在关键路径上。无 GPU 间隙。	需要修改每个受影响核函数的序言以提前退出。扩展性差。核函数的序言应该总是执行。不适用于非核函数工作。
将条件工作封装到条件节点的体图或设备启动的图中	CPU 不在关键路径上。无 GPU 间隙。无需修改 GPU 核函数；工作不限于核函数。	可能需要添加额外的 join/fork 图节点。

---

条件图节点 (Conditional Graph Nodes)

概述

• 一个条件节点包含：

  • 一个类型 (type)。
  • 一个条件 (通过条件句柄访问)。
  • 一个或多个与之关联的体图 (body graphs)。

• 时间线:

  • 核函数 A 执行。

  • 条件节点 B 执行：

    • 检查条件，并根据节点类型和条件值执行相应的体图（如果有的话）。

  • 核函数 C 在适用的体图完成之后执行。

---

条件节点类型及其体图

• IF: cond != 0 时执行。需要 1 个体图。
• IF/ELSE: cond != 0 时执行一个分支，cond == 0 时执行另一个。需要 2 个体图。
• WHILE: cond != 0 时循环执行。需要 1 个体图。
• SWITCH: 根据 cond 的值选择一个分支执行。需要 N 个体图（对应 N 个 case 语句）。

---

控制条件

通过 cudaGraphConditionalHandle 访问条件。

• 创建一个句柄:

cudaGraphConditionalHandle cond_handle;
cudaGraphConditionalHandleCreate(&cond_handle, graph, default_value, flags);

- `graph`: 使用 `cudaGraphCreate()` 创建的图。
- `default_value`: 可选，应用于每次图启动。
- `flags`: 0 (无默认值) 或 `cudaGraphCondAssignDefault` (使用默认值)。

• 在创建后设置条件:
  这是一个设备函数，只能从一个线程中调用。

__global__ void upstream_kernel(cudaGraphConditionalHandle handle, unsigned int new_cond_value, ...) {
    if (threadIdx.x == 0) {
        cudaGraphSetConditional(handle, new_cond_value); // device only function
    }
}

---

创建一个条件节点

1. 为此图创建一个条件句柄。

cudaGraphConditionalHandle cond_handle;
cudaGraphConditionalHandleCreate(&cond_handle, graph, default_value, flags);

1. 创建并添加一个与此句柄关联的条件节点到图中。

cudaGraphNodeParams params = {cudaGraphNodeTypeConditional};
params.conditional.handle = cond_handle; // 之前创建的句柄
params.conditional.type = cudaGraphCondTypeIf; // 或 ...While, ...Switch
params.conditional.size = 1; // 体图的数量，取决于节点类型
cudaGraphAddNode(&cond_node, graph, cond_node_deps.data(), cond_node_deps.size(), &params);

1. 填充条件节点的体图，通过 params.conditional.phGraph_out[i] 访问。

// 将节点添加为 cond 的体图的根节点
cudaGraphAddNode(&node, params.conditional.phGraph_out[0], nullptr, 0, &nodeParams);

1. 确保条件由一个上游核函数的一个线程填充。

---

设备图 (Device Graphs)

概述

• 设备图可以从设备端启动（也可以从主机端启动）

• 通过在一个特殊流上的图中的内核调用 cudaGraphLaunch 来启动设备图

  • 网格中只有一个线程应该启动该图！
  • 支持的流：cudaStreamGraphFireAndForget, cudaStreamGraphTailLaunch, cudaStreamGraphFireAndForgetAsSibling
  • 你可以控制图是否被启动

• 设备图:

  • 在图实例化 (cudaGraphInstantiate) 期间需要一个特殊标志 cudaGraphInstantiateFlagDeviceLaunch

  • 与主机图相比有额外的限制

    • 例如，一个设备图只能包含内核/内存拷贝/内存设置节点和子图节点等。如果设备图包含条件节点，图实例化将返回错误。

  • 在从设备启动之前需要上传到设备

  • 上传选项：显式调用 cudaGraphUpload，作为实例化的一部分通过特殊标志上传，或从主机进行一次额外的启动。

参考文献：
[1] https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__GRAPH.html#group__CUDART__GRAPH_1g0b72834c2e8a3c93c443c6c67626d0d9
[2] https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-graph-creation

CUDA Graphs 进一步阅读

• CUDA 编程指南中的 CUDA Graphs
• 图管理 API 参考: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__GRAPH.html 和 https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__GRAPH.html

• NVIDIA 技术博客文章和 GTC 演讲 (非完整列表):

  • https://www.nvidia.com/en-us/on-demand/session/gtcspring23-dlit52330/
  • https://developer.nvidia.com/blog/getting-started-with-cuda-graphs/
  • https://developer.nvidia.com/blog/constant-time-launch-for-straight-line-cuda-graphs-and-other-performance-enhancements/
  • https://developer.nvidia.com/blog/dynamic-control-flow-in-cuda-graphs-with-conditional-nodes/
  • If, while (自 CUDA 12.4), if-else 和 switch (自 CUDA 12.8)
  • https://developer.nvidia.com/blog/constructing-cuda-graphs-with-dynamic-parameters/
  • https://developer.nvidia.com/blog/optimizing-llama-cpp-ai-inference-with-cuda-graphs/
  • https://developer.nvidia.com/blog/optimizing-drug-discovery-with-cuda-graphs-coroutines-and-gpu-workflows/

• 相关的 CUDA 示例:

  • https://github.com/NVIDIA/cuda-samples/tree/master/Samples/simpleCudaGraphs, https://github.com/NVIDIA/cuda-samples/tree/master/Samples/jacobiCudaGraphs, https://github.com/NVIDIA/cuda-samples/tree/master/Samples/graphMemoryNodes, https://github.com/NVIDIA/cuda-samples/tree/master/Samples/graphMemoryFootprint, https://github.com/NVIDIA/cuda-samples/tree/master/Samples/graphConditionalNodes, https://github.com/NVIDIA/cuda-samples/tree/master/Samples/cudaGraphsPerfScaling

---

MIG, MPS, 和 Green Contexts

GPU 资源分区机制

动机

多实例 GPU (Multi-Instance GPU, MIG) 的案例

• 问题: GPU 越来越大，资源可能未被充分利用。
• 解决方案: 根据您的用例，将一个 GPU 分割成多个“较小的 GPU”，以提高不同应用程序同时使用时的资源利用率。
• 结论: MIG (Multi-Instance GPU) 可以提高您的 GPU 利用率。

多进程服务 (Multi-Process Service, MPS) 的案例

• 问题: 多个进程在默认计算模式下（无 MPS）同时运行时，由于时间分片和上下文切换开销，GPU 利用率可能很低。
• 解决方案: 使用 MPS，多个进程可以像在不同流中的单个进程一样并发运行，从而提高了整体 GPU 利用率。
• 结论: MPS (Multi-Process Service) 可以提高您的 GPU 利用率。

绿色上下文 (Green Contexts, GCs) 的案例

• 问题: 一个长时间运行的独立任务（任务 A）占用了所有 GPU 资源，导致一个时间关键型任务（任务 B）被延迟。
• 解决方案: 使用绿色上下文对 SMs 进行静态资源分区（例如 80%-20% 的划分），确保关键任务 B 能够及时获得资源并更快地完成。
• 结论: 绿色上下文的静态资源分区使关键工作 B 能够更早完成。

多租户选项 (单 GPU)

资源分区机制（可以组合使用）

机制	MIG (多实例 GPU)	MPS (多进程服务)	Green Contexts (绿色上下文)
示例
类型	静态地将 GPU 分区为多个 MIG 实例（“较小的 GPU”）	动态分区	SMs 的静态分区
目标	不同的应用程序可以使用不同的 MIG 实例	主要针对不同的进程	分区发生在应用程序内部
配置	在应用程序启动前配置	需要 MPS 服务运行	在应用程序启动前无需额外服务或配置

多实例 GPU (MIG)

概述

• 在支持的 GPU（计算能力 >= 8.0）上可以创建多个 MIG 实例（物理分区）。
• 每个实例根据所使用的 MIG 配置文件拥有预定比例的 GPU 资源。
  • 包括 SMs、数据交换单元 (Data xbar)、L2 缓存和内存 (DRAM)。

• 一个 MIG 配置文件 (MIG Xg.Ygb) 由 X 个计算切片 (SMs) + 内存切片 (L2, Mem.) 和 Y GB 总内存组成。

  • 示例：1g.12gb, 4g.48gb。

• MIG 提供（在 MIG 实例之间）：SM 性能隔离、错误隔离、内存带宽 QoS、内存保护。

如何使用

• 为 GPU 启用 MIG 模式：如果支持，使用 sudo nvidia-smi -i <GPU> -mig 1

  • 需要系统管理员权限；传递 0 以在最后禁用 MIG 模式。

• 列出支持的 GPU 实例配置文件：nvidia-smi mig -lgip

  • 来自 GH200 的部分输出（最小 + 最大配置文件）如下：

• 创建一个或多个实例：sudo nvidia-smi mig -cgi <profile ID1, ID2, ...> -C
• 可以列出设备：nvidia-smi -L

• 在特定实例上运行：

  • CUDA_VISIBLE_DEVICES=1 ./example_app 或 CUDA_VISIBLE_DEVICES=MIG-<UUID> ./example_app

• 销毁所有 MIG 实例：sudo nvidia-smi mig -dci && sudo nvidia-smi mig -dgi 并禁用 MIG 模式。

总结

项目	描述
分区类型	静态（仅 GPU 资源；不包括 PCI-e）
何时启用/配置	在原始 GPU 上启动任何应用程序之前
配置选项	使用的 MIG 配置文件；会影响应用程序性能
是否需要更改应用程序	否
使用案例	多用户或单用户运行不同应用程序且 GPU 未充分利用的情况、云服务提供商 (CSPs)；需要 QoS 和隔离

参考/进一步阅读:
- MIG: https://www.nvidia.com/en-us/technologies/multi-instance-gpu/
- MIG 用户指南: https://docs.nvidia.com/datacenter/tesla/mig-user-guide/index.html
- GTC 2022 演讲: "Optimizing GPU Utilization: Understanding MIG and MPS"
- NVIDIA Ampere 架构白皮书, "MIG (Multi-Instance GPU) Architecture" 部分: nvidia-ampere-architecture-whitepaper.pdf
- NVIDIA H100 Tensor Core GPU 架构, "第二代安全 MIG" 部分。

---

MPS (多进程服务)

MPS (Multi-Process Service) 概述

• MPS 允许多个进程同时在 GPU 上运行，无需时间分片。

  • 效果如同由单个进程在不同流中提交工作。

• 默认情况下，没有资源的静态分区或 QoS。

  • MPS 客户端竞争所有 SM 资源。

• 没有错误隔离（与 MIG 不同）。

  • 例如，一个非法的内存访问错误会影响所有进程。

• 每个物理 GPU 最多支持 48 个 MPS 客户端 (Volta+)，具体取决于：

  • CUDA_DEVICE_MAX_CONNECTIONS 环境变量、每个客户端的内存需求等。
  • 如果尝试连接的客户端超过支持数量，可能会收到 "CUDA-capable device(s) is/are busy or unavailable" 错误。

• 要使用 MPS，请在启动 MPS 客户端之前启动 MPS 守护进程。

如何使用 MPS

• 设置适当的环境变量:

  • export CUDA_VISIBLE_DEVICES=0 选择要使用的 GPU；0；也可以指定 GPU-UUID (包括 MIG 实例)。
  • export CUDA_MPS_PIPE_DIRECTORY=<accessible pipe path> 默认目录是 /tmp/nvidia-mps。
  • export CUDA_MPS_LOG_DIRECTORY=<accessible log path> 默认目录是 /var/log/nvidia-mps。

• 建议将相关 GPU 计算模式设置为 exclusive:

  • sudo nvidia-smi -i 0 -c EXCLUSIVE_PROCESS
  • 可能的 -c 或 --compute-mode 选项是 0/DEFAULT, 2/PROHIBITED, 3/EXCLUSIVE_PROCESS。
  • 如果您的 GPU 处于 EXCLUSIVE_PROCESS 模式且没有 MPS 服务在运行，那么一个进程将成功启动一个 GPU 内核，而对于其他进程，"cudaErrorDevicesUnavailable=46" 将被返回。

• 启动守护进程并利用 MPS:

  • sudo nvidia-cuda-mps-control -d
  • ...可以运行进程...
  • 如果一个应用程序已启动，您会看到一个由 CUDA MPS 控制守护进程启动的 nvidia-cuda-mps-server 进程，在 nvidia-smi 下。
  • 完成后关闭守护进程：sudo echo quit | sudo nvidia-cuda-mps-control。

• 从 CUDA 12.4 开始，您可以以编程方式检查此进程是否为 MPS 客户端，通过 mpsEnabled = 0; cuResult res = cuDeviceGetAttribute(&mpsEnabled, CU_DEVICE_ATTRIBUTE_MPS_ENABLED, device);

  • 假设 res = CUDA_SUCCESS，如果此进程是 MPS 客户端，mpsEnabled 将为 1。

使用 MPS 进行资源调配

何时使用

• 默认情况下，不同的 MPS 客户端会竞争所有 GPU 资源，如 SMs、内存等。

  • 如果进程 A 和 B 各自都未充分利用 GPU -> MPS 是完美的选择！
  • 如果 A 和 B 各自都充分利用 GPU，但没有延迟敏感性 -> 使用 MPS 时单个进程的持续时间可能更长，但它们仍然可以重叠。
  • 如果 A 和 B 合起来充分利用 GPU，并且至少有一个是延迟敏感的 -> 延迟敏感的应用可能会受到影响，没有 QoS。
  • 解决方案：MPS 资源调配。

• MPS 资源调配（活动线程百分比）对客户端进程可以使用的 SM 百分比设置了上限。

  • 可以为延迟敏感的案例提供一些 QoS，防止一个 MPS 进程（即其 GPU 工作）使用所有 SMs。

MPS 的活动线程百分比 (Active Thread Percentage)

如何设置

• 通过 CUDA_MPS_ACTIVE_THREAD_PERCENTAGE 环境变量设置（最高 100.0）。

设置时机	影响范围	注意事项
启动 MPS 控制守护进程之前	所有未来的 MPS 客户端	通过 sudo -E echo get_default_active_thread_percentage | sudo -E nvidia-cuda-mps-control 查询默认值
启动 MPS 客户端时 CUDA_MPS_ACTIVE_THREAD_PERCENTAGE=80 ./app	该客户端进程	限制不能大于 MPS 控制守护进程强制执行的限制

• cudaDevAttrMultiProcessorCount 属性将显示此进程的 活动线程百分比 * GPU 总 SMs 限制。
• 使用什么值？ 考虑每个进程独立运行时的活动线程百分比，以及所有进程的 ∑(活动线程百分比)，并根据您的工作负载进行实验。
  • 提醒：∑(活动线程百分比) > 100 & 单个活动线程百分比 < 100 => 超额订阅，同时没有单个进程可以使用所有 SMs。
  • 如果是同构处理，使用统一的百分比可能是一个好方法。
  • 如果应用程序即使在一起也未充分利用 GPU -> 保留默认值 (100)。
  • 如果没有应用程序是延迟敏感的 -> 可以保留默认值或设置为 < 100 以确保没有进程能使用所有 SMs。
  • 如果一个或多个进程是延迟敏感的 -> 考虑避免超额订阅以保持一定的 QoS。

Nsight Systems 分析说明：时间切片与 GPU 指标

• 在 Nsight Systems 中，可以通过设置 gpucxtsw=true 选项来观察时间切片（time-slicing）行为（在没有 MPS 且使用默认计算模式的情况下）。
• 下图展示了一个进程在一段时间内的运行情况。其中，GPU 上下文（GPU Contexts）显示了两个不同的上下文（[22934] 和 [22935]）在时间上交替运行。同时，SMs Active（活跃 SMs）指标也显示了 GPU 利用率的变化。

• 分析命令示例：

  • 通过 nsys profile --gpucxtsw=true --gpu-metrics-devices=0 ./launch_script 收集数据。
  • 示例中使用的默认时间切片：sudo nvidia-smi compute-policy --set-timeslice 0 (可能的值为 {0, 3})。

• SMs Active 指标解读：

  • SMs Active（通过 Nsight Systems 的 GPU 指标收集）可以提供关于 GPU 利用率的深刻见解。
  • 严格的指标定义：SMs Active 是指至少有一个 warp 在 SM 上运行的周期数与总周期数的比率（以百分比表示）。
  • 提示：通过 nvidia-smi 显示的 GPU 利用率（GPU Utilization）表示“在过去采样周期内，一个或多个内核在 GPU 上执行的时间百分比”。即使只有一个内核在单个 SM 上运行，这个值也可能是 100%。
  • 指标收集本身可能会带来一些开销，您可能需要调整采样频率。

MPS 总结

• 分区类型（Partition Type）: 动态（Dynamic）

• 启用/配置时机（When to enable/configure）:

  • 在启动应用程序之前启动 MPS 服务。
  • 可以在应用程序启动期间进一步配置。

• 配置选项（Config. Options）:

  • 按需设置活动线程百分比（Active thread percentage）及其他环境变量（例如，CUDA_MPS_PINNED_DEVICE_MEM_LIMIT 用于限制可分配的固定设备内存量）。

• 是否需要应用程序更改（Application changes needed）:

  • 不需要（除非您使用驱动 API 进行上下文创建）。

• 使用场景（Use cases）:

  • 不同的进程。
  • 不需要错误隔离的场景。

参考文献/进一步阅读

• Multi-Process Service
• GTC 2022 讲座：https://www.nvidia.com/gtc/session-catalog/?search=S41956&tab.scheduledsessions=1651212000000#/session/16431411202860012f2xhttps://www.nvidia.com/gtc/session-catalog/?search=S41956&tab.scheduledsessions=1651212000000#/session/16431411202860012f2x
• 另请参阅：https://developer.nvidia.com/gtc/2023/session/s72383

---

Green Contexts

Green Contexts (GCs) 概述

• 绿色上下文（Green Context, GC） 与一组 GPU 资源相关联，目前这些资源是 SMs（流式多处理器）。
• Green Contexts 允许我们对 SMs 进行空间分区，因此提交给某个 GC 的 GPU 工作（如内核、图）只能使用这些指定的 SMs。
• 可以在同一个应用程序内拥有多个 Green Contexts。

• Green Contexts 功能通过 CUDA Driver API (-lcuda) 提供。

  • 需要对应用程序进行少量更改（例如，创建绿色上下文），但不需要更改 GPU 代码（即无需修改内核）。

• 本节中的 GC 示例将假定使用 CUDA 12.8。

Green Contexts vs. MPS 对比

• MPS 主要针对不同的进程，而 Green Contexts 针对单个进程。

  • 使用 GCs 不需要启动 MPS 服务。

• 假设 MPS 设置了 80% 的活动线程百分比，而 Green Context 设置了 80 个 SMs 作为可用资源（GPU 共有 100 个 SMs）。

  • 两者最多都只能使用 80 个 SMs。
  • MPS 进程可以使用任意 80 个 SMs；使用的 SMs 可能会随时间变化。
  • GC 可以使用特定的 80 个 SMs，这些 SMs 在 GC 创建时被空间分区。

• 过度订阅（Oversubscribed）示例：3 个进程或 GCs（GPU 仍有 100 个 SMs）。

  • 3 个 MPS 进程，活动线程百分比分别为 80、20 和 40。
  • 对比 3 个 Green Contexts，分别可访问 80、20 和 40 个 SMs。
  • 用户可以通过创建 GCs 的方式，控制这 3 个 GCs 之间共享多少 SMs。而这在 MPS 中是无法实现的。

Green Contexts: 设备资源与资源描述符

• GCs 对 GPU SM 资源进行空间分区，允许通过内核/图等方式提交的 GC 工作只在这些目标 SMs 上执行。
• CUdevResource：
  • 一个用于表示设备资源的结构体。

struct {
    CUdevResourceType type; // enum with CU_DEV_RESOURCE_TYPE_INVALID=0, CU_DEV_RESOURCE_TYPE_SM=1
    union {
        CUdevSmResource sm; // struct with unsigned int smCount
    };
};

• CUdevResourceDesc：
  • 一个封装了资源的描述符。

Green Context 创建示例：概述

• 创建 Green Context 的步骤:

  1. 获取可用的 GPU 资源。
  2. 将资源（SMs）分割成一个或多个同构分区和一个剩余分区。
  3. 如果需要，创建一个资源描述符来组合不同的分区。
  4. 从描述符创建一个 Green Context。

• Green Context 创建之后:

  • 您可以创建属于该 Green Context 的 CUDA 流，或者将该 Green Context 设置为当前上下文。
  • 任何后续在该流上启动的工作将只能访问该上下文的 SM 资源。
    • 包括通过 <<<...>>> 语法启动的内核或使用任何 CUDA 驱动/运行时 API 的操作。

Green Context 创建：步骤 1 - 获取可用的 SM 资源

• 获取我们可以分区的 GPU SM 资源，并填充 CUdevResource 结构体。

  • 从设备获取: CUresult cuDeviceGetDevResource(CUdevice device, CUdevResource* resource, CUdevResourceType type)
  • 从上下文获取: CUresult cuCtxGetDevResource(CUcontext hCtx, CUdevResource* resource, CUdevResourceType type)
  • 从 Green Context 获取: CUresult cuGreenCtxGetDevResource(CUgreenCtx hCtx, CUdevResource* resource, CUdevResourceType type)

• 通常，您的起点将是 GPU 设备。

Green Context 创建：步骤 2 - 分割 SM 资源

• 使用 cuDevSmResourceSplitByCount() API 静态地将可用的 CUdevResource SM 资源分割成一个或多个同构分区，可能还会留下一些 SMs 在剩余分区中。

• API 调用:
  CUresult cuDevSmResourceSplitByCount(CUdevResource* result, unsigned int* nbGroups, const CUdevResource* input, CUdevResource* remaining, unsigned int useFlags, unsigned int minCount)

  • 请求创建 *nbGroups 个同构组，每组包含 minCount 个 SMs。
  • 结果: 更新后的 *nbGroups（可能小于请求值），每个组包含 N 个 SMs（N >= minCount）。

• API 详细说明：cuDevSmResourceSplitByCount(...)

  • *nbGroups (>= 1):

    • 作为请求的同构组数量开始。
    • 可能会在调用过程中被更新为较小的数值。

  • minCount (>= 0):

    • 请求每个分区的最小 SMs 数量。实际值（图中的 N）可能会因粒度（granularity）和最小值的要求而更大。
    • 例如，Hopper 架构要求最小 8 个 SMs 且为 8 的倍数（可通过 useFlags 更改）。

• 示例代码:
  • 示例：请求将资源分割成 5 组，每组 8 个 SMs。

• 可以使用 result=nullptr 来查询可以创建的组数。
• 如果您不关心剩余的 SMs，可以使用 remaining=nullptr。

Green Context 创建：步骤 3 - 生成描述符

• 在分割完资源后，您需要为您计划使用的每组资源创建一个资源描述符。

• API 调用：CUresult cuDevResourceGenerateDesc(CUdevResourceDesc* phDesc, CUdevResource* resources, unsigned int nbResources)

• 示例：生成一个封装了 3 组资源的资源描述符。

• 此处展示的 CUDA API 版本为 12.8（早期版本仅支持 num_resources=1）。

Green Context 创建：步骤 4 - 创建上下文

• 从资源描述符创建一个 Green Context。

• 要在该 Green Context 上提交工作，您可以：
  • 为其显式创建一个流。
  • 或者，将该 Green Context 转换为一个 CUcontext，然后将其设置为当前上下文并提交工作。

Green Contexts: 启动工作

• 要在 Green Context 上提交工作，您可以为其创建一个流。

• 或者:
  • 将 Green Context 转换为一个具有其 SM 资源的主上下文。
  • 设置该上下文为当前上下文。
  • 在该上下文下创建流。

更多 Green Contexts 驱动 API

• CUresult cuGreenCtxRecordEvent(CUgreenCtx hCtx, CUevent hEvent)

  • 记录一个事件，捕获指定 Green Context 在此调用时的所有工作/活动。
  • FAQ: 这与 cudaEventRecord 有何不同？
    • 如果您只有一个流，并且在该流上记录了一个事件，那么它们是等效的。
    • 但如果您在 Green Context 中有多个流呢？
    • 如果没有这个新的 API，您需要在每个 Green Context 的流上分别记录一个事件，然后让依赖的工作分别等待所有这些事件。

• CUresult cuGreenCtxWaitEvent(CUgreenCtx hCtx, CUevent hEvent)

  • 使 Green Context 等待一个事件。
  • 更方便，因为您只需要让 Green Context 中的所有流等待一个事件完成。替代方案是为 Green Context 中的每个流调用一次 cudaStreamWaitEvent。

• CUresult cuStreamGetGreenCtx(CUstream hstream, CUgreenCtx* phCtx)

  • 将 phCtx 更新为与 hstream 关联的 Green Context（如果有），否则设置为 NULL。

• CUresult cuGreenCtxDestroy(CUgreenCtx hCtx)

  • 销毁 Green Context。

---

Green Contexts 示例

静态资源分区使关键工作能够更早地开始和完成。

示例时间线：

• 在流 strm1 上启动一个 长时间运行的内核 (delay_kernel_us)，该内核在整个 GPU 上占用多个波次。
• 在 CPU 上等待一段时间，然后在流 strm2 上启动一个 较短的关键内核 (critical_kernel)。
• 测量 GPU 持续时间以及从 CPU 启动到两个内核完成的时间。

长时间运行的内核代理是一个延迟内核，其中每个 CTA（Cooperative Thread Array）运行 delay_us，并且 CTA 的总数大于 SM（Streaming Multiprocessor）的总数。

示例运行场景对比

代码将在此处提供

Nsight Systems 时间线（无 Green Contexts）

下图展示了在没有使用 Green Contexts 的情况下，critical_kernel（高优先级）的启动被 delay_kernel_us 阻塞，导致了约 0.9ms 的“损失时间”。delay_kernel 运行约 10ms，而 critical_kernel 运行约 50us。

Nsight Systems 时间线（有 Green Contexts）

通过 Green Contexts 对 GPU 资源进行分区：
- 为 critical_kernel 分配 N 个 SM，为长时间运行的内核分配 7*N 个 SM（以及一些剩余的 SM），其中 N 是在给定 Green Context 约束下支持的最大值。
- 示例展示了在 H100 上（总共 132 个 SM），N=16 的情况。

如下图所示，critical_kernel 几乎在启动后立即执行，几乎没有损失时间。其执行时间约为 95us。而 delay_kernel 的执行时间增加到约 12ms（比之前增加了 2ms），因为它使用了更少的 SM。

在 Nsight Compute 中显示的 Green Context 资源

• 前面的示例有两个 Green Contexts，分别包含 16 个 SM 和 7*16 个 SM。
• Nsight Compute 的 Session 页面视图显示，在这个示例中，Green Contexts 按预期使用了不同的 SM。

Green Contexts 总结

特性	描述
分区类型	静态 (SMs)
何时启用/配置	在应用程序内部，启动工作之前。
配置选项	SM 数量和 SM 重叠由分区方式决定。
是否需要应用程序更改	是的，但仅在内核/GPU 代码之外。
使用场景	具有不同工作负载类型的单个进程；需要确保关键工作有可用的 SM 资源。

参考文献:
- https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__GREEN__CONTEXTS.html

---

Cluster Launch Control (集群启动控制)

线程块 (Thread blocks)

• 线程块 (CTA) - 内核执行中问题“工作分解”的基本单位。
• CTA 彼此完全独立，不进行交互（合作式启动除外）。
• 内核启动参数: kernel<<<CTA_COUNT, THREAD_COUNT>>>(ARGUMENTS)
• CTA 以任意顺序被调度执行。
• 驱动/调度器可以交错或重叠来自不同内核的 CTA，特别是基于流优先级。
• 调度一个 CTA 会产生开销。
• 一旦 CTA 开始执行，它会不间断地运行直到完成（进程之间的时间切片是可能的）。
• 例如，一个 CTA 不能为了一个更高优先级#内核的 CTA 而被换出。

  • 指与更高优先级流关联的内核。

如何确定线程块数量？

选择 CTA 数量的两种主要方法：

1. 基于问题大小 (Problem size - based):
   • 每个 CTA 处理固定/有限的工作量。
   • CTA 的数量与问题大小成比例。
   • 示例代码:

_global_ void kernel(float* data, float alpha, int n)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    data[i] *= alpha;
}
kernel<<<(n + 1023) / 1024, 1024>>>(data, alpha, n);

1. 基于硬件资源 (HW resources - based):
   • 又称持久化内核 (persistent kernels) 或 grid-stride loop。
   • 固定/有限数量的 CTA。
   • 每个 CTA 的工作量与问题大小成比例。
   • 示例代码:

_global_ void kernel(float* data, float alpha, int n)
{
  for (int i = blockIdx.x * blockDim.x + threadIdx.x;
       i < n; i += gridDim.x * blockDim.x)
    data[i] *= alpha;
}
kernel<<<sm_count * 2, 1024>>>(data, alpha, n);

基于问题大小的线程块数量控制

• ✅ 抢占 (Preemption): 高优先级内核的执行可以“插入”到低优先级内核的执行中。

• ✅ 负载均衡 (Load balancing): 尽管单个运行时间存在差异，但 SM 的总运行时间相似，从而产生较低的尾部效应。

• ❌ 调度器开销 (Scheduler overhead)

• ❌ 调度器开销 (Scheduler overhead) 和 公共代码开销 (common code overhead)

持久化内核 (Persistent Kernels)

这是对基于硬件资源方法的分析。

• ❌ 抢占 (Preemption): 在已启动的低优先级内核的 CTA 完成之前，无法执行高优先级内核。
• ❌ 负载均衡 (Load balancing): 单个运行时间的可变性导致 SM 运行时间的可变性。
• ✅ 调度器开销 (Scheduler overhead) 和 公共代码开销 (common code overhead) 较小。

自定义负载均衡 (Custom load balancing)

• 典型方法 - 基于原子操作的 CTA 计数器。

if (threadIdx.x == 0)
  shared_counter = bid.fetch_add(1, cuda::memory_order_relaxed);
__syncthreads();
bx = shared_counter;

reset_counter<<<1, 1>>>();
kernel<<<sm_count, 1024>>>(n_blocks);

• 缺点:
  • 解决方案繁琐，没有专门的硬件支持。
  • 需要在内核调用前重置计数器。
  • 需要以 sm_count 个 CTA 启动，并将实际的 CTA 数量作为参数传递。
  • 不支持抢占。
  • 如果内核可以并行启动，则需要多个计数器。
  • 需要共享内存同步。
  • 对于 2D/3D，需要使用除法来解码 CTA 索引：
    • bx = shared_counter % BX;
    • by = (shared_counter / BX) % BY;
    • bz = shared_counter / (BX * BY);

集群启动控制 - 优缺点总结

下表总结了三种方法的优缺点：

概念与优势

集群启动控制（Cluster Launch Control）是一种结合了多种调度策略优点的新方法。下表比较了基于问题规模、基于硬件资源、自定义（原子计数器）以及集群启动控制这四种方法的特性。集群启动控制在抢占（Preemption）、负载均衡（Load balancing）、开销（Overhead）和易用性（Ease of use）四个方面均表现出色，实现了“两全其美”的效果。

核心特性:

• 可用性: 从 Blackwell 架构（及 CUDA 12.8）开始提供支持。
• 工作窃取（Work-stealing）方法:
  • 内核（Kernel）请求一个已排队的协作线程数组（CTA）的索引。
  • 请求成功（SUCCESS）时，该索引会从可用 CTA 池中移除。
  • 请求失败（FAILURE）时，内核通常会退出。它也可以在一个工作窃取循环中执行一个被“窃取”（已取消）的 CTA 的工作。

下图通过时间线直观展示了集群启动控制的工作模式。CPU 发起两次启动（LAUNCH）操作，流式多处理器（SM）则持续处理任务，展示了动态和持续的工作分派流程，从而实现高效的负载均衡。

---

API 介绍：取消单个 CTA

集群启动控制提供了一套 API 来动态管理任务。取消一个 CTA 的基本流程如下：

1. 从单个线程异步请求取消，并将结果存入 __shared__ 内存。
2. 基于事务计数（transaction count），使用 __shared__ 内存屏障（barrier）同步该请求。
3. 检查同步结果以确认操作是否成功。
4. 从同步结果中提取被取消的 CTA 的索引。

注意：虽然可以从多个线程发起取消请求，但这在典型工作流中并不推荐，也非必需，因为取消操作本身是低延迟的。

API 代码示例

以下 PTX 代码展示了取消 CTA 的具体实现。代码逻辑分为前序（PROLOGUE）、线程块计算（THREAD BLOCK bx COMPUTATION）和后序（EPILOGUE）三个部分。

代码关键点解析：

1. 单线程发起（Single Arrival）：通常由线程块中的单个线程（例如 threadIdx.x == 0）发起异步取消请求，以避免冗余操作。

   

2. 基于事务计数的完成机制：同步的完成与否是通过事务计数来判断的。此处的 tx_count 基于结果数据结构 uint4 的大小，用于 mbarrier 的同步。

   

3. 异步请求：clusterlaunchcontrol_try_cancel 是一个异步（“in flight”）请求，它可以在前一个 CTA 仍在计算时被提交，从而实现计算和控制的重叠。

   

4. 取消完成：代码通过查询 clusterlaunchcontrol_query_cancel_is_canceled 的返回值来判断取消操作是否已成功完成，并据此决定是否跳出循环。

   

5. 内存栅栏（Fence）：__syncthreads() 不足以保证异步代理（async proxy）操作的可见性。必须使用专门的 fence_proxy_async_generic_sync_restrict 指令来确保所有线程都能观察到异步操作的结果。

   

API 示例优化

1. 双缓冲（Double-buffering）：使用双缓冲（result[2] 和 phase 变量）可以避免使用 __syncthreads() 来保护结果的覆写，从而提升性能。

   

2. 避免循环剥离（Loop Peeling）：使用 cg::invoke_one(cg::coalesced_threads(), ...) 代替 if (threadIdx.x == 0) 可以让所有线程执行统一的指令路径，避免了因条件分支导致的线程束发散（divergence），这是一种常见的性能优化技巧。

   

3. 多维 CTA 适配：在使用一维、二维或三维 CTA 时，需要相应地调整代码逻辑，例如发起请求的线程判断以及获取 CTA ID 的方式。

   

---

API 介绍：集群（Cluster）情景

在集群范围内取消 CTA 的流程与单个 CTA 类似，但引入了多播（multicast）机制：

1. 从集群中的任意一个 CTA 的单个线程异步请求取消。
2. 取消结果被多播到集群中每个 CTA 的本地 __shared__ 内存中。
3. 在每个 CTA 内部使用本地 __shared__ 内存屏障进行同步。
4. 从同步结果中提取根（root）CTA 的索引。
5. 将本地 CTA 的偏移量添加到根 CTA 索引上，得到全局索引。
6. （待办）在首次取消前，需通过屏障确保集群内所有 CTA 都已启动。

下图展示了一个 2x2 集群中，根 CTA (0,0) 将取消结果多播到其他 CTA 的示意图。

集群情景下的 API 代码示例

以下 PTX 代码展示了在集群范围内进行取消操作的实现。关键改动包括使用集群组同步 cg::cluster_group::sync()、通过 cg::cluster_group::thread_rank() == 0 选择发起线程，以及调用多播版本的取消指令 clusterlaunchcontrol_try_cancel_multicast。

代码关键点解析：
* 集群同步：在循环开始处的 cg::cluster_group::sync() 调用至关重要。它确保了在第一次迭代时集群内所有 CTA 都已准备就绪，并保护了共享数据在后续迭代中不被覆写。

以下代码片段展示了在集群情况下使用集群启动控制 API 的一个示例。

• 在序言（Prologue）部分之后，代码进入一个循环，通过 cg::cluster_group::sync() 进行同步。
• if (cg::cluster_group::thread_rank() == 0)：由单个集群线程发起请求，每个 CTA（Cooperative Thread Array，线程块）完成请求。
• bx = cg::this_cluster().block_index(J)：从多播的根节点获取当前 CTA 的索引。

集群启动控制：负载均衡示例

以下是在 NVIDIA B200 上使用 4GB 数组进行的计时测试，展示了集群启动控制在负载均衡方面的优势。

• 基准情况：当块（block）数量等于 SM（Streaming Multiprocessor）数量时：

  • 持久化内核（Persistent kernel）：0.030 秒
  • 集群启动控制：0.031 秒

• 有另一个内核并行运行（占用一个 SM）：

  • 持久化内核：0.059 秒 - 由于需要第二波（second wave）调度，时间翻倍。
  • 集群启动控制：0.031 秒 - 性能未受影响，展示了更好的资源利用率。

• 线程块在某个 SM 上运行时间长 25%（负载不均）：

  • 持久化内核：0.037 秒 - 执行时间因最慢的块而延长。
  • 集群启动控制：0.031 秒 - 能够有效处理负载不均衡，总时间未受影响。

集群启动控制：优先级示例

传统持久化内核的问题

下图展示了传统持久化内核在处理优先级任务时的局限性。

• 启动一个持久化内核 kernel_p。
• 随后启动一个高优先级内核 kernel_h。
• 结果：持久化内核 kernel_p 会占据整个 GPU 资源，导致高优先级内核 kernel_h 无法提前开始执行，必须等待 kernel_p 结束。

使用启动控制内核的优势

下图展示了使用集群启动控制的内核如何解决优先级问题。

• 启动一个启动控制内核 kernel_w。
• 随后启动一个高优先级内核 kernel_h。
• 结果：高优先级内核 kernel_h 可以立即开始执行。这是因为启动控制内核 kernel_w 允许“让出”（yield）计算资源给更高优先级的任务。

CUDA 开发者会议

该幻灯片列出了一系列与 CUDA 相关的开发者会议主题，涵盖了通用 CUDA、CUDA Python、CUDA C++、开发者工具、多 GPU 编程和性能优化等领域。

---