怎样管理网站,iis搭建网站教程,口碑好企业网站建设,搭建网站内链系统C17 parallel algorithms and HIPSTDPAR — ROCm Blogs (amd.com) C17标准在原有的C标准库中引入了并行算法的概念。像std::transform这样的并行版本算法保持了与常规串行版本相同的签名#xff0c;只是增加了一个额外的参数来指定使用的执行策略。这种灵活性使得已经使用C标准…C17 parallel algorithms and HIPSTDPAR — ROCm Blogs (amd.com) C17标准在原有的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求解器依赖于以下函数来检查各个城市排列并选择成本/距离最低的那一个。以下是不使用任何并行性的详细实现
templateint 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_iteratorN it(i);// first city visitedint from it.first();// visited all other cities in the chosen route// and compute costwhile (!it.done()){int to it.next();cost distances[to N*from];from to;}// update best_route - reductionreturn 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!个路径彼此独立计算它们的各自成本是一种令人尴尬的并行操作。C17允许开发人员仅通过在算法调用时传递执行策略作为第一个参数来轻松并行化前面的代码。C17标准定义了三种可能的执行策略
std::execution::sequenced_policy及相应的策略对象作为参数std::execution::seqstd::execution::parallel_policy及相应的策略对象作为参数std::execution::parstd::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 CHANGEcounting_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-31Checking results...
Expected 9.546235158221428e01
Actual 9.546235158231138e01
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-31Checking results...
Expected 9.546235158221428e01
Actual 9.546235158231144e01
This run PASSED (Difference is within 0.00000000%)
两个版本之间的性能差异源于处理非常驻内存初次分页的开销。为了“使事情更公平”可以调整HIP版本也使用 hipMallocManaged() 而不是 hipMalloc()。这种特定的配置已经在 TeaLeaf 的 HIP 版本中可用并且可以通过在编译时传递一个简单的标志来启用。以下是当使用 hipMallocManaged() 和 XNACK 为所有GPU分配时 TeaLeaf 的 HIP 版本的输出。
Timestep 10CG: 3679 iterationsWallclock: 39.573sAvg. time per cell: 2.473331e-06Error: 9.962546e-31Checking results...Expected 9.546235158221428e01Actual 9.546235158231144e01This 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_XNACK1之上的可以重试的页错误处理机制可以进行页面迁移。这种模式被称为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 vectorint v) {return any_of(execution::par_unseq, cbegin(v), cend(v), [](auto x) {return x foo;});
}bool only_in_hmm_mode(const vectorint 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 vectorint 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-archgfx90a 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上与我们联系。
