AUTOMATED PERFORMANCE IMPROVEMENT USING CUDA LINK TIME OPTIMIZATION

MIKE MURPHY, COMPILER DEVELOPMENT
作者信息未提供

为什么需要链接时优化 (LTO)

链接时优化 (LTO) 旨在解决单独编译带来的优化限制，从而提高程序的运行时性能。

单独编译的问题

最初，CUDA 程序规模小，为单文件，不支持单独编译。
CUDA 5.0 增加了单独编译支持，使得设备代码能够跨多个文件。
但这限制了优化的范围，导致更少的内联和更多的内存流量，从而影响运行时性能。
问题图示

链接时优化 (LTO)

LTO 将单独编译的灵活性与全程序优化的性能相结合。

全程序编译 (Whole program compilation)：单个 whole.cu 文件，通过 nvcc 编译优化，生成完全优化的可执行文件 a.Out。
单独编译 (Separate compilation)：a.cu 和 b.cu 分别通过 nvcc -dc 编译优化生成 a.o 和 b.o（在各自范围内优化），然后链接生成 a.out（部分优化可执行文件）。
链接时优化 (Link time optimization)：a.cu 和 b.cu 分别通过 nvcc -dc -dlto 编译生成 a.o 和 b.o（LTO 中间表示），然后 nvlink -dlto 将其合并为 nvlink 中的中间表示，进行优化，最终生成完全优化的可执行文件 a.out。
LTO工作流程

如何使用 LTO

只需在 nvcc 命令行中添加 -dlto，例如：
- nvcc -arch=sm_80 -dc *.cu -dlto
- nvcc -arch=sm_80 *.o -dlto
如果使用 -gencode 并希望在编译时保存多个中间文件，可以使用 code=lto_NN 作为目标，例如：
- nvcc -gencode arch=compute_80,code=sm_80 -gencode arch=compute_60,code=lto_60 -dc *.cu
LTO 在 CUDA 11.2 中引入；NVIDIA 是首家将 LTO 应用于 GPU 编译的公司。
LTO 中间文件可以混淆代码，有助于保护知识产权。

部分 LTO

可以对部分文件使用 LTO 构建，而对其他文件不使用。
- 使用 LTO 的文件获得完全优化，其他文件获得常规优化。
- 这在链接您未构建的库时很有用。
部分LTO流程
- a.cu 和 b.cu 通过 -dc -dlto 生成 LTO 中间表示 (a.o, b.o)。
- libc.a 包含非 LTO 设备代码。
- nvlink -dlto 对 a+b 执行 LTO，然后与 libc 链接，生成可执行文件。

LTO 结果

LTO 带来的性能提升因应用程序而异，范围从 5% 到 2x。
LTO性能提升图表
- quicksilver：LTO 性能提升达 91%。
- lawa：LTO 性能提升达 91%。
- defense：LTO 性能提升达 32%。
- LLNL：LTO 性能提升达 27%。
劳伦斯利弗莫尔国家实验室 (LLNL) 在一个大型蒙特卡洛风格应用程序上使用了 LTO，该程序包含数百个文件和大约 300,000 行代码。
他们的论文链接：https://www.osti.gov/biblio/1798430-enhancements-supporting-ic-usage-pem-libraries-next-gen-platforms
“LTO 在所有情况下都提供了加速；最大加速为 27.1%。”

即时链接时优化 (JIT LTO)

在 JIT LTO 中，链接在运行时执行。
NVIDIA 是首家将 LTO 应用于 JIT 编译的公司。
JIT LTO 是 CUDA 11.4 中引入的预览功能，API 名称可能会发生变化。
LTO IR 的生成可以通过 nvcc 离线进行，也可以通过 nvrtc 在运行时进行。

离线与运行时生成的 JIT LTO IR

离线 (Offline)：使用 nvcc -dc -arch=lto_70 *.cu -fatbin 命令生成 LTO IR。它包含为多种架构生成的代码。
运行时 (Runtime)：使用 nvrtc API，通过 nvrtcCompileProgram、nvrtcGetNVM、cuLinkCreate 等调用生成 LTO IR。
无论哪种方式，LTO IR 都会被 JIT LTO 处理，最终生成 SASS (GPU 汇编代码)。
JIT LTO IR生成流程

JIT LTO 调用

使用 nvrtc 生成 LTOIR (C++ 字符串)：

nvrtcProgram prog;
nvrtcCreateProgram(&prog, input, name, 0, nullptr, nullptr);
const char *options[2] = {"-dlto", "-dc"};
nvrtcResult result = nvrtcCompileProgram(prog, 2, options);
size_t irSize;
nvrtcGetNVMSize(prog, &irSize);
char *ltoIR = (char*)malloc(irSize);
nvrtcGetNVM(prog, ltoIR); // returns LTO IR

