CUDA: New Features and Beyond

Stephen Jones, GTC 2022

引言：从单线程到数据中心
1.1. 软件开发的第一时代
1.2. 现代计算的规模：从节点到数据中心
1.3. 核心挑战：管理局部性
CUDA 编程模型与架构演进
2.1. GPU 架构的演进：十年九倍的流式多处理器
2.2. CUDA 编程模型：Grid → Blocks → Threads
2.3. 新层次：线程块集群 (Thread Block Cluster)
Hopper 架构新特性
3.1. Tensor Memory Accelerator (TMA)：异步数据移动
3.2. Cooperative Groups：面向硬件层级的编程
CUDA C++ 编译器与库
4.1. CUDA C++ 和编译器更新
4.2. JIT链接与链接时优化 (JIT Linking with LTO)
4.3. 数学库设备端扩展 (Maths Libraries Device Extensions)
4.4. CUTLASS：高性能矩阵运算模板库
CUDA 性能分析与调试工具
5.1. CUDA 开发者工具家族
5.2. Nsight Systems: 新的网络性能分析
5.3. Nsight Compute: 内核性能深度分析
5.4. NVTX v3: 代码注解与可视化
CUDA 应用优化技术
6.1. CUDA 二进制文件剖析与加载优化
6.2. CUDA Graphs 中的执行管理

1. 引言：从单线程到数据中心

1.1. 软件开发的第一时代

软件开发的第一个时代是单线程时代（持续至约2007年），其特点是顺序执行代码（Straight-Line Code）。

在这个时代，程序指令按顺序一步一步执行。即便是在现代，许多复杂算法的逻辑流程也可以被视为顺序执行代码。例如，一个典型的深度学习模型（如下图所示的Transformer架构）就包含了一系列顺序执行的层，如输入嵌入、多头注意力、前馈网络等。

1.2. 现代计算的规模：从节点到数据中心

高密度计算节点

下图展示了一个高密度计算节点的内部结构，这通常是现代数据中心和AI超级计算机的基本构建单元。图中可以清晰地看到多个GPU模块、高速互联组件以及先进的散热系统（如铜制散热片和热管）。这种设计旨在将强大的计算能力集成在单个紧凑的服务器单元中。

数据中心规模计算

计算需求已从单个节点扩展到整个数据中心。通过将前述的高密度计算节点集成到机架中，并部署大量机架，可以构建出用于处理大规模AI训练和高性能计算（HPC）工作负载的数据中心级计算集群。

1.3. 核心挑战：管理局部性

在计算架构中，管理数据局部性（Locality）是一个由来已久的核心问题。下图示意了这一基本概念：处理器（方块）访问与其物理位置更近的内存（矩形）时，速度更快、延迟更低。图中高亮的蓝色部分展示了一组处理器及其紧密耦合的本地内存。在扩展到数据中心规模时，如何高效地管理跨节点、跨机架的数据局部性，对于系统整体性能至关重要。

2. CUDA 编程模型与架构演进

2.1. GPU 架构的演进：十年九倍的流式多处理器（SM）

在过去的十年里，GPU 的并行处理能力实现了巨大的飞跃。以 NVIDIA Hopper H100 GPU 为例，其完整芯片包含 132 个流式多处理器（SMs），而十年前的 Kepler GK110 完整芯片仅有 15 个 SMs。这意味着，整个 GK110 芯片的计算核心数量，如今仅相当于 H100 芯片的一小部分。这种规模的增长要求编程模型也随之演进，以便开发者能够有效地利用如此大规模的并行硬件。

Page 31: 展示了 Hopper H100 芯片（132个SMs）与 Kepler GK110 芯片（15个SMs）在SM数量上的巨大差异，突显了GPU并行处理能力的指数级增长。

2.2. CUDA 编程模型：Grid → Blocks → Threads

CUDA 编程模型采用层次化的结构来组织和管理并行任务，其核心概念是 Grid（网格）、Block（线程块） 和 Thread（线程）。

工作网格 (Grid of Work)

一个完整的计算任务，在 CUDA 中被称为一个 Grid。可以将其想象成需要处理的整个工作负载，例如处理一张完整的图像。

划分工作为线程块 (Blocks)

为了实现并行处理，整个 Grid 被划分为一个由多个大小相等的 Blocks 组成的网格。每个 Block 负责处理整个任务的一部分。

