Evaluating Application Characteristics for GPU Portability Layer Selection

本文通过分析来自主要高能物理实验的代表性异构应用，旨在识别影响不同 GPU 可移植性层性能的关键特征，从而指导开发者根据其特定需求选择最合适的技术。

要点

想象一下你是一位正试图烹饪一场盛大宴会的厨师。你的厨房里有三种不同类型的高功率烤箱：一台由 NVIDIA 制造，一台由 AMD 制造，还有一台由 Intel 制造。每台烤箱烹饪食物的方式都不同，使用的旋钮也不同，并且需要不同的食谱才能发挥最佳性能。

如果你专门为 NVIDIA 烤箱编写食谱（使用一种叫做 CUDA 的语言），你就不能直接把同样的食谱放入 AMD 或 Intel 的烤箱中。你必须重写整个东西。这是一个问题，因为你并不总是知道明天你的厨房里会有哪种烤箱。

为了解决这个问题，论文讨论了“可移植层”（portability layers）。把它们想象成 通用翻译器 或 智能适配器。它们能让你编写一个主食谱，然后由翻译器将其转换为每台烤箱都能理解的特定语言。论文研究了几个这样的翻译器（如 Kokkos、SYCL、OpenMP 和 Alpaka），以观察哪一个最适合不同类型的烹饪任务。

以下是作者在用高能物理实验（例如用于研究亚原子粒子的实验）中的真实“食谱”测试这些翻译器时发现的结果：

1. “启动时间”问题

启动 GPU（烤箱）并不是瞬间完成的。它需要几毫秒的时间来唤醒并做好准备。

• 问题所在： 有些翻译器在启动烹饪过程时非常缓慢。例如，在使用 AMD 烤箱时，Kokkos 会导致显著的延迟。如果你的烹饪任务非常短（比如煮一个 10 秒钟的鸡蛋），而翻译器光是启动炉灶就花了 5 秒钟，那么你就浪费了一半的时间。
• 教训： 如果你的任务非常微小且快速，请避开那些会让启动变慢的翻译器。

2. “拥挤厨房”问题

在真实的物理实验室中，GPU 并不是独自工作的。它是更大系统的一部分，许多人（线程）试图同时使用这台烤箱。

• 问题所在： 一些翻译器处理不好人群。例如，Kokkos 有一条规则说：“一次只能有一个人与烤箱交谈”，这会导致当多个厨师试图同时启动任务时出现交通堵塞。SYCL 则有些不稳定；有时它允许所有人同时烹饪，有时又会强迫他们在排队等待，这取决于你使用的是哪个版本的翻译器。
• 教训： 如果你的应用程序需要许多人同时工作，你需要一个懂得如何管理繁忙厨房而不锁上门的翻译器。

3. “工具箱兼容性”问题

物理食谱通常使用特殊的工具（库，如 ROOT 或 Eigen）来辅助数学运算和数据处理。

• 问题所在： 一些这些工具与翻译器配合得并不好。例如，一个流行的数学工具 Eigen 在使用 NVIDIA 编译器（许多翻译器依赖于此）时经常会出错。此外，尝试在一个项目中同时使用两个不同的编译器（一个用于 CPU，一个用于 GPU）就像是用两套不匹配的蓝图来盖房子——这会让建筑过程（构建软件）变成一场噩梦。
• 教训： 在选择翻译器之前，请检查你喜爱的工具是否能装进它的框架里。

4. “家具摆放”问题

GPU 喜欢简单、平坦的布局。它们更喜欢数据像一排整齐的盒子那样排列。然而，物理数据往往呈现出复杂、凌乱的形状（比如一堆大小不一的行李箱）。

• 问题所在： 翻译器试图通过将数据包装在特殊的容器中来修复这种混乱。虽然这使得代码具有可移植性，但它增加了“开销”——就像在移动任何东西之前，都要先把每一件物品都装进行李箱，即使你只需要移动一只袜子。这会减慢速度。此外，没有任何一个翻译器擅长处理“锯齿状”数据（长度不一的行），而这在物理学中非常常见。
• 教训： 如果你的数据复杂且凌乱，翻译器可能会为了整理它们而拖慢你的速度。

5. “专业工具”问题

有时你需要特定的工具，比如随机数生成器（RNG）或快速傅里叶变换（FFT）。

