C++17并行算法与HIPSTDPAR

C++17 parallel algorithms and HIPSTDPAR — ROCm Blogs (amd.com)

C++17标准在原有的C++标准库中引入了并行算法的概念。像std::transform这样的并行版本算法保持了与常规串行版本相同的签名，只是增加了一个额外的参数来指定使用的执行策略。这种灵活性使得已经使用C++标准库算法的用户只需对代码进行最小程度的修改，就能利用多核架构的优势。

从ROCm6.1开始，只要用户愿意添加一两个额外的编译器标志，这些并行算法就能通过HIPSTDPAR无缝卸载到AMD加速器上执行。尽管HIPSTDPAR提供的功能适用于所有AMD GPU（包括消费级显卡），但本篇博客主要聚焦于使用ROCm6.1的AMD CDNA2™和CDNA3™架构（分别对应MI200和MI300系列卡）。作为示例代码，我们将关注这里提供的旅行商问题（TSP）求解器。

旅行商问题

旅行商问题旨在回答这样一个问题：“给定一系列城市及其之间的距离，访问每个城市恰好一次并返回出发城市的最短可能路线是什么？”由于指数级的复杂性，这一问题特别难以解决（属于NP难问题），每增加一个城市，需要检查的组合数量就会呈指数增长。对于超过17或18个城市的实例，仅通过枚举所有可能组合并检查每一种的计算成本是不可行的。在实际应用中，会采用高级方法（如切割平面法和分支定界技术），但为了本文的目的，我们关注的是TSP的一个极度并行化的暴力求解实现。

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. 由作为最后一个参数传递的lambda函数实现的转换（相当于map操作）；
2. 由作为第四个参数传递的lambda函数表示的缩减操作。

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

在AMD Zen4处理器上，这个串行代码计算涉及12个城市的TSP实例的最佳路径大约需要11.52秒。同样的代码计算涉及13个城市的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秒。与串行版本所需的11.52秒相比，这种并行运行速度提高了近34倍！对于涉及13个城市的TSP实例，并行版本大约需要5秒。最后，对于涉及14个城市的更大问题，48个Zen4逻辑核心大约需要77秒。
std::execution::parallel_unsequenced_policy
此策略确保用户提供的代码满足最严格的要求。使用此策略调用的算法可能会以无序且未排序的方式执行各个步骤。这意味着各种操作可以在同一线程上相互交错。此外，任何给定的操作都可以在一个线程上开始并在另一个线程上结束。在指定并行未排序策略时，用户保证不采用涉及调用与另一个函数同步的函数的操作。实际上，这意味着用户代码不执行任何内存分配/释放，仅依赖于std::atomic的无锁特化，并且不依赖于诸如std::mutex之类的同步原语。

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

在MI300A上，只需切换策略并使用上述标志重新编译，涉及13个城市的TSP实例的执行时间就减少到0.5秒。当使用14个城市时，使用MI300A的GPU部分将执行时间从并行版本的48个Zen4逻辑核心所需的77秒减少到4.8秒。因为每个人都喜欢一个好的表格，让我们通过总结从CPU上的顺序执行到卸载到加速器的并行未排序执行的进展来结束本节：

14-city TSP	Timing (s)
`seq`	2337
`par`	77
`par_unseq` on CPU	75
`par_unseq` on GPU	4.8

TeaLeaf

一个展示HIPSTDPAR使用和性能的更复杂例子是TeaLeaf。这个代码是来自英国布里斯托大学的TeaLeaf热传导小应用程序的一个C++实现。多个实现示例展示了各种并行编程范式，包括HIP和并行化标准算法。这使得我们能够在优化的基于HIP实现和一个基于HIPSTDPAR的实现之间，进行公平的性能比较。为了这个测试，我们选择了`tea_bm_5.in`基准测试，包含一个4000x4000单元的2D网格和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()。这种特定的配置已经在 TeaLeaf 的 HIP 版本中可用，并且可以通过在编译时传递一个简单的标志来启用。以下是当使用 hipMallocManaged() 和 XNACK 为所有GPU分配时 TeaLeaf 的 HIP 版本的输出。

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的核心机理

(“nuts and bolts”是一个英语习语，意思是“基本的工作原理”或者“详细的实际细节”。它通常用于描述某件事情的具体、基础的组成部分，而非它的广义或概念性内容。这个习语源自于实际的螺母（nuts）和螺栓（bolts），它们是用来物理连接各种构件的基础硬件。因此，在这种情况下，提及 “HIPSTDPAR的nuts and bolts” 就是指想要深入了解HIPSTDPAR如何实际工作的所有基础和技术细节。翻译成“HIPSTDPAR的核心机理”是为了捕捉这种想要详细了解它的工作原理和组件的意图。因此，“Nuts”（螺母）在这里不是指实际的螺母，而是泛指细节或基本部分；“bolts”（螺栓）也是如此。这个习语与硬件本身无关，而是关于理解事物的基本组成部分和工作方式。 )

HIPSTDPAR的C++标准并行算法执行能够卸载到GPU，取决于LLVM编译器、HIPSTDPAR以及rocThrust之间的交互。从ROCm6.1开始，用来编译常规HIP代码的LLVM编译器将可以在传递了`--hipstdpar`标记的情况下，将调用带有`parallel_unsequenced_policy`执行策略的标准算法转发至HIPSTDPAR头文件库。这个仅头文件的库负责将C++标准库使用的并行算法映射为等价的rocThrust算法调用。这种非常简单的设计为并行标准算法的卸载实现提供了低开销。此时一个自然的问题是：“计算是很好，但是它操纵的内存怎么办？”默认情况下，HIPSTDPAR假设底层系统启用了HMM（异构内存管理），并且通过实现在XNACK（例如，导出HSA_XNACK=1）之上的可以重试的页错误处理机制，可以进行页面迁移。这种模式被称为HMM模式。