Page 33: 将工作网格划分为许多个大小相等的线程块 (Blocks)，每个块处理一部分任务。

线程块内的线程 (Threads)

每个 Block 内部由许多个 Threads 组成。一个 Block 可以看作一个独立的程序，其内部的多个线程协同执行。这种模型允许将大规模问题分解为可以在 GPU 的众多核心上并行执行的小任务。

2.3. 新层次：线程块集群 (Thread Block Cluster)

为了更好地利用现代 GPU 架构的物理局部性，Hopper 架构引入了一个新的层次：线程块集群 (Thread Block Cluster)。

概念与定义

一个线程块集群是一组被共同调度到相邻多处理器上的线程块的集合。这在原有的 Grid → Block → Thread 层次结构中增加了一个新的层级，变为 Grid → Cluster → Block → Thread。

Page 35: 线程块集群是介于网格 (Grid) 和线程块 (Block) 之间的新层次，它是一组线程块的集合。

利用 GPU 规模的局部性

线程块集群的设计初衷是为了充分利用 GPU 硬件的物理局部性。通过将一个集群内的所有线程块调度到物理上相邻的 SMs 上，可以实现以下优势：
- 保证协同定位的线程块 (Guaranteed co-located blocks)：集群内的块在物理上彼此靠近。
- 新的保证并发层级 (New tier of guaranteed concurrency)：集群内的所有块保证同时运行。
- 快速数据交换与同步 (Fast data exchange & sync)：由于物理上的邻近性，集群内的线程块可以更高效地进行数据共享和同步。

Page 36: 将线程块集群的概念映射到GPU硬件上，它占据了芯片上一组相邻的SMs，从而利用硬件局部性。

编程实现

在程序中构建和使用线程块集群需要遵循特定的编程规范。

定义：一个集群最多可以包含 16 个线程块。
保证：
集群内的线程块保证被分配到不同的 SMs 上。
集群内的所有线程块保证同时运行。
集群可以是 1D、2D 或 3D 结构，与线程块类似。
实现方式：
通过一个新的注解 __cluster_dims__(x, [y, [z]]) 来为一个 CUDA 核函数 (kernel) 指定其所需的集群维度。
以下是一个代码示例，定义了一个 4x2x1（即 8个）线程块组成的集群：

```c++
// 定义一个 4x2x1 的8块集群
cluster_dims(4, 2, 1)
global void hellocluster()
{
// 获取当前集群的 cooperative group
cooperative_groups::cluster_group cluster = this_cluster();
// 在集群内进行同步
cluster.sync();

  printf("Hello from cluster elem %d\n", cluster.cluster_rank());

}
```
- 启动配置：此外，还引入了新的可扩展启动 API，允许在运行时配置集群参数。

Page 37: 线程块集群的编程细节，包括其定义、保证、以及通过`__cluster_dims__`注解和cooperative_groups API在代码中进行实现。

3. Hopper 架构新特性

3.1. Tensor Memory Accelerator (TMA)：异步数据移动

Tensor Memory Accelerator (TMA) 是一种用于异步数据移动的硬件加速单元。它支持硬件加速的双向批量拷贝，可以在全局内存（Global Memory）与共享内存（Shared Memory）之间，以及在集群内分布式共享内存（DSMEM）之间进行数据传输。TMA 使用异步事务屏障（Asynchronous Transaction Barrier）来跟踪操作的完成情况。

低延迟单边数据传输与光环交换（Halo Exchange）

TMA 的一个关键应用是实现低延迟的单边数据传输。它支持在集群内部进行单边批量数据拷贝，并在拷贝到全局内存时支持元素级的规约（reduction）操作。

一个典型的例子是在一个集群内进行单边光环交换（halo exchange）。如下图所示，在一个2x2的线程块集群（Cluster of Blocks）中，TMA可以通过自同步事务（self-synchronizing transactions）在集群内的不同线程块之间高效交换边界数据（halo区域），其速度比传统方法快7倍。

更快、更灵活的异步拷贝

TMA 带来了更快、更灵活的异步拷贝功能，它作为对现有 cuda::memcpy_async() 函数的直接增强。其核心优势包括：
- 单线程触发：单个线程即可触发任意大小的数据拷贝，无需编写循环或使用协作式拷贝（collective copying）。
- 聚合屏障：多个拷贝操作可以汇集到一个集群范围的异步事务屏障上。