• 问题所在： 每个烤箱制造商都有自己专属的、超快速的专业化工具。通用翻译器通常不包含这些专门的工具，或者使用的是它们自己较慢的版本。虽然你可以强制翻译器使用烤箱的原生工具，但这会破坏“可移植性”，因为那个工具只适用于那台特定的烤箱。
• 教训： 如果你严重依赖这些特定工具，你可能必须在速度（使用烤箱的原生工具）或可移植性（使用翻译器的通用工具）之间做出选择。

6. “建设时间”与“搬家”问题

• 构建食谱： 一些翻译器会让“烹饪时间”（编译时间）变得更长。对于大型项目，使用某些翻译器可能会让构建过程从几分钟变成几小时。
• 搬迁厨房： 如果你为特定的烤箱（例如 NVIDIA V100）构建了软件，它可能无法在更新的烤箱（例如 NVIDIA A100）上运行。这要求你为可能遇到的每一种类型的烤箱都构建一个单独的版本，这为向不同实验室分发软件带来了巨大的物流难题。

最终结论

论文得出结论：不存在“完美”的翻译器。

• Kokkos 对某些任务表现出色，但在特定硬件上的并发处理和启动时间方面存在困难。
• SYCL 功能强大，但根据编译器版本的不同，其表现可能并不一致。
• OpenMP 及其他翻译器在处理内存和不同硬件方面各有优劣。

核心要点： 你不能仅仅因为某个翻译器很流行就去选择它。你必须观察你的特定“食谱”（你的应用程序）。如果你的代码短促且快速，请选择启动时间低的翻译器。如果你的代码复杂且使用了许多工具，请选择能与这些工具良好配合的翻译器。

作者还指出，这些技术正在迅速发展，就像每年都会推出新模型的烤箱一样。今天的最佳选择在明天可能会发生变化，因此开发者需要持续关注这一领域的发展。虽然未来的新标准可能会让这些选择变得更容易，但目前，仔细的测试是找到合适方案的唯一途径。

技术摘要

技术摘要：评估应用程序特性以进行 GPU 可移植性层选择

问题陈述

高能物理（HEP）计算正日益依赖 GPU 来提升性能，但硬件格局正在向包括 AMD 和 Intel 在内的多样化方向发展。虽然特定厂商的语言（如 CUDA）能提供高性能，但它们限制了在特定硬件上的部署，降低了可移植性。相反，可移植性层（如 Kokkos、SYCL、HIP、OpenMP、Alpaka 和 std::par）承诺能够在多种 CPU 和 GPU 架构上执行。然而，这些层在功能上并不等同；它们在性能、并发支持和与现有代码库的兼容性方面表现出截然不同的特征。核心问题在于，缺乏一个清晰的框架，帮助开发者根据特定的应用程序特性、内核结构和运行时环境来选择最合适的移植层。

方法论

高能物理计算卓越中心（HEP-CCE）通过使用来自主要 HEP 实验的代表性异构应用程序进行了对比研究：

• ATLAS: FastCaloSim（参数化量热计模拟）。
• CMS: Patatrack（包含 40 多个内核、多线程和内存池的异构像素重建）以及 P2R（单内核径迹重建）。
• DUNE: Wire-Cell Toolkit（栅格化、scatter-add 和 FFT 卷积内核）。
• 事件生成器: Sherpa（Pepper）和 Madgraph 的 GPU 移植版本（利用自动生成的代码）。

这些测试平台从串行 C++ 被移植到各种可移植技术，旨在识别显著影响性能和兼容性的代码特征。该研究在多个维度对这些应用程序进行了评估，包括内核启动延迟、并发处理、库兼容性、内存管理和编译开销。

关键发现与结果

1. 内核启动延迟

内核启动并非瞬时完成，其范围从微秒到数十微秒不等。

• Kokkos: 被观察到会增加显著的开销（数十微秒），特别是在 AMD GPU 上，而 AMD GPU 本身的原生启动延迟已经高于 NVIDIA 或 Intel。
• SYCL, HIP, OpenMP, Alpaka: 与原生实现相比，并未表现出巨大的额外惩罚。
• AMD 特定情况: 在 AMD GPU 上调用 rocRAND 会导致极高的首次调用延迟（超过 1 秒），这表明如果使用会加剧启动开销的可移植层，短运行的内核可能会遭受严重的性能下降。

2. 并发与线程池

HEP 应用程序通常运行在利用多线程（如 Intel TBB）的更大框架内。

• Kokkos: 展示了与外部线程池的显著不兼容性。其串行后端使用的锁会使调用序列化，而其线程后端则禁止外部线程调用。目前，仅通过使用特定架构 API 的 CUDA/HIP 后端，并发内核启动才是可行的，这限制了可移植性。
• SYCL: 实现情况不一致。根据编译器制造商和版本的不同，某些实现会对来自不同线程的并发调用进行序列化处理。研究时，AMD 架构似乎缺乏并发的 SYCL 解决方案。