当这两个要求都得到满趀时，卸载到GPU的代码（通过rocThrust实现）触发页面迁移机制，数据会自动从主机迁移到设备。在MI300A上，虽然物理迁移既不需要也没有用处，但通过XNACK处理页面错误仍然是必要的。有关页面迁移的更多详情，请参考以下博客文章。("Post"在这个上下文中通常被用来指代网站或者博客上发布的一篇文章或记录。在网络术语中，"post"可以是一个名词，指的就是所发布的内容；也可以是一个动词，指的是发布信息的行为。所以，链接AMD Instinct™ MI200 GPU memory space overview - amd-lab-notes - AMD GPUOpen引向的内容是AMD官方网站上的一个页面，其中包含了关于MI200系列加速器内存空间的详绽解读和如何启用页面迁移的信息，因此将其译为"博客文章"是合理的，因它确实是一篇公开发布、围绕具体主题的文章。翻译为"博客文章"是为了传达这是一篇可能相对非正式或以教育、信息共享为目的的文章，它与学术论文或新闻报道的语调和格式不同。在中文中，博客文章也常被简称为"博文"。在这种语境下，直接翻译为"文章"也是可以接受的，只是"博客文章"提供了关于文章发布平台的额外信息。)

在没有启用HMM/XNACK的系统上，我们仍然可以通过传递一个额外的编译标记：`--hipstdpar-interpose-alloc`来使用HIPSTDPAR。这个标记会指导编译器用在HIPSTDPAR头文件库中实现的兼容的hipManagedMemory分配来替换所有的动态内存分配。例如，如果正在编译的应用程序或其传递性包含的内容通过`operator new`分配了自由存储内存，那么这个调用会被替换为对`__hipstdpar_operator_new`的调用。通过查看该函数在HIPSTDPAR库中的实现，我们可以看到实际的分配是通过`hipMallocManaged()`函数执行的。这样在一个没有启用HMM的系统上，主机内存被固定并且能够直接被GPU访问，而不需要任何页面错误驱动的迁移到GPU内存。这种模式被称为"Interposition模式"。

限制

对于HMM和Interposition模式，以下限制适用：

1. 函数指针和所有相关功能，例如动态多态性，不能被传递给算法调用的用户提供的可调用对象（直接或传递性地）使用；
2. 全局/命名空间范围/静态/线程存储期变量不能被用户提供的可调用对象（直接或传递性地）以名称使用；
- 当在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模式还带来以下额外的限制：

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](GitHub - ROCm/roc-stdpar)的目标是让任何使用标准算法的C++开发者都能够在不增加认知负担的情况下利用GPU加速。应用程序开发者可以坚定地留在标准C++世界中，而不必跨入HIP或SYCL等GPU特定语言的崭新世界。幸运的是，我们的特定例子能够提供一些有限的、定量的洞察，让我们了解我们离这个目标有多近。Tealeaf的作者已经通过多种编程接口实现了求解器，这意味着我们可以使用`cloc`工具计算`tsp.cpp`实现所需的代码行数：

Programming Interface	LoC
Kokkos	145
OpenACC	142
OpenMP	116
Standard C++ Serial	112
Standard C++ Parallel Algorithms	107
SYCL	169

显然，使用由编译器标志驱动的卸载（如HIPSTDPAR所启用）可以节省相当多的打字工作 - 例如，与SYCL相比，最多可以节省57%。这使向GPU加速执行的转变过程更自然。因此，程序员至少在初始阶段可以专注于算法/问题解决，并发现适用于GPU的通用算法优化，而无需深入GPU“奥秘”。

太长了，没读。直接告诉我怎么快速进行吧

要快速地利用GPU加速，可以在Linux环境下使用HIPSTDPAR（未来也会支持Windows）。假设你已经根据ROCm的快速开始教程设置好了环境，通过包管理器安装`hipstdpar`包就可以得到所有必需的组件。由于标准库实现细节问题（参见例如备注3），可能还需要安装TBB库。例如，在Ubuntu上安装`libtbb-dev`。接着，如果你有一个主程序文件`main.cpp`，它使用标准算法解决了某个问题，那么只需简单地运行以下编译命令：

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

就可以自动将使用`std::execution::parallel_unsequenced_policy`执行策略的所有算法调用进行GPU加速，假定你的目标GPU与`gfx906`兼容（即Vega20）。如果是另外的GPU目标，则需要指定：

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

结论

我们在本文中提供了一个关于ROCm支持C++标准并行算法进行卸载加速的高层次概述，展示了现有的C++开发者如何利用GPU加速，而无需采用任何新的、特定于GPU的语言（例如HIP）或指令（例如OpenMP）。我们相信，这种标准且极其易于访问的方式，以利用硬件并行性，对于针对MI300A加速器的应用程序将特别有益，其中CPU和GPU共享同一池的HBM。虽然今天没有演示，但APU架构和HIPSTDPAR的结合可以实现CPU与GPU之间的细粒度合作，通过统一的编程接口使它们成为真正的对等实体。

如果想更深入了解HIPSTDPAR的编译器支持，可以阅读相关的AMD-LLVM文档。

本文作者感谢Bob Robey和Justin Chang对文章的有益审阅。如果有任何问题，可以在GitHub Discussions上与我们联系。