如下图所示，TMA 单元在 H100 SM 内部，负责处理从 GPU 内存到共享内存的数据传输。通过一个新的异步事务屏障机制（Arrive -> Wait），可以有效跟踪数据传输的完成状态。编程模型上，开发者可以使用带有屏障对象的新版 cuda::memcpy_async 语法来显式启用 TMA 的高级功能，而旧的语法在 H100 上也会自动利用 TMA 进行增强。

硬件加速的1D-5D张量内存拷贝

TMA 能够硬件加速多维张量的拷贝，最高支持5阶张量（rank 5）。其特性包括：
- 自动地址生成：自动处理多维张量的步长（stride）和地址计算。
- 边界填充：为越界访问提供边界填充（Boundary padding）。
- “发射后不管”：仅需单个线程发起拷贝指令，后续所有复杂操作均由 TMA 硬件处理，开发者无需编写迭代或边界检查代码。

下图展示了 TMA 如何从一个大的多维张量中拷贝一个子区域（sub-region）。TMA 会自动处理张量访问步长、块宽度/高度以及必要的自动填充。

3.2. Cooperative Groups：面向硬件层级的编程

Cooperative Groups 是一个用于显式表达线程间协作和同步的编程模型。它将GPU的执行层级结构暴露给开发者，允许在不同粒度上进行编程。

GPU 自然执行层级

GPU 的自然执行层级从粗到细依次为：
- Grid of work: 整个内核的工作网格。
- Cluster of Blocks: 线程块的集群。
- Block of Threads: 线程块。
- Warp of Threads: 线程束。

开发者可以通过相应的API获取每个层级的句柄，例如 this_grid()、this_cluster()、this_thread_block() 和 tiled_partition<32>()。

各层级的协作操作 (Collective Operations)

Cooperative Groups 在不同的执行层级上支持不同的协作操作：
- Grid 和 Cluster 层级: 主要支持同步（Synchronization）。需要注意的是，要使用 grid 范围的协作操作，Grid 必须以协作方式启动（cooperatively launched）。
- Block 和 Warp 层级: 支持更丰富的操作，包括同步（Synchronization）、规约（Reduction）和前缀和（Prefix-sum）。
- Warp 及更小粒度的组: 除了上述操作外，还支持 shfl, any, all, ballot, match 等 warp 内置函数。

支持多线程束组 (Multi-Warp Groups) 的协作操作

该模型还支持由多个 warp 组成的组进行协作操作。支持的 thread_block_tile 尺寸包括64、128、256和512。例如，可以通过 cooperative_groups::thread_block_tile<128> 来创建一个由128个线程（即4个warp，一个 quadwarp）组成的组，并在此组上执行协作操作。

类型安全实现可组合的并行函数库

Cooperative Groups 通过C++的类型系统，可以创建类型安全的并行函数库。函数可以明确声明其期望的线程数量。例如，一个为128个线程设计的基数-128 FFT（Radix-128 FFT）函数，可以将其接口参数类型定义为 const thread_block_tile<128>&。这样，编译器就能确保该函数只能被一个包含128个线程的 thread_block_tile 对象调用，从而实现了编译时的安全性，避免了运行时错误。

// 内核函数将线程块划分为128线程的tile
__global__ void kernel(...) {
    thread_block_tile<128> tile128 = tiled_partition<128>(this_thread_block());
    fft128(tile128, ...); // 调用FFT函数
}

// 一个必须由128个线程调用的设备函数
__device__ void fft128(const thread_block_tile<128>& group, ...) {
    do_fft(...);
    group.sync(); // 同步组内所有128个线程
}

应用示例

示例：分层执行、数据交换与同步

Cooperative Groups 可用于实现复杂的分层计算模式。以下是一个流体模拟的例子：
1. 启动一个由 4x4 个 cluster 组成的协作式 Grid，每个 cluster 包含 4x2 个 block。
2. 使用协作式 cluster 索引，从邻近的 cluster 拉取数据。
3. 在每个 block 内部进行局部求解。
4. 在 cluster 层面进行同步，并利用分布式共享内存（distributed shared memory）进行块间通信和求解。
5. 在 grid 层面进行同步和求解，然后返回第2步，开始下一次迭代。

