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工作流程

如何使用 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流程
    部分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性能提升图表
    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 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 调用代码

数学库中的 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 回调流程

JIT LTO 结果: 数学库回调

  • 间接回调 vs JIT LTO 回调 (A100 (40 GB))
  • 在不同的 FFT 大小下，JIT LTO 回调显示出显著的性能提升 (以 GFLOPS 衡量)。
  • JIT LTO 数学库回调性能
    • 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）。

JIT LTO 用例：数学库组合
JIT LTO 用例：数学库组合

何时使用 (或不使用) 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!”这样的混淆名称。
LTO 包含引用信息
LTO 包含引用信息

相关出版物链接

欲了解更多信息，请参考以下链接：

博客：

手册：
* https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#optimization-of-separate-compilation -- nvcc
* https://docs.nvidia.com/cuda/cuda/nvrtc/index.html -- nvrtc
* https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MODULE.html#group__CUDA__MODULE -- cuLink APIs
* https://docs.nvidia.com/cuda/deploy/cuda-compatibility/ -- 兼容性保证

用例：
* https://www.nvidia.com/en-us/on-demand/session/gtcfail21-a31155?playlistId=playlist-ead11304-9931-4e91-9d5a-fb0e1ef27014 -- “JIT LTO Adoption in cuSPARSE/cuFFT: Use Case Overview” in GTC Fall 2021.
* https://www.osti.gov/biblio/1798430-enhancements-supporting-ic-usage-pem-libraries-next-gen-platforms -- 提及使用 LTO 的 LLNL 论文。