C++17 并行算法与 HIPSTDPAR

C++17 标准在现有的 C++ 标准库中添加了并行算法的概念。像std::transform这样的算法的并行版本，除了增加一个指定要使用的执行策略的额外参数外，其签名与常规串行版本保持一致。这种灵活性允许已经在使用C++ 标准库算法的用户，通过引入最小的代码更改来利用多核架构。

从 ROCm 6.1 开始，只要用户愿意添加一两个额外的编译器标志，并行算法就可以通过 HIPSTDPAR 无缝地卸载到 AMD 加速器。

虽然 HIPSTDPAR 引入的功能适用于所有 AMD GPU（包括消费级显卡），但本文将重点介绍使用 ROCm 6.1 的 AMD CDNA2™ 和 CDNA3™ 架构（分别是 MI200 和 MI300 系列显卡）。作为代码示例，我们重点关注可在此处找到的旅行商问题 (TSP) 求解器。

旅行商问题

旅行商问题试图回答以下问题：“给定一个城市列表以及每对城市之间的距离，要找到一条访问每个城市一次并返回原点的最短可能路线？”。由于指数级复杂度，这个问题尤其难以解决（NP-hard）；在列表中增加一个城市会导致需要检查的组合数量呈指数级增长。对于超过 17 或 18 个城市的 TSP 问题，仅通过枚举所有可能的组合并逐一检查来解决该问题在计算上是不可行的。对于实际应用，会使用高级方法（割平面和分支定界技术），但在本文中，我们专注于一种非常易于并行实现的蛮力方法。

我们分析的 TSP 求解器依赖于以下函数来检查各种城市排列并选择成本/距离最低的路线。这是一个未使用任何并行化的详细实现：

template<int N>
route_cost find_best_route(int const* distances)
{
  return std::transform_reduce(
    counting_iterator(0),
    counting_iterator(factorial(N)),
    route_cost(),
    [](route_cost x, route_cost y) { return x.cost < y.cost ? x : y; },
    [=](int64_t i) {
    int cost = 0;

    route_iterator<N> it(i);

    // first city visited
    int from = it.first();

    // visited all other cities in the chosen route
    // and compute cost
    while (!it.done())
    {
      int to = it.next();
      cost += distances[to + N*from];
      from = to;
    }

    // update best_route -> reduction
    return route_cost(i, cost);
  });
}

std::transform_reduce 算法执行两个操作：

1. 一个转换（等效于 map 操作），由作为最后一个参数传递的 lambda 函数实现；

2. 一个归约操作，由作为第四个参数传递的 lambda 函数表示。

上述函数遍历从 0 到 N! 的所有元素，每个元素代表一个特定的城市排列，计算该特定路径的成本，并返回一个 `route_cost` 对象实例，该对象包含特定路径的 ID 和相关成本。最后，通过比较各种路径的成本并选择成本最低的路径来进行归约。

在 AMD Zen4 处理器上，此串行代码大约需要 11.52 秒来计算涉及十二个城市的 TSP 实例的最佳路径。对于涉及十三个城市的 TSP 实例，相同代码大约需要 156 秒。这是由于 TSP 带来的搜索空间指数增长的正常后果。

执行策略和 HIPSTDPAR

由于 N! 条路径中的每一条都是独立的，计算它们的个体成本是一个非常易于并行的操作。C++17 允许开发人员通过将执行策略作为算法调用的第一个参数来轻松并行化之前的代码。C++17 标准定义了三种可能的执行策略：

• std::execution::sequenced_policy 和对应的策略对象，用作参数 std::execution::seq

• std::execution::parallel_policy 和对应的策略对象，用作参数 std::execution::par

• std::execution::parallel_unsequenced_policy 和对应的策略对象，用作参数 std::execution::par_unseq

执行策略允许用户向实现者传达用户代码应强制执行/维护的不变量信息，从而使后者能够可能采用更有利/更高效的执行方式。

std::execution::sequenced_policy

顺序策略限制实现者在调用算法的线程上执行所有操作，从而抑制了可能的并行执行。所有操作都在调用线程内以不确定顺序执行，这意味着在同一线程内，对同一算法的后续调用可能会以不同的顺序执行其操作。