LTO IR 随后传递给 cuLink 驱动 API，链接在运行时执行:

CUlinkState linkState;
CUjit_option jitoptions[] = {CU_JIT_OPTION_JIT_INPUT_TYPE};
void *jitoptionValues[] = {(void*)1};
cuLinkCreate(1, jitoptions, jitoptionValues, &linkState);
cuLinkAddData(linkState, CU_JIT_INPUT_NVVM,
            ltoIR, irSize, name, 0, NULL, NULL);
cuLinkAddData( /* other input */ );
size_t size;
void *linkedCubin;
cuLinkComplete(linkState, &linkedCubin, &size);
cuModuleLoadData(&mod, linkedCubin);

数学库在其 CreatePlan API 中隐藏了 cuLink 的细节。
JIT LTO 调用代码

数学库中的 JIT LTO 回调

CUDA 数学库 (cuFFT, cuSPARSE 等) 正在开始使用 JIT LTO。
请参阅 GTC Fall 2021 演讲 “JIT LTO Adoption in cuSPARSE/cuFFT: Use Case Overview”。
旧方法 (Previous)：用户提供的设备函数通过间接回调调用，不涉及模块间内联或跨模块优化。
使用 JIT LTO (With JIT LTO)：用户设备函数和库本身的代码都生成 LTO-IR 片段，JIT LTO 对这些片段进行跨用户和库代码的优化，最终生成 SASS。
JIT LTO 回调流程

JIT LTO 结果: 数学库回调

间接回调 vs JIT LTO 回调 (A100 (40 GB))
在不同的 FFT 大小下，JIT LTO 回调显示出显著的性能提升 (以 GFLOPS 衡量)。
- FFT 大小 128：JIT LTO 性能提升约 1.88x。
- FFT 大小 256：JIT LTO 性能提升约 1.23x。
- FFT 大小 512：JIT LTO 性能提升约 1.27x。
- FFT 大小 1024：JIT LTO 性能提升约 1.76x。
- FFT 大小 2048：JIT LTO 性能提升约 1.80x。
- FFT 大小 4096：JIT LTO 性能提升约 1.92x。
- FFT 大小 8192：JIT LTO 性能提升约 1.78x。

JIT LTO 用例：数学库组合

数学库目前包含数千种手写内核的排列，但不能涵盖所有尺寸。通过将构建块与 JIT LTO 结合，我们可以生成用户所需的任何配置，同时获得最佳性能和最小库大小。

现状：

用户配置 A, B, C： 每个配置都涉及独立的加载（Load）、存储（Store）、计算（Compute）过程。
设备功能（Device Function）： 每个操作都对应一个设备功能。
SASS： 最终生成多个 SASS（GPU汇编语言）二进制文件。
问题： 导致二进制文件膨胀（Exploding binaries）。

使用 JIT LTO：
- 单一用户配置： 统一的用户配置。
- 加载、存储、计算： 操作转换为 LTO-IR（链接时间优化中间表示）。
- JIT LTO： 在运行时（Runtime）通过 JIT LTO 处理 LTO-IR。
- SASS： 最终生成单一的 SASS 二进制文件。
- 优势： 更好的性能（Better performance），更小的二进制文件占用空间（Smaller binary footprint）。

何时使用 (或不使用) LTO

LTO 提供了一种生成更优代码的简便方法。
但 LTO 将编译时完成的优化移至链接时进行。
示例： LAWA 应用程序约有 38K 行代码，分布在 300 多个文件中。
- 在单独编译中，链接时间从 1 秒增加到使用 LTO 的 49 秒。
- 但编译时间减少了 18 秒，整体编译 + 链接时间仅增加了 5%。
使用场景： 在应用程序需要反复运行时使用，而不是在迭代开发时使用。
- 根据早期的 LLNL 论文：“考虑到 LTO 提供的加速，链接时间减慢对于生产构建是可以接受的。”
使用 JIT LTO，较慢的链接时间可能通过参数化函数范围的缩小来抵消。
我们正在开发编译器增强功能，以在未来版本中加快此过程。
CUDA 11.x 中 JIT LTO 不支持 CUDA 次要版本兼容性；12.0 将具有更强的兼容性。

包含引用信息的 LTO

从 CUDA 11.7 开始，nvcc 将跟踪主机对设备代码的引用，LTO 可以利用此信息来移除未使用的代码。
JIT LTO 需要用户告知此信息，因此引入了新的 cuLinkCreate 选项：
- CU_JIT_REFERENCED_KERNEL_NAMES
- CU_JIT_REFERENCED_VARIABLE_NAMES
- CU_JIT_OPTIMIZE_UNUSED_DEVICE_VARIABLES
- “NAMES”字符串使用隐式通配符，因此“foo”将匹配像“Z3foo!”这样的混淆名称。