GPU负载均衡的研究与应用

79→一种GPU负载均衡穆罕默德·奥萨马mosama@ucdavis.edu加州大学戴维斯分校美国塞尔班·D加州sdporumbescu@ucdavis.edu大学戴维斯分校，美国加州John D.加州jowens@ucdavis.edu大学戴维斯分校戴维斯，加利福尼亚，美国摘要我们提出了一个GPU细粒度的负载平衡抽象，从工作处理中的负载平衡，旨在支持静态和动态调度与可编程接口，以实现新的负载平衡计划。在我们的工作之前，释放GPU在不规则问题上的潜力的唯一方法是通过特定于应用程序的紧耦合负载平衡技术来通过我们的开源负载平衡框架，我们希望在GPU上开发不规则并行算法时提高程序员的生产力，并通过允许快速试验各种现有负载平衡技术来提高此类应用程序的整体性能特征。此外，我们还希望通过在抽象中将负载平衡的关注点与工作处理分离开来，管理现有代码并将其扩展到未来的架构中变得更加容易。CCS概念：·计算方法共享内存算法。关键词：负载均衡，稀疏计算，GPU，调度1介绍图形处理单元（GPU）擅长并经常设计用于常规细粒度并行问题，例如通用矩阵乘法（GEMM）。在像GEMM这样的常规问题中，相邻线程具有相似或相同的工作负载，并且通常实现接近100%的峰值GPU理论性能。更具挑战性的是具有足够细粒度并行性但不规则并行性的应用程序。在这样的应用中，相邻线程分发声明本作品采用知识共享署名国际4.0许可协议进行许可PPoPP©2023版权归所有者/作者所有。ACM ISBN979-8-4007-0015-6/23/02。https://doi.org/10.1145/3572848.3577434以锁步方式运行将具有不同的工作负载--可能是不同的工作量--这使得在像GPU这样的高度并行机器上的高效实现成为一个重大挑战。考虑稀疏矩阵向量乘法（SpMV），其中稀疏矩阵A和密集向量x作为输入。 SpMV计算输出向量y = Ax，并且是不规则细粒度并行的示例。 与GEMM不同，SpMV中的稀疏矩阵可以在矩阵的行内包含不规则性：矩阵的行可以具有不同数量的非零条目。 一行到每个GPU线程的简单映射可能会暴露这种不规则性，其中相邻线程可能被分配不同数量的非零来处理，导致同一线程束1内的线程等待具有大量非零的线程。 由于这种不规则性而产生的不平衡--具体地说，当工作不均匀地分布在并行参与者之间时，因此，一些参与者空闲，而另一些参与者做更多的工作--被定义为负载不平衡问题。当前的实现使用特定于应用程序的负载平衡技术来解决GPU上的负载不平衡问题，该技术旨在均匀分布工作，使得每个线程获得相同数量的工作项以实现最大性能（例如，Merrill和Garland的负载平衡SpMV实现[ 20 ]）。这些负载平衡技术通常与应用程序本身紧密耦合。这些实现中的负载平衡组件既复杂又通常是应用程序性能的最重要贡献者。我们在这里的工作概括了今天的应用程序特定的负载平衡算法到一个干净的，模块化的，功能强大的抽象，可以应用于许多复杂的不规则的工作负载。在构建抽象的过程中，我们确定了常见的负载平衡方法目前部署在GPU上的稀疏，不规则的应用 程 序 中 ： 特 定 于 应 用 程 序 的 框 架 ， 如 GraphIt [5] ，Gunrock [29]和Graph-BLAST [31];来自低级CUDA库的技术，如ModernGPU [3]和CUB [24];以及应用程序中负载平衡算法的其他手工编码[13，16]和广度优先搜索[6，21]。我们证明了CUDA线程束是32个线程的集合，它们以锁步方式执行指令. 线程束中的线程是无发散的，并且以单指令多线程（SIMT）方式运行。PPoPPM. Osama，S.D. Porumbescu，和J.D. 欧文斯80×通过简单、直观、强大的抽象，这些负载平衡调度可以扩展到支持不规则的工作负载，这些不规则的工作负载比它们设计的特定问题更普遍我们证明了这一点，通过使用稀疏线性代数为中心的数据图遍历内核的负载平衡。编写高性能负载平衡代码是复杂的，很大程度上是因为这些代码必须执行许多角色。在其他任务中，它必须从特定的数据结构中摄取数据，对该数据执行用户定义的计算，并以负载平衡的方式调度该计算。在我们的抽象中，关键的洞察力是将工作负载映射（负载平衡任务）和工作执行（用户定义的计算）之间的关注点分开，我们将稀疏格式（例如压缩稀疏行（CSR））映射到简单的抽象组件，称为工作原子，瓦片和集合。这些基本组件表示为可组合的C++范围和基于范围的for循环，并用于构建负载平衡调度。然后，程序员可以使用这些API来构建负载平衡的高性能应用程序和原语。 以这种方式表示，我们可以重建现有的应用程序相关的负载平衡技术，解决更一般的，可移植的，和可编程的不规则性。我们的工作贡献如下：1. 我们提出了一种新的抽象GPU上的不规则并行工作负载 我们的抽象在一个高层次上允许程序员开发稀疏，不规则并行算法与最少的代码，同时提供高性能。2. 我们设计并实现了一组直观的API，可在我们的开源GPU负载平衡框架工作，建立在建议的抽象使用CUDA-C++范围和范围为基础的循环。3. 我们通过实现一个新的基于协作组的负载平衡调度来实现新的负载平衡调度，如第5.2节所述，这是对以前的线程级、线程束级和块级负载平衡调度的概括[30]。4. 我 们 提 供 了 最 先 进 的 SpMV 性 能 作 为 基 准 ，SuiteSparse矩阵集合[ 11 ]的几何加速比为2.7，而cuSparse的最先进的实现使用简单的算法和3 GPU负载平衡调度。2设计目标我们的编程模型侧重于细粒度嵌套数据并行的广泛类别负载平衡任务级并行性需要不同的方法，超出了本工作的范围。本节重点介绍负载平衡抽象的设计目标：实现高性能。 首先，我们的工作目标是实现现有的高性能负载平衡算法的不规则应用程序。我们的抽象不能以显著的开销或性能下降为代价。 我们衡量我们的成功achieev-ING高性能比较我们的抽象对现有的硬连线实现的性能的性能。可组合和可编程的接口。重要的是，我们不希望将用户限制在控制更大系统的库接口上。程序员更喜欢采用新的软件组件来适应他们的控制结构，而不是要求他们采用新的控制结构。我们希望允许用户（1）保持对GPU内核边界（内核启动）的控制，（2）能够添加新的负载平衡算法，以及（3）从现有的负载平衡API组成新的负载平衡原语我们衡量我们的工作的可编程性比较paring我们的抽象代码行（代码行）对现有的实现，并通过实现一个新的负载平衡算法在我们现有的API显示组合。可扩展到新的应用程序。 我们的目标是解耦和扩展应用程序特定的负载平衡技术，新的不规则并行域。我们的抽象旨 在 促 进 现 有 的 负 载 平 衡 技 术 的 重 用 ， 为 新 的applications。我们使用SpMV作为使用三种不同负载平衡技术实现的基准应用程序，其中一些技术以前用于实现并行图形分析内核[5，6，10，29]。促进对优化的探索我们抽象的一个关键目标是通过切换用于平衡工作的底层负载平衡算法来促进对给定应用程序的优化探索。我们希望鼓励我们的用户尝试新的负载平衡技术，以发现什么最适合他们的应用程序需求。 我们衡量这一目标的成功，通过优化SpMV的性能响应一个大型语料库的稀疏矩阵在几个不同的负载平衡技术。非目标除了上述设计目标，我们还定义了非目标：针对其他并行架构。虽然我们相信这些经验教训应该适用于其他并行架构，但我们明确针对NVIDIA的CUDA架构和编程模型[ 23 ]。我们抽象的许多组件都利用了CUDA的线程、线程束和映射到物理流多处理器上的块的计算层次结构，超额订阅模型分配了比处理器数量更多的工作，以完全饱和一种GPU负载均衡编程模型PPoPP81底层硬件和CUDA多GPU支持。 这项工作主要关注单个GPU的负载不平衡问题，而不考虑多GPU单节点或多节点系统，尽管这些是未来工作的内部方向。3我们的负载平衡抽象我们的GPU负载平衡抽象背后的关键洞察是工作项到处理单元的映射和工作执行之间的关注点分离。我们将抽象分为三个关键概念（如图1所示），每个概念描述了实现的不同方面：（1）定义工作;（2）定义GPU线程、线程束或块之间的工作负载平衡;（3）定义平衡工作上的每个线程的工作执行和计算。 这种分离允许我们在应用程序开发人员和负载平衡库开发人员之间干净地划分工作，并通过混合不同的负载平衡技术和稀疏不规则算法来促进优化的探索。侧栏1给出了一个实际的例子，说明我们的负载平衡抽象的动机。3.1稀疏数据结构的输入我们从以某种形式的稀疏数据结构表示的输入数据开始。这种数据结构的示例包括但不限于压缩稀疏行（CSR）和坐标（COO）格式。我们抽象的第一阶段的目标是将输入数据格式映射到公共数据框架和词汇表，这是下一阶段的输入。这个词汇表有三个简单的组成部分，它们共同表达输入数据：1. 一个工作原子，一个要被调度到处理器上的单个工作单元（例如，稀疏矩阵的非零元素我们假设所有的工作原子在执行过程中具有相同的成本2. 工作瓦片，表示为一组工作原子的逻辑实体（例如，稀疏矩阵的一行工作瓦片在执行期间可能具有不同的成本正如我们在介绍中强调的那样，工作在逻辑上最并行化在工作瓦片上，但通常在工作原子上最有效地并行化，并且工作瓦片和工作原子之间的映射可能是昂贵和复杂的。3. 图块集，一组工作图块，它们共同构成整个工作问题（例如，稀疏矩阵）。在我们的抽象中，瓦片集中的瓦片必须是独立的（因此可以在多个处理器上并行运行）。稀疏格式和原子/图块/图块集之间的映射由用户定义虽然我们还没有实施考虑开源CUDA CUB库[24]中提供的GPU上的SpMV实现CUB实现并维护Merrill和Garland [20]论文中提出的SpMV算法。基于合并的SpMV（在5.2.1节中详细解释）是一种基于CSR的、完美负载平衡的SpMV，其中每个线程获得均匀的工作份额，工作量由矩阵行的总数和非零的总数相加定义。 在参考文献中，这种高效的最先进的实现在3个文件中占用了1，100行代码（LoC）（或503 LoC的内核代码）（不包括额外234LoC的分段修复步骤所需的第4个文件）。相比之下，该参考实现中的SpMV的实际计算在单个for循环和4 -5LoC内表示！以负载平衡的方式将工作项映射到处理单元所需的LoC与表达所需计算所需的LoC之间的这种差异是我们工作背后的关键动机。此外，CUB实现专门用于SpMV算法，并且需要进行大量重写才能将其应用于其他算法，即使是在同一计算域内。这种精确重写的一个这样的例子是Yang等人， 他们将合并路径负载平衡从SpMV扩展到稀疏矩阵密集矩阵乘法（SpMM）实现[30]。这两个作品中的负载平衡算法是相同的，但应用于不同的计算，这激发了重用的需要。边栏1现有的主要方法的一个实际例子，用于负载平衡稀疏不规则的工作负载。所有这些，我们相信我们这里的映射抽象足够灵活，可以表达文献[12]中各种各样的现有稀疏数据格式，以这种方式，它们适合于我们抽象的下一阶段的负载平衡同样，我们已经在我们的负载平衡库实现中包含了几种常见的稀疏格式（CSR，CSC，COO），这样用户就可以简单地选择和使用它们，而不必实现它们。给定到原子/瓦片/瓦片集的映射，我们接下来可以实现负载平衡算法，该算法可以从计算3.2定义负载平衡通过通过抽象来表达工作负载，该抽象以不同的粒度级别捕获工作（即，图块集、原子和图块），我们可以更轻松地在GPU的可用资源中均匀分布计算给定用户定义的输入图块集和原子和图块的相关联序列以及用户选择的分区算法，我们的负载平衡阶段输出分配给处理器id的原子和图块的顺序（即， 其中将处理原子或瓦片）。PPoPPM. Osama，S.D. Porumbescu，和J.D. 欧文斯82稀疏数据结构迭代器表示负载平衡工作执行工作原子和瓷砖Thread0Thread1用户自定义计算计算（kernel）（内核）图1.负载平衡作为我们抽象的三个关键概念的简单管道：（1）表示为迭代器的稀疏数据结构，（2）将工作划分到线程上的负载平衡算法，以及（3）用户定义的计算消耗平衡的工作并在每个线程上执行。由此产生的对处理器ID的优先级分配对于跨处理元素有效地平衡工作负载是至关重要的，并且通常是特定于问题和处理器的。用户必须指定必要的序列。理想情况下，一个或-acle将采取这些序列，并为每个处理元素选择最佳的顺序。找到这样一个预言是一个开放的问题，因此我们提供了下一个最好的事情：用户能够选择和实验从一组预定义的时间表和实现自己的时间表的能力。一般来说，负载平衡算法设计者必须在调度成本和更好的调度带来的好处之间进行平衡 调度可以像将处理元件分配给具有任意数量的原子的瓦片（例如，在稀疏矩阵中具有任意数量的非零的行）到采用更全面的方法工作的更复杂/昂贵的东西（例如， 考虑在稀疏矩阵中具有变化数量的非零的多个行上的工作）。3.3定义工作执行我们的负载平衡抽象的最后一个组成部分表示不规则并行计算本身。 该阶段输入负载不平衡的工作并对其进行负载平衡;然后该阶段通过对其执行计算来消耗负载平衡的工作。可以表达的计算范围是广泛的，并且仅受负载平衡工作（表示为序列）如何在CUDA内核中使用的限制。 由于框架不承担内核的控制权，因此您可以在CUDA内核中编写的任何内容都可以在我们的框架中工作。例如，程序员可以表达对每个原子或工作的每个瓦片执行的数学运算，或者构建协作算法，该算法不仅消耗分配给每个线程的工作，而且还将结果与相邻线程组合以实现更复杂的算法，例如并行减少或扫描。我们已经在我们的框架中使用这种抽象实现的实际示例（参见第4.3节和第5.3节）包括但不限于稀疏线性代数内核，例如稀疏矩阵和稀疏张量收缩，以及以数据为中心的并行图算法，如单源最短路径（SSSP）和广度优先搜索（BFS）建立在一个邻域遍历内核。我们希望我们库的典型用户只为这个抽象阶段编写他们自己的代码，并使用标准数据结构和负载平衡调度，这些都是我们库的一部分。但是，这些用户也可以实现自定义数据格式和负载平衡计划。4高级别框架实施我们的GPU负载平衡框架使用C++ 17和CUDA实现了第3节中描述的抽象在我们的系统中，程序员使用CUDA/C++开发不规则并行算法和实现新的负载平衡计划。根据我们的可组合API、可扩展性和重用的设计目标，本节和下一节介绍了我们的API的实现细节，以及如何使用它来开发新的应用程序，以促进框架中可用的高性能负载平衡技术的重用。我们还探索了一种基于CUDA的协作组模型的新的负载平衡方法（第5.2此外，我们确定了我们的工作如何可以用来促进探索的优化为一个给定的应用程序，如SpMV。4.1实现稀疏数据结构我们的框架将稀疏数据结构（例如，COO，CSR，CSC）转换为工作原子，工作瓦片和瓦片集（第3.1节），使用简单的C ++迭代器。 C ++迭代器是指向一个元素范围内的某个元素的对象，它可以使用一组运算符来迭代该范围内的元素。例如，counting_iterator是一个迭代器，它表示指向序列值范围的指针[1]。我们的框架要求用户使用C++定义三个重要的迭代器：（1）所有工作原子的迭代器;（2）工作瓦片的迭代器;（3）每个工作瓦片中原子数量的迭代器（我们的库已经支持几种常见的稀疏数据结构。）使用这些迭代器，负载平衡调度可以确定负载平衡的工作并将其分布到负载平衡调度程序线程0线程1atoms iter= 0，1，2，3tile iter = 0，1，2，3atoms/tile =0，1，3，0值= 1，3，6，21362一种GPU负载均衡编程模型PPoPP831//原子和瓦片的简单迭代器。2counting_iteratorint> atoms_iter（0，nnz）;3counting_iteratorint> tile_iter（0，rows）;4//遍历tile i中的原子。5autoatoms_per_tile = make_transform_iterator（6tile_iter，7[tile_iter，row_offsets]8个主机public int findDuplicate（inti）{9return（row_offsets[tile_iter[i +1]]-10row_offsets[tile_iter[i]]）;11}）;清单1.压缩稀疏行（CSR）格式在我们的框架中使用C++17表示。CSR格式使用三个数组描述矩阵：（1）非零值的列索引;（2）行的范围（行偏移量）;以及(3)非零值。由于CSR数据结构不包含指向原子和瓦片（非零和行）索引的数组，在上面的列表中，我们将原子和瓦片迭代器定义为简单的计数迭代器，分别从0到非零的总数（nnz）和从0到矩阵中的总行数（rows）（第2使用transform iterator来表达每个工作瓦片的原子的迭代器，其计算针对每个瓦片id的所提供的函数内的表达式。 对于CSR，这仅仅是当前区块的行偏移量减去下一区块的偏移量（第5-11行）。底层硬件清单1显示了我们的抽象如何将常用的CSR格式表示为框架中的tile集。4.2实现负载平衡计划也许最直接的调度是将每个工作瓦片调度到一个GPU线程上。这种方法在文献和实践中很常见[3，10，21，26，30];尽管这种策略在存在跨瓦片的显著负载不平衡的情况下无效，但我们在这里将其用作示例来说明如何在我们的框架中定义负载平衡输入是最后一个阶段的三个迭代器加上原子和瓦片计数。然后，负载平衡算法开发人员实现tiles（）和atoms（）处理调用，其返回要由当前线程处理的tiles和atoms的C ++范围，有效地创建分配的处理器id和工作负载段之间的映射。清单2显示了线程映射调度的完整示例。 虽然是一个简单的算法，但它可以通过粗粒度并行（每个瓦片的原子数量较少）为平衡良好的工作负载提供高性能，例如将稀疏向量乘以密集向量。此外，我们的抽象不仅限于简单的调度算法，因为5.2节提供了更复杂的负载平衡算法的示例。1class_t {2//构造一个线程映射的调度。3主机设备4schedule_t（atoms_it_t atoms_it，5tiles_it_t tiles_it，6atoms_it_t atoms_per_tile_it，7size_tnum_atoms，size_tnum_tiles）：8m_atoms_it（atoms_it），m_tiles_it（tiles_it），9m_atoms_per_tile_it（atoms_per_tile_it），10m_num_atoms（num_atoms），11m_num_tiles（num_tiles）{}12//要在“this”线程中处理的tiles的范围。13//按网格维度步幅。14hostdeviceautotiles（）{15自动开始= m_tiles_it（blockDim.x * blockIdx.x16+ threadIdx.x）;17autoend = m_tiles_it（m_num_tiles）;18返回范围（开始，结束）19.step（gridDim.x * blockDim.x）;20}21//在“this”线程中处理的原子范围。22主机设备自动原子（conststd：：mysql_list）{24autobegin = m_atoms_per_tile_it[tile];25autoend = m_atoms_per_tile_it[tile +1];26返回范围（开始，结束）.步骤（1）;27}28};29使用schedule_t = thread_mapped_schedule_t;清单2. 一个线程映射的负载平衡算法，表示为C++范围，包含清单1中定义为迭代器的原子和瓦片。 每个图块被映射到线程，其中线程id对应于图块集合中的图块的索引。瓦片内的所有原子由线程顺序处理在处理完一个图块之后，线程被映射到下一个图块，这是通过将索引跨越内核的网格大小来4.3实施工作执行我们的框架被设计为显式地让用户拥有内核启动边界。拥有CUDA内核边界意味着用户负责维护和配置启动参数，并实现用于定义应用程序的CUDA内核。虽然这种设计决策是以方便和简单为代价的，但它在用户可以通过我们的抽象表达什么方面提供了很大的灵活性这一设计决策是由以下原因驱动的（1）用户不需要向现有的工作流/库添加复杂的依赖关系，因此代码维护更简单，可扩展性更强，因为他们不必依赖我们的框架来整合新的CUDA结构和功能。（2）用户可以自由地在内核中表达CUDA允许的任何内容，同时使用我们的负载平衡C ++范围。这允许PPoPPM. Osama，S.D. Porumbescu，和J.D. 欧文斯84因为用户现在可以指定多个负载平衡的工作域，基于范围的for循环，甚至融合多个计算以在单个内核中构建更复杂的算法。(3)高级API可用于构建更简单的高级抽象，这些抽象拥有内核边界，并以灵活性为代价提供更简单的API。作为此阶段的输入，用户使用负载平衡的C++范围来实现其计算。 这可以通过多种方式来实现，但最常见的一种模式是基于嵌套范围的for循环，它在所有分配的tile和atoms范围上循环。 清单3显示了CUDA内核的一个简单示例，它使用清单1和清单2中描述的CSR格式和线程映射负载平衡算法实现SpMV算法。在这个例子中，每个线程中的外部for循环迭代稀疏矩阵（tiles）的指定行，内部循环顺序处理每行中在5.3节中，我们实现并讨论了更复杂的内核和计算。5实现细节5.1支持CUDA的灵活、可组合范围使用我们的API的负载平衡原语和应用程序的可组合性是我们通过使用支持CUDA的C ++范围支持的框架内的一个有意识的设计选择。 我们的框架不拥有内核边界（内核启动），这迫使我们的API集中并包含在内核中。这允许程序员构建和维护他们自己的内核，同时仍然受益于我们框架这在很大程度上是使用设备范围的C ++函数和标记有CUDA关键字device的类来实现的。2我们实现并公开了几种不同类型的专用范围，它们在实现负载平衡调度时特别有用：step_range：一个范围，以step的步骤从开始到结束进行迭代。对于定义需要自定义步进范围或每个线程处理固定数量的工作项（可以使用step定义）的负载平衡计划很有用infinite_range：从开始到无穷大迭代的范围。 用于在持久内核模式下定义负载平衡计划[32]，其中内核持久运行，直到所有工作都被消耗或算法收敛。grid_stride_range：步长范围的一种特殊情况，使用CUDA内核的网格大小从开始到结束迭代步长。还支持块和1//实现负载平衡的SpMV内核。2globalvoidspmv（const size_trows，3const size_t变量，const size_tnnz，4const int* offsets，const int* indices，5const float* values，const float* x，int x（intx）{7//配置负载均衡。8// Input：为CSR格式定义的迭代器。9schedule_tconfig（10atoms_iter，tile_iter，11atoms_per_tile_it，12nnz，rows）;13//使用基于范围的for循环来连接行。14for（autorow：config.tiles（））{int sum =0;16//使用基于范围的for循环连接原子。17for（autonz：config.atoms（row））18sum += values[nz] * x[indices[nz]];19y[row] = sum;20}21}22//启动SpMV内核。23constexpr size_t块=256;24size_tgrid =（rows + blocks -1）/blocks;25spmv网格，块>（行，行，nnz，26偏移、索引、值、x、y）;清单3.稀疏向量矩阵乘法（SpMV）在我们的负载平衡抽象中使用基于范围的嵌套for循环实现。 稀疏矩阵使用基于CSR的格式表示，其中x是密集输入向量，y是密集输出向量（y = Ax）。 第9 - 12行使用清单2中实现的负载平衡调度和清单1中定义的迭代器来构造要处理的负载平衡工作。第14行和第17行显示了每个线程中的for循环，它遍历稀疏矩阵的指定行，并顺序处理每行中的签名原子。第18行显示了对每个工作原子（非零）执行的实际计算，第19行将结果写入密集输出向量y。分别以块或扭曲大小为步长的扭曲步幅变量5.2实现非平凡的负载平衡正如我们在5.1节中所描述的，我们可以将现有的负载平衡技术解耦并表示为一组C++范围。为了说明这种抽象的潜力，我们首先解耦和表达一种最先进的负载平衡算法，称为合并路径[17]，以前用于平衡基于CSR的SpMV和SpMM[20， 30] ， 并 实 现三个额外的负载平衡算法（warp-，block-）。2使用device关键字修饰的方法和组映射），所有这些都可以在我们的库中找到编译器生成设备可调用入口点。 这允许从内核内部调用代码[23]。供程序员使用。我们新的组映射算法是一种基于每个组切片的计划，其中定义···一种GPU负载均衡编程模型PPoPP85作为任意大小的线程的集合（不限于线程束或块大小）。我们的组映射调度是使用CUDA的协作组编程模型[ 18 ]的每线程瓦片、线程束或块调度[5，21]的一般化5.2.1合并路径负载平衡。 在稀疏矩阵的语言中，merge-path假设矩阵中的每个非零和矩阵中的每个新行都是等量的工作，然后将nnz+行的工作均匀地分配给工作线程集。然后，每个线程在CSR矩阵的非零索引和行偏移量内执行2-D二进制搜索，以找到需要处理的行和非零行的起始位置。然后线程从起始位置开始顺序处理行和非零值，直到它们到达分配给它们的工作的末尾[20]。我们通过两个步骤来表达它，在我们的抽象中将此算法实现为负载平衡调度：（1）设置：C ++调度类的初始化步骤计算每个线程的工作单元数，如上所述进行二分搜索，并将每个瓦片和原子的起始位置存储在线程局部变量中。（2）范围：算法的第二步为每个线程构建范围，以作为“完整”瓦片和“部分”瓦片进行处理如果线程的原子范围完全位于一个瓦片内，则它是“完整的如果一个线程由于我们将负载平衡方法（第4.2节及以上）与工作执行（第4.3节）分离，因此我们不仅可以使用此合并路径实现来实现SpMV，还可以实现任何其他算法，其工作可以分为瓦片和原子，例如， 图邻域遍历算法用于实现广度优先搜索[29]。同样重要的是，合并路径调度现在不再局限于基于CSR的稀疏格式。支持其他格式只需要构建必要的稍微复杂一点的迭代器，这些迭代器能够对每个瓦片的原子进行计数（CSR实现使用清单1中的行偏移数组实现的计算）。5.2.2Warp和块级负载平衡。 线程束级或块级负载平衡调度的目标是为每个线程束或块分配相等份额的切片，然后依次处理这些切片。每个图块中的工作原子将由线程束或块中的可用线程并行处理。 每个线程以线程束或块的大小为步长处理一个新的工作原子，直到工作结束。不同处理单元之间的不平衡留给硬件调度器来处理。 此调度程序依赖于CUDA的超额订阅模型，在该模型中，程序员可以启动比GPU在任何给定时间物理调度的更多数量的线程束或块。作为处理当单元完成处理它们的工作时，从超额预订的池中调度新的单元[5，21]。5.2.3组级负载平衡。组级负载平衡概括了经线级和块级调度。与要求组大小是线程束或块的大小不同，如上所述，该方法利用CUDA的协作组（CG）编程模型[18]来允许程序员指定任意大小的动态大小的组。在这些组中，CG模型允许对组的同步行为进行详细控制，我们利用这个强大的工具来实现一个通用的组级负载平衡计划，有效地给我们的翘曲和块级计划以上免费时，组大小等于一个翘曲或一个块。我们的时间表将工作切片分配给一个组，每个组查看其相等份额的瓦片，并计算每个瓦片的原子数，并将其存储在暂存器（CUDA的共享内存）中。然后，该小组执行paral- lel prefix-sum，这是一种广泛使用的并行算法，它输入一个数组并产生一个新数组，其中任何位置的元素都是所有先前元素的总和[4]。我们使用这个前缀和数组有两个目的：（1）前缀和数组的最后一个元素表示一个组必须处理的工作原子的聚合数量，以及（2）前缀和数组中每个和的位置对应于这些原子所属的工作瓦片调度的设置阶段在暂存器存储器中构建每组的前缀和数组，并且调度的范围循环将原子返回到每个线程中进行处理。如果需要，则通过简单的get_tile（atom_id）操作来获得对应的瓦片，该简单的get_tile（atom_id）操作在前缀和数组内执行二进制搜索以找到对应于正被处理的原子的瓦片。此负载平衡计划依赖CG模型ULE具有根据问题的形状和底层硬件体系结构配置组大小（有效地，直接映射到硬件上的软件构造）的独特优点。例如，现在可以通过简单的编译时间常数，或者配置组大小以完美地与问题的结构保持一致，来针对线程束大小不是32个线程的GPU（AMD的GPU架构支持64个线程束大小[2]）。5.3应用空间我们的工作定义（第3.1节），可组合的API（第5.1节）和多个复杂的，高性能的负载平衡计划（第5.2节）一起提供了一个通用的和可扩展的框架，有足够的空间进行特定于应用程序的优化。 在清单3中，我们已经演示了如何使用PPoPPM. Osama，S.D. Porumbescu，和J.D. 欧文斯861//我想...在CUDA内核中。2//遍历所有分配的行。3publicvoidrun（）{4//遍历矩阵B的所有列。5for（autocol：range（size_t（0），B. size））6.stride（size_t（1）{/新建循环return0;8//遍历所有赋值的非零值。9for（autonz：config.atoms（row））sum += values[nz] * B（nz，col）;11//将和输出到Matrix-C。12C（row，col）= sum;13}14}清单4. 清单3中引入的SpMV周围的一个简单循环允许我们表示稍微复杂一些的SpMM负载平衡计算。我们的框架。一个简单而自然的扩展是实现稀疏矩阵乘法（SpMM）。清单4显示了必要的微小更改，它在清单3中现有代码周围的B矩阵列上添加了另一个循环，以实现SpMM。这种实现也可以扩展到支持Gustavson的通用稀疏矩阵-矩阵乘法（SpGEMM），使用两个内核和一个分配阶段;第一个内核将计算用于为输出稀疏矩阵分配内存的输出行的大小，第二个内核将执行乘法-累加。除了稀疏线性代数之外，我们还可以使用我们的框架来解决其他领域的应用。清单5使用我们的组级负载平衡调度实现了图原语Single-SourceShortestPath（SSSP） SSSP在GPU上的性能很大程度上取决于良好的负载平衡[5，29]，但如果程序员从我们的库中选择负载平衡计划，则负载平衡的细节完全可以1//我想...在CUDA内核中。2//遍历所有指定的边进行处理。3public void run（）{4autosource = config.get_tile（edge）;5// G是图数据结构6自动邻居= G.get_neighbor（source，edge）;7autoweight = G.get_edge_weight（edge）;8floatsource_dist = dist[source];9floatneighbor_dist = source_dist + weight;10//检查目标节点是否已被第11章被当成别人的孩子12floatrecover_distance =13atomicMin（（dist[neighbor]），neighbor_dist）;14//将邻居添加到边界。15if（neighbor_dist recover_distance）16out_frontier[neighbor] =true;17}1819//我想...在CUDA内核之外。20//循环直到边界为空。清单5.并行单源最短路径（SSSP）图形原语表示使用我们的负载平衡计划。1002010210.20.10.020.010.0020.001隐藏此外，在一次会议应用领域（例如，稀疏线性代数）在该不同应用领域中容易地可重用。6评价我们的目标是表明，我们的框架，建立在我们的负载平衡-ING抽象，使高性能和更好的可编程性稀疏不规则的问题。我 们 下面的评估使用我们的SpMV 实现作为基准， 与NVIDIA我们考虑（并实现）了几个额外的应用程序进行评估，包括SSSP，BFS和SpMM。 我们发现他们得出了类似的高层次结论。因此，我们的评估重点是SpMV。我们的测试语料库由大约整个SuiteSparse Matrix Collection [11]组成，具有广泛的稀疏矩阵1101001，00010，0001，000，000 1，000，00010,000,000一亿非零数图2. SpMV运行时比较：我们的合并路径SpMV实现与所有SuiteSparse数据集的CUB。 我们的运行时几乎完美地匹配CUB的所有数据集。 CUB更快的数据集数量较少是由于CUB用于单列稀疏矩阵的简单启发式（即，稀疏向量）。来自许多不同的高性能计算领域。我们在基于Ubuntu 20.04LTS的工作站上运行所有实验，该工作站具有NVIDIATesla V100 GPU和CUDA 11.7。6.1性能开销我们的首要目标是确保我们的抽象中的元素不会给现有的负载平衡技术和使用它们开发的算法增加任何额外的性能为了验证这一点，我们内核幼崽合并路径时间（ms）一种GPU负载均衡编程模型PPoPP87××∼××将我们使用合并路径调度的SpMV实现的运行时性能与NVIDIA的CUB库[24]（也用于Merrill和Garland的合并路径SpMV论文[20]）在Suite-Sparse集合上提供的实现进行比较。 如前所述，与我们的设计相反，CUB包含合并路径调度算法的硬连线实现，并且不从实际的SpMV计算中分离工作负载平衡。CUB图2绘制了非零的数量（即， 总工作量）与我们工作的运行时间与CUB当使用我们的抽象时，我们的实现具有最小的性能开销：2.5%的几何减速，CUB，92%的数据集至少达到CUB性能的90%。图2显示了我们的实现几乎完美地匹配所有数据集的CUB，除了一些非零值少于100，000的数据集在进一步的研究中，我们发现CUB使用一个简单的启发式来启动一个线程映射的SpMV内核，其中给定输入矩阵的列数等于1（即，稀疏向量）。与我们更通用的实现不同，CUB的简单（但专用）线程映射SpMV内核对于完美平衡的工作负载（如SpVV计算）没有负载平衡开销。6.2改进的性能响应我们还将我们的工作与NVIDIA 图3显示了我们的SpMV实现的性能响应，分别使用我们的每种调度算法与cuSparse的最新实现。 在我们的任何实现之间切换需要很少的代码更改;在合并路径和线程映射的情况下，我们只需要更新一个C ++枚举（标识符）来选择所需的负载平衡调度。然后，我们将我们的调度算法合并到一个SpMV的实现中（图4），证明了cuSparse的显著性能改进。 这是primar-ily可能的，因为我们能够快速实验不同的启发式计划与各种可用的负载平衡计划。在这里，我们使用merge-path，除非行数或列数小于阈值α，并且给定矩阵的非零值小于阈值 β （ 对 于 SuiteSparse ， 我 们 选 择 α = 500 ， β =10000）。在这种情况下，我们使用线程映射或组映射负载平衡，而不是合并路径。 我们的系统显示了39的峰值性能加速比和2.7的几何性能加速比，cuSparse。我们的框架不仅允许程序员表达计算高效且简单（即，而不用担心负载平衡算法），而且速度很快负载均衡算法NVIDIA/CUB我们的工作合并路径50336线程映射2221组映射N/A30扭曲映射N/A30（免费）块映射N/A30（免费）表1. NV