示例：任意规模的生产者/消费者任务并行

生产者/消费者（Producer/Consumer）模型可以在 Grid、Cluster 或 Block 等任意层级上实现。
- Grid 级别: 将整个 GPU 上的 Clusters 划分为 "生产者" 和 "消费者" 角色。它们通过全局内存交换数据，并使用 this_grid().sync() 进行同步。
- Cluster 级别: 在一个 Cluster 内部，将 Blocks 划分为 "生产者" 和 "消费者" 角色。它们通过分布式共享内存或全局内存交换数据，并使用 this_cluster().sync() 进行同步。
- Block 级别: 在一个 Block 内部，将 Warps 划分为 "生产者" 和 "消费者" 角色。它们通过共享内存或全局内存交换数据，并使用 this_thread_block().sync() 进行同步。

该模式在 Transformer 模型的计算中也有应用，例如可以将模型中的一部分层指定为生产者，另一部分为消费者。

示例：Longstaff Schwartz 定价模型

这是一个将生产者/消费者任务并行应用于量化金融领域的实例。
- Longstaff Schwartz 模型: 一种通过蒙特卡洛模拟和反向迭代来确定金融期权价值的定价技术。
- 分析: 分布式共享内存（Distributed shared memory）的更大容量为高效内核的编程带来了新方法。线程块级别的生产者/消费者模型现在可以在共享内存中实现，并利用快速的 cluster 同步。硬件加速的 cluster 同步机制易于使用且高效。
- 性能: 实验表明，在 H100 平台上，与不使用 cluster 的方案相比，使用 cluster 同步（Cluster.sync()）可以将 Longstaff Schwartz 模型的吞吐量提升 2.7倍。如果进一步使用 cluster 和异步屏障（Barrier），性能还能再提升 10%。

Page 57: 集群对生产者/消费者模型在Longstaff Schwartz期权定价中的影响

4. CUDA C++ 编译器与库

4.1. CUDA C++ 和编译器更新

CUDA C++ 对128位整数的支持

CUDA 11.7 在整个工具链和库中增加了对 __int128 类型的支持。
- 兼容性: 可与兼容的主机编译器（如 gcc, clang, icc）一起使用，或通过运行时编译（NVRTC）使用。
- 完整支持: 支持算术、逻辑、位运算、数学运算、库函数和开发工具等。
开发者可以在主机代码和设备代码（kernel）中无缝使用128位整数。

NVRTC 多线程编译

为了缩短应用的构建时间，NVRTC 和 PTX JIT（即时编译）的锁机制从全局锁（global locking）改进为分阶段锁（per-stage locking）。这一改进使得在使用多个CPU线程进行编译时，能够实现并发编译。
- 效果: 从顺序编译流程转变为流水线式（Pipelined）编译流程，显著提升了编译速度。
- 性能: 测试表明，在 CUDA 11.5 中，使用4个CPU线程进行编译的速度比使用单个线程快一倍以上（从44秒缩短至20秒）。

近期其他编译器更新

C++20 支持: CUDA 11.7 提供了对 C++20 的预览支持，兼容 GCC 10+, Clang 10+, nvc++ 20.7+ 等主机编译器。对 MSVC 的支持将在后续版本中提供。
Grid私有的“硬”常量: 新增 __grid_constant__ 注解，用于内核参数。该注解允许编译器避免为只读参数创建每个线程的私有副本，从而优化资源使用。
诊断信息的程序化管理: 提供了一系列 #pragma 指令（如 nv_diag_suppress, nv_diag_warning 等）来在代码中动态地控制编译器的诊断信息（错误、警告）。
新的 nvcc 目标架构选项:
- -arch=all: 为所有支持的架构生成代码。
- -arch=all-major: 为所有支持的主要架构版本生成代码。
- -arch=native: 为系统上所有可见的GPU生成代码。

4.2. JIT链接与链接时优化 (JIT Linking with Link-Time Optimizations)

即时编译链接（JIT Linking）与链接时优化（LTO）相结合，能够在运行时将对象文件与库进行组合和优化。该技术允许在a.out（可执行文件）与cuFFTDx内核库之间进行调用，最终生成一个完全优化的内核。