std::execution::parallel_policy

并行策略允许实现者采用并行执行。操作可以在调用算法的线程上执行，也可以在标准库实现创建的线程上执行。对于执行算法调用所描述的计算的所有线程，所有操作都在一个线程内以不确定顺序执行。此外，元素访问函数调用本身没有排序保证。与顺序策略相比，对算法使用的各种组件施加了额外的约束。特别是，对迭代器、值和可调用对象的*操作*以及它们的传递闭包必须是数据竞争自由的。

在前面的示例中，可以通过将 `std::execution::par` 策略作为第一个额外参数传递来并行化 `find_best_route` 函数，如下所示：

return std::transform_reduce(
  std::execution::par, // THE SIMPLE CHANGE
  counting_iterator(0),
  counting_iterator(factorial(N)),
  route_cost(),
  [](route_cost x, route_cost y) { return x.cost < y.cost ? x : y; },
  [=](int64_t i)

通过这个简单的更改，代码现在将在所有可用的 CPU 核心上运行。在配备 48 个 Zen4 逻辑核心的 MI300A 的 CPU 部分，解决具有 12 个城市的 TSP 实例大约需要 0.34 秒。这次并行运行比串行版本快了近 34 倍（串行版本需要 11.52 秒）！对于具有 13 个城市的 TSP 实例，并行版本大约需要 5 秒。最后，对于涉及 14 个城市的更大问题，48 个 Zen4 逻辑核心大约需要 77 秒。

std::execution::parallel_unsequenced_policy

此策略保证了用户提供的可调用对象满足最严格的要求。使用此策略调用的算法可以以无序和不序列化（ względlə to one another）的方式执行步骤。这意味着各种操作可以在同一线程上交错执行。此外，任何给定的操作都可能在一个线程上开始，在另一个线程上结束。指定并行无序策略时，用户保证不使用需要调用与另一个函数同步（synchronizes-with）的函数的操作。实际上，这意味着用户代码不执行任何内存分配/释放，仅依赖于无锁（lock-free）的 std::atomic 特化，并且不依赖于 std::mutex 等同步原语。

此策略是当前唯一可以选择以卸载 到 AMD 加速器的并行性。要触发使用并行无序策略调用的所有并行算法的 GPU 卸载，必须在编译时传递 —hipstdpar 标志。此外，对于当前默认值（gfx906）以外的 GPU 目标，用户还必须传递 —offload-arch= 来指定正在使用的 GPU。

在 MI300A 上，只需切换策略并使用上述标志重新编译，对于具有 13 个城市的 TSP 实例，执行时间将降至 0.5 秒。使用 14 个城市时，与在 48 个 Zen4 逻辑核心上运行的并行版本所需的 77 秒相比，使用 MI300A 的 GPU 部分可将执行时间缩短到 4.8 秒。由于每个人都喜欢一张好的表格，让我们通过总结从 CPU 上的顺序执行到卸载到加速器的并行无序执行的进展来结束本节。

TeaLeaf

一个更复杂的示例，展示了 HIPSTDPAR 的使用和性能，是 TeaLeaf。该代码是英国布里斯托大学 TeaLeaf 热传导迷你应用的 C++ 实现。多种实现方式说明了各种并行编程范式，包括 HIP 和并行化标准算法。这使我们能够对优化后的、基于 HIP 的实现与 HIPSTDPAR 实现进行公平的性能比较。在本测试中，我们选择了 `tea_bm_5.in` 基准测试，它包含一个 4000×4000 单元的二维网格和 10 个时间步。

对于 HIPSTDPAR 版本，在 MI300A 卡上，获得以下输出：

Timestep 10
CG:                    3679 iterations
Wallclock:             40.884s
Avg. time per cell:    2.555271e-06
Error:                 9.805532e-31

Checking results...
Expected 9.546235158221428e+01
Actual   9.546235158231138e+01
This run PASSED (Difference is within 0.00000000%)

对于 HIP 版本，其表现如下：

Timestep 10
CG:                    3679 iterations
Wallclock:             34.286s
Avg. time per cell:    2.142853e-06
Error:                 9.962546e-31

Checking results...
Expected 9.546235158221428e+01
Actual   9.546235158231144e+01
This run PASSED (Difference is within 0.00000000%)

两个版本之间的性能差异源于处理初始页面加载非驻留内存所产生的开销。为了“拉平”，HIP 版本可以通过使用 hipMallocManaged() 而不是 hipMalloc() 来进行调整。这种特定配置已在 HIP 版本的 TeaLeaf 中提供，并且可以通过在编译时传递一个简单的标志来启用。以下是使用 hipMallocManaged() 和 XNACK 进行所有 GPU 分配时，HIP 版本的 TeaLeaf 的输出。

Timestep 10
 CG:                    3679 iterations
 Wallclock:             39.573s
 Avg. time per cell:    2.473331e-06
 Error:                 9.962546e-31

 Checking results...
 Expected 9.546235158221428e+01
 Actual   9.546235158231144e+01
 This run PASSED (Difference is within 0.00000000%)

正如预期的那样，引入 hipMallocManaged() 后 HIP 版本的性能与 HIPSTDPAR 版本观察到的性能相当。最后，我们注意到正在进行的工作有望减少开销，从而使卸载版本更接近 HIP 版本。

HIPSTDPAR 的细节

将 C++ 标准并行算法执行卸载到 GPU 的能力依赖于 LLVM 编译器、HIPSTDPAR 和 rocThrust 之间的交互。从 ROCm 6.1 开始，用于编译常规 HIP 代码的 LLVM 编译器能够在传递 —hipstdpar 标志时，将采用 `parallel_unsequenced_policy` 执行策略的标准算法的调用转发给 HIPSTDPAR 头文件库。该头文件库负责将 C++ 标准库使用的并行算法映射到等效的 rocThrust 算法调用。这种非常简单的设计使得并行标准算法的卸载实现具有低开销。这时一个自然的问题是：“计算固然好，但它操作的内存呢？” 默认情况下，HIPSTDPAR 假定底层系统启用了 HMM (Heterogeneous Memory Management)，并且页面迁移可以通过建立在 XNACK 之上的可重试页面错误处理来实现（例如，`export HSA_XNACK=1`）。这种特定模式被称为HMM 模式。

当满足这两个要求时，卸载到 GPU 的代码（通过 rocThrust 实现）会触发页面迁移机制，数据将自动从主机迁移到设备。在 MI300A 上，尽管物理迁移既不需要也不有用，但通过 XNACK 处理页面错误仍然是必要的。有关页面迁移的更多详细信息，请参阅以下博文。

在没有 HMM / XNACK 的系统上，我们仍然可以通过传递一个额外的编译标志：—hipstdpar-interpose-alloc 来使用 HIPSTDPAR。此标志将指示编译器将所有动态内存分配替换为 HIPSTDPAR 头文件库中实现的兼容的 hipManagedMemory 分配。例如，如果正在编译的应用程序或其任何传递包含通过 operator new 进行内存分配，则该调用将被替换为对 `__hipstdpar_operator_new` 的调用。通过查看HIPSTDPAR 库中的该函数实现，我们可以看到实际的分配是通过 hipMallocManaged() 函数执行的。在非 HMM 启用的系统上这样做，主机内存会被固定（pinned），并且 GPU 可以直接访问，而无需通过页面错误驱动的迁移到 GPU 内存。这种特定模式被称为交织模式 (Interposition Mode)。

限制

对于 HMM 和交织模式，均适用以下限制：

1. 函数指针及其所有相关功能（例如，动态多态）不能被用户提供的可调用对象（直接或间接）使用，该可调用对象传递给算法调用；

2. 全局/命名空间作用域/`static`/`thread` 存储持续时间变量不能被用户提供的可调用对象（直接或间接）通过名称使用；

   • 当在HMM 模式下执行时，它们可以在地址中使用，例如。

     namespace { int foo = 42; }
     
     bool never(const vector<int>& v) {
     return any_of(execution::par_unseq, cbegin(v), cend(v), [](auto&& x) {
     return x == foo;
     });
     }
     
     bool only_in_hmm_mode(const vector<int>& v) {
     return any_of(execution::par_unseq, cbegin(v), cend(v),
     [p = &foo](auto&& x) { return x == *p; });
     }

3. 只有使用 `parallel_unsequenced_policy` 调用的算法才是卸载的候选算法；

4. 只有使用模拟 random_access_iterator 的迭代器参数调用的算法才是卸载的候选算法；

5. 用户提供的可调用对象不能使用异常；

6. 用户提供的可调用对象不能使用动态内存分配（例如，operator new）；

7. 无法进行选择性卸载，即无法指示仅执行使用 `parallel_unsequenced_policy` 调用的某些算法到加速器。

除了上述之外，使用交织模式 (Interposition Mode) 还会施加以下额外限制：

1. 所有期望互操作的代码都必须使用 `—hipstdpar-interpose-alloc` 标志重新编译，即组合独立编译的库是不安全的；

2. 自动存储持续时间（即栈分配）变量不能被用户提供的可调用对象（直接或间接）使用，例如。

   bool never(const vector<int>& v, int n) {
   return any_of(execution::par_unseq, cbegin(v), cend(v),
   [p = &n](auto&& x) { return x == *p; });
   }

为什么选择它？

在经历了如此快速的介绍之后，问“但这对我，一个 C++ 开发者有什么好处？”是可以理解的。 HIPSTDPAR 的目标是让任何使用标准算法的 C++ 开发人员能够利用 GPU 加速，而无需任何认知过载。应用程序开发人员可以牢牢地留在标准 C++ 世界中，而无需进入 GPU 特有语言（例如 HIP 或 SYCL）的全新领域。对我们来说幸运的是，我们选择的示例允许我们对该目标有多接近进行一些有限的、定量的洞察。Tealeaf 的作者通过多种编程接口实现了求解器，这意味着我们可以使用 cloc 工具来计算 `tsp.cpp` 实现所需的代码行数：

很明显，使用编译器标志驱动的卸载（由 HIPSTDPAR 启用）可以节省大量的打字量——例如，与 SYCL 相比节省了 57%。这使得走向 GPU 加速执行的旅程更加自然。因此，程序员至少在开始时可以专注于算法/问题解决，并发现对 GPU 有利的通用算法优化，而无需一头扎进 GPU 的“奥秘”中。

简而言之，直接告诉我如何提速

最初，HIPSTDPAR 在 Linux 上得到官方支持，Windows 支持将在未来推出。从已为 ROCm 设置的环境开始，使用包管理器安装 `hipstdpar` 包通常会带来所有必需的功能。此外，在撰写本文时，由于标准库实现细节（参见例如 注 3），存在对 TBB 的依赖。因此，有必要安装系统的 TBB 包（例如，Ubuntu 上的 libtbb-dev）。有了这些工具，并假设我们有一个 `main.cpp` 文件，该文件使用一些标准算法来解决给定问题，编译器驱动程序调用

clang++ --hipstdpar main.cpp -o main

如果我们的目标是与 gfx906 ISA（即 Vega20）兼容的 GPU，则会自动卸载所有使用 `std::execution::parallel_unsequenced_policy` 执行策略的算法调用。否则，我们还必须指定卸载目标：

clang++ --hipstdpar --offload-arch=gfx90a main.cpp -o main

结论

在本文中，我们提供了对 ROCm 支持 C++ 标准并行算法卸载的高层概述，旨在展示现有 C++ 开发人员如何在无需采用任何新的、特定于 GPU 的语言（例如 HIP）或指令（例如 OpenMP）的情况下，利用 GPU 加速。

我们认为，这种标准、极其易于访问的硬件并行性利用方式，对于针对 MI300A 加速器的应用程序尤其有益，因为 CPU 和 GPU 共享同一池 HBM。尽管今天没有演示，但 APU 架构和 HIPSTDPAR 的结合可以实现 CPU 和 GPU 之间的细粒度协作，它们成为真正的对等体，可以通过统一的编程接口进行访问。

有关 HIPSTDPAR 编译器方面支持的深入了解，有兴趣的读者应查阅相关的 AMD-LLVM 文档。

如果您有任何疑问，请在 GitHub 讨论区与我们联系。