3. 外部库与编译器兼容性

• Eigen: 许多版本与 NVIDIA 的 nvcc 不兼容，从而影响了依赖于它的可移植层（如 Kokkos, Alpaka）。它目前也与 OpenMP offload 不兼容。
• ROOT: 无法使用 NVIDIA 的 nvc++ 进行完整编译。混合编译（使用 g++ 和 nvc++）是可能的，但会使构建规则复杂化，并需要精细的内存分配管理（GPU 数据必须由 nvc++ 分配）。
• 构建复杂度: 在一个项目中支持多个可移植层需要复杂的构建规则（例如使用 #ifdef 或文件分离），以避免像 Kokkos 和 CUDA 这样的层之间的头文件冲突。

4. 数据结构与内存传输

• 面向对象 vs. 平坦布局: HEP 中常见的复杂面向对象数据结构在 GPU 上表现不佳，GPU 更倾向于平坦布局或结构体数组（SoA）。
• 可移植性开销: 诸如 Kokkos 和 SYCL 之类的层提供了抽象（例如 Kokkos::Views, sycl::buffers），这增加了分配和传输开销，特别是对于小对象而言。
• 初始化: Kokkos::Views 的默认内存初始化可能成为性能瓶颈；这可以通过使用 Kokkos::ViewAllocateWithoutInitializing 来缓解。
• 统一共享内存 (USM): 观察到 std::par 使用的自动 USM 传输比显式传输更慢，并且可能不会在最佳时机触发。
• 锯齿状数组 (Jagged Arrays): 没有 API 原生支持大小不一的多维向量（锯齿状数组），需要进行手动构建、维度缩减或填充。

5. RNGs、FFTs 与原子操作

• 原生库: 厂商特定的库（cuRAND, rocRAND, oneMKL）提供调优后的性能，但缺乏可移植性。虽然 Kokkos 提供了跨架构的 RNG/FFT 实现，但它们可能无法达到原生性能。
• HEP-CCE 贡献: 该小组开发了一个头文件库，封装了设备原生 RNG，具有统一的 API 和硬件原生的 oneMKL 后端，允许 SYCL 调用跨厂商的原生实现。
• 原子操作: 原子操作在 GPU 硬件上的映射效果较差。性能因编译器和硬件的不同而差异巨大（例如，在 NVIDIA 上使用 OpenMP/clang13 的 atomic<int> 比 CUDA 慢 35 倍）。整数原子操作通常比浮点/双精度原子操作性能更差。

6. 编译时间

• 编译器开销: nvc++ 通常比 g++ 慢 2 到 3 倍。Kokkos 增加了 10% 到 20% 的开销。
• 重度模板代码: Madgraph 利用了大量模板化的自动生成代码，在处理大型内核时，其使用 SYCL 的 icpx 后端的编译时间比其他后端慢了 1000 倍以上。

7. 运行时配置

• Kokkos: 需要使用特定的硬件标识符进行构建。为一种 GPU（如 V100）构建的二进制文件可能无法在另一种 GPU（如 A100）上运行。支持共享库和多种架构会产生复杂的构建配置矩阵，给二进制分发带来挑战。
• Alpaka: 需要针对每个后端编译器进行单独编译。
• SYCL/OneAPI: 多平台二进制文件正在涌现，但通常局限于特定架构或需要编译时识别。
• OpenMP: 某些实现（LLVM Clang）允许在编译时列出多种架构，并在运行时动态选择后端。

重要性与结论

论文得出结论：没有任何一种单一的 GPU 可移植性解决方案适用于所有场景。可移植层的选择必须由应用程序的具体特性驱动，包括内核运行时、并发需求、数据结构复杂度和外部库依赖。

作者强调，生态系统正在快速演进。无论是可移植层（Kokkos, SYCL 等）还是原生编译器（clang, nvcc, hipcc）都在不断改进，趋势是向着增加功能支持和优化性能方向发展。论文主张频繁监控这些技术，以确保代码库与可移植层之间的最佳匹配。此外，作者指出，新兴的 C++ 标准（如用于异步执行的 P2300）未来可能会简化这些选择，尽管广泛的硬件支持仍需数年时间。

最终，如果不充分了解代码的本质、外部依赖关系及运行时环境就选择可移植层，可能会在性能和部署灵活性方面面临重大风险。