性能对比：
在A100（40GB）GPU上，JIT-LTO回调相较于间接回调（Indirect Callback）在不同FFT规模下展现了显著的性能提升，最高可达1.92倍。

工作原理：对象文件在运行时被合并和优化。nvcc -dlto 编译 .cu 文件生成可重定位的对象文件，运行时JIT LTO将此对象文件与cuFFTDx内核库链接并进行联合优化，生成一个高效的、完全优化的内核。

图表展示了JIT-LTO回调（绿色）相比间接回调（深蓝色）在不同FFT尺寸下的性能提升（GFlops）

4.3. 数学库设备端扩展 (Maths Libraries Device Extensions)

设备端代码直接调用数学库（如cuFFTDx）可以显著加速小规模问题的处理。

性能对比：设备端调用 vs. 主机端启动
在A100（80GB）GPU上，与传统的主机端启动cuFFT相比，从设备端代码调用的cuFFTDx在处理小尺寸FFT时性能优势尤为明显，TFlop/s性能提升显著。

图表对比了设备端调用的cuFFTDx（绿色）和主机端启动的cuFFT（蓝色）在不同FFT尺寸下的性能，设备端调用在小尺寸问题上优势巨大

支持设备端调用的接口库：
通过在高性能数学库中调用自定义内核代码，可以实现更复杂的计算流程。

已支持:
cuFFTDx: 快速傅里叶变换。
未来支持:
cuBLASDx: GPU加速的基础线性代数子程序库 (BLAS-1, 2, 3)。
cuSOLVERDx: GPU加速的稠密线性代数库。

应用案例：带内联卷积的FFT
在带内联卷积的FFT计算中，cuFFTDx的性能优于标准的cuFFT以及使用回调机制的cuFFT。

图表展示了在带内联卷积的FFT任务中，cuFFTDx（亮绿色）相较于cuFFT（蓝色）和带回调的cuFFT（深绿色）的性能优势

4.4. CUTLASS：高性能矩阵运算模板库

CUTLASS 是一个用于深度学习和高性能计算的 CUDA C++ 模板库。它提供了在各种范围和规模下最优化的 CUDA C++ 矩阵运算模板。

核心特性:
Tensor Core 加速: 支持矩阵乘法（GEMM）、卷积、归约以及融合的输入/输出操作。
开源: 是一个仅包含头文件的 C++ 模板库。
底层接口: CUTLASS 编译到 CUDA 内核中，直接与 GPU 硬件交互。

Tensor Core 加速操作的粒度

CUTLASS 在不同抽象层次上提供了 Tensor Core 加速的操作。
- Device: GEMM、卷积、归约等，支持所有数据类型、SIMT、Tensor Core及所有架构。
- Kernel: GEMM、批处理GEMM、卷积、归约、融合操作等。
- Thread Block: 流水线矩阵乘法、Epilogue、到张量的集合操作、卷积矩阵访问。
- Warp: 张量核乘加操作、高效访问置换内存布局。
- Thread: 数值转换、<functional> 操作符、复数<T>、快速数学算法。
- Architecture Intrinsic: 包装了架构特定的PTX指令（如 mma, cp.async, ldmatrix, cvt）。

使用Tensor Cores加速单精度计算

CUTLASS 采用一个三步TF32序列，在A100 GPU上实现了48 TFLOPs的FP32精度计算性能。

优势:
该方法通过补偿舍入误差，能够有效利用TF32 Tensor Cores。
性能超过峰值单精度浮点乘加（FFMA）性能的两倍。
精度优于单精度（但与IEEE标准不完全兼容）。
为GEMM和卷积提供了示例实现。

性能与精度分析:
- 性能: 3步TF32 GEMM的性能远超A100的FP32峰值性能（19.5 TFLOPs）。
- 精度: 相较于标准的FP32计算，3xTF32方法的相对误差略高，并随问题规模（GEMM-K维度）增大而增加。

左图展示了3步TF32 GEMM的FP32性能，远超硬件FP32峰值；右图对比了3xTF32与FP32的相对误差

5. CUDA 性能分析与调试工具

5.1. CUDA 开发者工具家族

CUDA生态系统提供了一系列丰富的开发工具，涵盖调试、性能分析、代码修正和IDE集成。

IDE 集成 (IDE Integration):
Eclipse
Visual Studio Code
Visual Studio
调试器 (Debugger):
Nsight Visual Studio Edition
Nsight Visual Studio Code Edition
Nsight Eclipse Edition
cuda-gdb
性能分析器 (Profiler):
Nsight Systems
Nsight Compute
NVTX
CUPTI
修正工具 (Sanitizer):
Memcheck (内存错误检查)
Initcheck (未初始化检查)
Racecheck (竞态条件检查)
Synccheck (同步错误检查)

CUDA开发者工具生态图，展示了调试器、性能分析器、修正工具和IDE集成四大类工具

5.2. Nsight Systems: 新的网络性能分析

新的网络性能分析功能可以拦截并追踪对UCX（Unified Communication X）协议层的调用。这使得开发者能够深入了解MPI通信背后的底层UCP（Unified Communication Protocol）操作。

追踪 UCX 调用

关键特性: 追踪非阻塞UCP通信操作的完成情况，从而精确定位网络瓶颈。

Nsight Systems 时间线视图，展示了MPI调用与底层UCX发送/接收操作的对应关系

Nsight Systems中的网卡性能指标 (NIC Performance Metrics)

Nsight Systems现在可以直接展示网卡（NIC）的性能指标，帮助分析多节点应用的通信性能。

可监控指标:
InfiniBand (IB) 接收字节数
InfiniBand (IB) 发送字节数
InfiniBand (IB) 发送等待时间
启用方式:
通过命令行参数 $ nsys profile --nic-metrics=[true|false] ... 来控制是否采集网卡指标。

Nsight Systems 界面截图，显示了多个NIC的IB接收/发送速率和等待时间等性能指标

5.3. Nsight Compute: 内核性能深度分析

Nsight Compute 提供了深入分析CUDA内核性能的强大工具。

占用率计算器 (Occupancy Calculator): 帮助分析和可视化寄存器数量和共享内存大小对SM（Streaming Multiprocessor）占用率的影响。
独立源码查看器 (Standalone Source Viewer): 将性能数据直接关联到源代码和汇编代码，方便定位瓶颈。
寄存器依赖可视化 (Register Dependency Visualization): 可视化汇编指令间的寄存器依赖关系，帮助理解指令级并行性和延迟。

Nsight Compute 工具界面展示，包括占用率计算器、源码查看器和寄存器依赖可视化功能

5.4. NVTX v3: 代码注解与可视化

NVTX (NVIDIA Tools Extension) 是一个用于在代码中插入标记以供工具进行分析的注解库。

简介:
功能: 插入标记，用于工具分析。
开销: 开销极小，可忽略不计。
依赖: 无依赖，仅需包含头文件。
工作原理:
标记 (Markers): 记录一个时间点。
范围 (Ranges): 使用 push/pop 对记录一个时间段。范围形成一个栈结构，可以展示嵌套的细节。
工具集成: 性能分析工具（如Nsight Systems）会追踪NVTX调用，并在时间线上显示这些标记和范围。

#include <nvtx3/nvToolsExt.h>

void LaunchKernel() {
    nvtxRangePush(__FUNCTION__);
    Kernel<<<20000, 256>>>();
    nvtxRangePop();
}

void Example() {
    nvtxRangePush("Creating CUDA context");
    cudaFree(0);
    nvtxRangePop();

    nvtxRangePush("Launching & waiting for kernels");
    for (int i = 0; i < 3; ++i) {
        LaunchKernel();
    }
    cudaDeviceSynchronize();
    nvtxRangePop();
}

NVTX在性能分析工具中的可视化效果

NVTX的注解代码会直接映射到性能分析工具的时间线视图中，形成层次化的事件范围，极大地增强了代码执行流程的可读性。

nvtxRangePush("Creating CUDA context"); 在时间线上生成一个名为 "Creating CUDA context" 的范围。
nvtxRangePush("Launching & waiting for kernels"); 创建了一个父范围。
在 LaunchKernel 中调用的 nvtxRangePush(__FUNCTION__); 则在父范围内部生成了多个嵌套的 "LaunchKernel" 子范围。

Nsight Systems 时间线视图展示了NVTX代码注解如何生成层次化的事件范围，清晰地标识出内核启动等不同阶段

6. CUDA 应用优化技术

6.1. CUDA 二进制文件剖析与加载优化

二进制文件基本结构

一个简单的CUDA "Hello world"程序经过nvcc编译后，会生成一个包含主机和设备代码的可执行文件。

// Hello world example
__global__ void kernel() {
    printf("Hello, CUDA\n");
}

void main() {
    kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
}

编译命令：nvcc example.cu -o example

与传统的CPU程序不同，编译生成的CUDA可执行文件包含多个部分：
- CPU代码区 (CPU code section): 包含主机端代码，如main()函数。
- GPU代码区 (GPU code section): 包含编译为特定GPU架构机器码的设备端代码，如kernel()。
- PTX JIT区 (PTX JIT section): 包含PTX（Parallel Thread Execution）中间代码。这使得程序可以在运行时即时编译（JIT），以兼容未来或编译时未指定的GPU架构。

CUDA可执行文件结构示意图，包含CPU代码、GPU代码和PTX JIT三个部分

标准加载流程与瓶颈

当链接大型库（如 cuDNN）时，该库会向最终的二进制文件中添加大量的内核 (kernels) 和 PTX 文件，导致二进制文件体积显著增大。
nvcc example.cu -o example -lcudnn

当一个CUDA应用程序启动时，其标准加载流程如下：
1. 从磁盘加载到主机内存：操作系统将整个二进制文件从磁盘加载到CPU主机内存中。由于所有代码段都被完整加载，这会导致较大的主机内存占用（例如，1,866MB）。

2. 从主机内存上传到设备内存：CUDA 驱动程序在初始化阶段，会将二进制文件中包含的所有 GPU 内核从主机内存上传到 GPU 设备内存中。这个过程可能非常耗时（例如，1.6秒），并且会占用大量的设备内存（例如，1,245MB），即使应用程序在整个生命周期中可能只使用到其中的一小部分内核。

延迟加载：按需上传函数 (Lazy Loading)

为了优化上述标准加载流程带来的高昂开销，引入了延迟加载（Lazy Loading）机制。其核心思想是：仅在函数（内核）首次被引用或访问时，才将其从主机内存上传到设备内存。

采用延迟加载可以带来显著的优化效果：
- 减少二进制加载时间：由于初始阶段无需上传所有内核，程序启动速度加快。
- 降低设备内存占用：GPU 内存中只保留当前活跃的内核，大幅减少了静态内存占用。
- 降低主机内存占用：同样可以观察到主机内存占用的减少。

激活延迟加载与性能对比

延迟加载功能可以通过设置一个环境变量来激活：
set CUDA_MODULE_LOADING=LAZY

在一个简单的 cuDNN 卷积测试中（环境：Ubuntu 20.04, A100 40GB），延迟加载与标准加载的性能对比如下：
- 二进制加载时间 (s)：标准加载为 1.6秒，延迟加载为 0.8秒，时间减少 50%。
- 设备内存占用 (MB)：标准加载为 1,245MB，延迟加载为 435MB，占用减少约 65%。
- 主机内存占用 (MB)：标准加载为 1,866MB，延迟加载为 1,229MB，占用减少约 34%。

实验数据明确显示，延迟加载能够有效优化应用程序的启动时间和内存资源利用率。

6.2. CUDA Graphs 中的执行管理

CUDA Graphs 提供了对复杂工作流进行高效管理和执行的能力。新的特性进一步增强了其灵活性和控制力。

图节点启用/禁用 (Graph Node Enable/Disable)

该功能允许用户在图启动（launch）之前动态修改执行流程，通过禁用图中的某些节点来实现。

示例：在一个原始图中，存在两条并行的执行路径 (B->D->E 和 X->Y)。通过禁用节点 "D"，当图执行时，节点 "D" 将被跳过。因此，左侧分支的实际执行路径将变为 "BCE"。这为条件执行提供了高效的控制机制，而无需重新构建整个图。

节点特定优先级 (Node-Specific Priority)

该功能允许为任务图中的特定分支或路径创建高优先级，从而影响调度和执行顺序。

示例：在一个图中，可以将路径 "BCDE" 设置为高优先级（图中以白色边表示）。这意味着在执行时，调度器会优先执行 "BCDE" 路径上的节点，使其相对于路径 "XY" 具有更高的执行优先级。这对于管理具有不同服务质量（QoS）要求的任务非常有用。