WASP: Exploiting GPU Pipeline Parallelism with Hardware-Accelerated Automatic Warp Specialization

作者/机构: Neal C. Crago, Sana Damani, Karthikeyan Sankaralingam, Stephen W. Keckler (均来自 NVIDIA)

A1 主要贡献

本文介绍了一种名为WASP的架构和编译器，旨在解决现代GPU中warp专业化（warp specialization）技术的三大局限性，以提升对内存延迟和带宽敏感的应用性能。

核心问题: 尽管GPU拥有高计算吞吐量和内存带宽，但许多应用（如机器学习、高性能计算等）因无法有效重叠内存访问与计算，导致GPU资源利用率低下，性能受限。现有的warp专业化技术虽能通过流水线并行来缓解此问题，但存在以下局限：
1. 实现复杂且手动：warp专业化是一种复杂的程序转换，需要程序员手动完成，这超出了许多开发者和应用的能力范围。
2. 仅支持粗粒度内存模式：现有技术主要支持全局内存到共享内存（SMEM）之间的大块数据传输，对细粒度的内存访问模式（如流式和收集操作）支持不佳。
3. 硬件不可知：GPU硬件对程序员实现的流水线并行结构一无所知，无法利用这些信息在运行时做出更优的决策。

研究目标与创新点: 为了克服上述局限，本文提出了WASP，一个集成了硬件和编译器支持的warp专业化解决方案。其核心贡献如下：
* 识别现有局限：明确指出现代GPU在warp专业化方面的不足之处。
* 显式流水线阶段命名与流水线感知硬件：
* 引入了显式的warp到流水线阶段的命名机制，使GPU硬件能够感知流水线结构。
* 基于此，设计了新颖的流水线感知的warp映射、寄存器分配和调度方案，以利用warp的异构性，实现更高效的资源利用和执行重叠。下图展示了warp专业化如何将串行的内存-计算阶段（a）转化为可重叠的流水线（b）。

A3 背景知识与关键观察

A. GPU架构概述

现代GPU的并行结构。现代GPU是高度并行的处理器，包含多个流式多处理器（Streaming Multiprocessors, SM），这些SM连接到共享的片上L2缓存和片外DRAM。图2展示了一个SM的组织结构，它由一组称为处理块（processing blocks）的处理器构成，这些处理块共享一个L1缓存和共享内存暂存器（SMEM）【22, Nvidia a100 tensor core gpu architecture, 2020】。一个处理块作为单指令多线程（SIMT）处理器运行，即一条指令被获取并在一个称为warp的线程向量上执行。处理块拥有每线程和warp集体的功能单元（如TensorCore）、一个大的寄存器文件以及用于管理一组warp上下文执行的warp调度器。warp调度器通过交错执行处理块上的多个warp来隐藏流水线停顿。Nvidia的GPU通过使用CUDA模型和接口【25, Cuda c++ programming guide, 2023】开发内核进行编程。每个内核被组织为线程的层次结构，每个线程执行整个程序的一部分工作。CUDA模型中的线程块是一个基本单元，由一个或多个warp组成，这些warp在SM上并发且协作地执行。

流水线并行是解决GPU利用率不足的方案。如今，跨越多种计算领域的应用利用了GPU提供的高计算和内存吞吐量。然而，尽管存在丰富的并行性，一些应用对长内存延迟敏感，或在重叠内存访问和计算方面存在困难，导致GPU利用率不足【16, S.-Y. Lee, A. Arunkumar, and C.-J. Wu, “Cawa: Coordinated warp scheduling and cache prioritization for critical warp acceleration of gpgpu workloads,” in Proceedings of the 42nd Annual International Symposium on Computer Architecture, ser. ISCA ’15. New York, NY, USA: Association for Computing Machinery, 2015, p. 515–527.】【17, A. Li, B. Zheng, G. Pekhimenko, and F. Long, “Automatic horizontal fusion for gpu kernels,” in 2022 IEEE/ACM International Symposium on Code Generation and Optimization (CGO), 2022, pp. 14–27.】【36, T. G. Rogers, M. O’Connor, and T. M. Aamodt, “Cache-conscious wavefront scheduling,” in 2012 45th Annual IEEE/ACM International Symposium on Microarchitecture, 2012, pp. 72–83.】【39, A. Sethia, D. A. Jamshidi, and S. Mahlke, “Mascar: Speeding up gpu warps by reducing memory pitstops,” in 2015 IEEE 21st International Symposium on High Performance Computer Architecture (HPCA), 2015, pp. 174–185.】。解决这个问题的一个特别流行的方案是以warp专业化的形式利用流水线并行。

B. Warp专业化

Warp专业化的基本原理。GPU流水线并行的warp专业化方法首次在CudaDMA【1, M. Bauer, H. Cook, and B. Khailany, “Cudadma: Optimizing gpu memory bandwidth via warp specialization,” in Proceedings of 2011 International Conference for High Performance Computing, Networking, Storage and Analysis, ser. SC ’11. New York, NY, USA: Association for Computing Machinery, 2011.】中被描述，并被广泛用于CUTLASS【12, A. Kerr, D. Merrill, J. Demouth, and J. Tran. (2017) Cutlass: Fast linear algebra in cuda c++. [Online]. Available: https://devblogs.nvidia. com/cutlass-linear-algebra-cuda/】，这是一个用于机器学习的、流行的、最先进的GEMM库。在warp专业化中，CUDA中的单个线程块被编程为在SM上实现一个流水线，使用warp作为流水线阶段。这些warp是“专业化的”，并执行整个CUDA内核的一个独特部分。在当前的GPU上，warp（流水线阶段）之间的通信是通过SM级的同步和数据存储来促进的。与线程内软件流水线等替代技术相比，warp专业化是一种解耦方法，允许在warp级别进行动态调度和细粒度交错。

Warp专业化实例分析。图1a和1b展示了一个warp专业化的实例。图1a描绘了一个线程块在一个常见GPU模式下的执行：将内存块从全局内存传输到SMEM共享内存，以用于软件管理的缓冲区。在这种情况下，一个线程块的warp协同工作，同步缓冲区的状态并在两级内存层次结构之间传输数据。首先，线程块中的warp到达一个屏障，表示SMEM中的缓冲区已准备好被写入。接下来，warp继续发出全局内存加载（LDG）指令。在长延迟的LDG指令完成后，使用共享存储（STS）指令将缓冲区写入SMEM。使用另一个线程块屏障来确保内存块的所有数据已成功写入SMEM，并准备好使用。然后，warp可以自由使用共享加载（LDS）指令从缓冲区读取数据并执行计算。最后，使用一个屏障来表示所有warp都已完成对该缓冲区的使用，从而允许数据被修改或替换。在这种常见模式下，计算阶段必须等到内存块传输阶段成功完成后才能开始，反之亦然。因此，内存访问与计算没有重叠，阶段的严格双峰性质确保了全局内存带宽和计算功能单元（例如TensorCore，浮点单元）不会同时处于活动状态。图1b描绘了如何使用warp专业化来创建一个两阶段流水线。原始的warp被分成两个新的warp，一个用于内存访问，一个用于计算。缓冲区和屏障的数量扩展到两个（A和B），以实现双缓冲。与CudaDMA类似，这里使用了到达/等待屏障【1, M. Bauer, H. Cook, and B. Khailany, “Cudadma: Optimizing gpu memory bandwidth via warp specialization,” in Proceedings of 2011 International Conference for High Performance Computing, Networking, Storage and Analysis, ser. SC ’11. New York, NY, USA: Association for Computing Machinery, 2011.】。在一个到达屏障上，执行的warp记录它已到达该屏障，但继续执行。回到图1b的例子中，一旦阶段0完成填充缓冲区A，它就到达该缓冲区的FilledA屏障，向阶段0发出信号，表示数据块已准备好使用。然后阶段0继续执行，首先等待缓冲区B变空并准备好被填充（使用EmptyB屏障），然后发出相应的LDG指令。因此，时间步N+1的内存块传输阶段与时间步N的计算阶段重叠，warp的执行在SM上交错进行。本质上，通过这种流水线并行提取了更多的内存级并行性，降低了整体的内存敏感性。

C. 现有GPU的局限性

现有技术的局限性。Warp专业化是一种强大的技术，今天的GPU通过硬件支持快速的到达/等待屏障和用于全局内存与SMEM之间内存传输的硬件卸载单元（Hopper集成了一个新的Tensor Memory Accelerator，或TMA）【22, Nvidia a100 tensor core gpu architecture, 2020】【23, Nvidia h100 tensor core gpu architecture, 2022】来采用它。然而，当前GPU硬件和软件对warp专业化流水线的支持由于三个关键问题而相当有限。

硬件对流水线并行性的无知。首先，当前的GPU硬件对软件中表达的流水线并行性是不可知的。特别是GPU SM被设计为作为SIMT执行模型的一部分来执行数据并行的线程块，使用的warp（在很大程度上）执行相同的程序。因此，SM硬件假定资源利用是均匀的。另一方面，具有warp专业化流水线的线程块本质上是异构的，不同的warp需要不同的资源。由warp实现的每个流水线阶段都需要一个独特的程序、计算资源（例如功能单元）、寄存器分配，并且在某些情况下更倾向于不同的warp调度优先级【1, M. Bauer, H. Cook, and B. Khailany, “Cudadma: Optimizing gpu memory bandwidth via warp specialization,” in Proceedings of 2011 International Conference for High Performance Computing, Networking, Storage and Analysis, ser. SC ’11. New York, NY, USA: Association for Computing Machinery, 2011.】【2, M. Bauer, S. Treichler, and A. Aiken, “Singe: Leveraging warp specialization for high performance on gpus,” in Proceedings of the 19th ACM SIGPLAN symposium on Principles and practice of parallel programming, 2014, pp. 119–130.】【12, A. Kerr, D. Merrill, J. Demouth, and J. Tran. (2017) Cutlass: Fast linear algebra in cuda c++. [Online]. Available: https://devblogs.nvidia. com/cutlass-linear-algebra-cuda/】【19, N. Maruyama and T. Aoki, “Optimizing stencil computations for nvidia kepler gpus,” 2014.】【48, H. Wei, E. Liu, Y. Zhao, and H. Yu, “Efficient non-fused winograd on gpus,” in Advances in Computer Graphics: 37th Computer Graphics International Conference, CGI 2020, Geneva, Switzerland, October 20–23, 2020, Proceedings. Berlin, Heidelberg: Springer-Verlag, 2020, p. 411–418.】。如果GPU硬件能够感知所表达的流水线并行性，就可以设计新的硬件来利用warp专业化流水线的异构性，从而带来更好的性能。

缺乏对细粒度内存模式的支持。其次，最先进的GPU硬件支持在全局内存和SMEM之间进行粗粒度内存块传输的流水线。虽然这种支持对于一些重要的用例（如CUTLASS中使用的GEMM数据流）是足够的，但我们发现，如果提供对细粒度内存访问模式的支持，更大范围的应用可以从warp专业化中受益。例如，图3展示了Pointnet++，它在点集上执行深度学习，并包含一次性使用的收集（gather）和流式（streaming）内存操作【33, C. R. Qi, L. Yi, H. Su, and L. J. Guibas, “Pointnet++: Deep hierarchical feature learning on point sets in a metric space,” arXiv preprint arXiv:1706.02413, 2017.】。尽管存在丰富的并行性，Pointnet++在Ampere上由于无法重叠交替的计算和内存访问阶段，难以维持高的TensorCore和L2带宽利用率（图3a）。WASP（本文）能够利用重叠这些阶段的机会，实现更好的L2带宽利用率（图3b）。

手动优化的复杂性。最后，当前warp专业化的一大限制是它是一种复杂的优化，由程序员在CUDA源代码级别手动实现。程序员负责将原始内核划分为流水线，将线程分配给流水线阶段，并管理阶段之间的数据移动和同步。手动优化的复杂性限制了可以利用warp专业化的应用数量。然而，这种优化结构良好，适合编译器自动化，可以使用先前研究中的类似技术【9, N. C. Crago and S. J. Patel, “Outrider: Efficient memory latency tolerance with decoupled strands,” SIGARCH Comput. Archit. News, vol. 39, no. 3, p. 117–128, jun 2011.】【21, Q. M. Nguyen and D. Sanchez, “Phloem: Automatic acceleration of irregular applications with fine-grain pipeline parallelism,” in 2023 IEEE International Symposium on High-Performance Computer Architecture (HPCA), 2023, pp. 1262–1274.】。在本文中，我们介绍了WASP，这是一种新的GPU硬件和编译器支持，旨在克服现有的warp专业化限制，并提高各种应用的性能。WASP使得线程块的warp专业化流水线结构能够传达给SM并更高效地执行。WASP使用新的寄存器文件队列和硬件卸载单元扩展了对细粒度内存访问模式的硬件支持。最后，WASP由一个编译器启用，该编译器减轻了程序员的负担，并自动对现有的CUDA代码进行warp专业化。

A2 方法细节

III. WASP架构

WASP架构概览。本节详细阐述WASP架构如何改进当前GPU中发现的两个局限性（第三个局限性在第四节中详述）。在架构层面，WASP表现为对GPU的SM和处理块的增强。图4描绘了WASP SM架构，它更新了SM以更好地利用通过warp专业化发现的可用流水线并行性。首先，更新了线程块规范，以显式地将warp分配给流水线阶段（第III-A节），这使得SM能够利用warp专业化流水线的异构性。WASP实现了新的硬件，以利用线程块warp映射器和warp调度器中每个流水线阶段的行为，这些是正确平衡warp执行的关键机制。warp映射器包括对新的每阶段寄存器分配和流水线感知的warp到处理块映射的支持（第III-B节）。WASP还包括一个新的warp调度器，以强调执行重叠（第III-D节）。其次，WASP增加了对细粒度流式和收集内存访问模式的支持，以更好地支持可能从warp专业化中受益的广泛CUDA应用。架构可见的队列及相关的ISA扩展被添加到每个处理块的共享寄存器文件中（第III-C节）。这些寄存器文件队列（RFQs）提供低延迟数据访问，并自然支持当前GPU上不支持的一次性使用流式和收集模式。最后，我们增强了SM中的TMA单元以支持流式和收集内存访问模式，通过减少在处理块上执行的warp地址指令流进一步提高效率（第III-E节）。

A. 显式的Warp到流水线阶段命名

通过命名机制使硬件感知流水线。显式命名在硬件中将warp分配给流水线阶段，允许程序员或编译器指定warp专业化流水线的配置。其核心是，线程块中的每个warp都被分配一个流水线阶段ID。显式命名对性能至关重要，因为它启用了WASP采用的其他流水线感知特性，使SM硬件能够感知流水线并做出更好的映射、寄存器分配、调度和数据移动决策。为了实现显式命名，WASP扩展了线程块规范，增加了一个额外的维度来表示warp专业化流水线的深度。表I展示了这一新规范，其中还包括了由WASP编译器在第四节中生成的其他元信息。这个新维度在CUDA内核调用时由程序员或代码生成框架显式声明。新的线程块维度规范形式为：{dim.x, dim.y, dim.z, num pipeline stages}。在SM上执行时，一个线程会被分配一个pipe_stageId。这个新的硬件状态存在于SM的特殊寄存器中，可以被线程块的线程查询以确定其分配的流水线阶段，类似于查询其他线程块参数如threadId.x。

表I：WASP线程块规范

B. 流水线感知的Warp映射和寄存器分配

流水线感知的Warp映射算法。当一个线程块被调度到SM上执行时，每个warp必须被映射（即分配）到一个处理块上执行。warp到处理块的映射由SM内的一个硬件映射器控制，该映射器首先从全局GPU调度器接收一个线程块规范，然后以轮询方式将线程块的warp分配到各个处理块上执行。WASP的warp映射器利用线程块规范中的新信息来改进warp到处理块的映射算法和warp寄存器分配。warp映射算法专注于通过最小化将相似流水线阶段映射到同一处理块来平衡计算资源的使用。group_pipeline算法将属于同一warp专业化流水线切片的所有warp一起映射到同一个处理块上。图5展示了WASP映射算法相对于原始轮询映射算法的潜在价值。WASP线程块由一个两阶段的warp专业化流水线组成，每个阶段有四个warp。想象一下S0是内存访问阶段，S1是计算阶段的场景。轮询算法一次一个地交替将warp映射到处理块，导致执行不平衡，S0阶段被映射到处理块0和2，而S1阶段被映射到处理块1和3。group_pipeline算法通过将每个流水线切片的warp在同一个处理块上执行，提供了更好的平衡。

流水线感知的寄存器分配。在warp映射过程中，通过划分处理块共享寄存器文件的一个连续段来进行warp的寄存器分配。在warp专业化中，不同的流水线阶段运行不同的程序，因此使用不同数量的寄存器。然而，在当前的GPU中，warp专业化流水线被分配一个统一的warp寄存器分配值，该值是所有流水线阶段中的最大使用量。在我们评估的内核中，一些流水线阶段每线程需要的寄存器数量几乎与原始程序相同，这可能导致一个N阶段流水线的线程块寄存器使用量增加近N倍。WASP使用表I中线程块规范里的每流水线阶段寄存器值来分配大小独特的寄存器段，这大大节省了寄存器文件空间，是对现有技术的改进。图7展示了一个例子，比较了原始程序、现有GPU上的warp专业化程序以及WASP提供的寄存器节省的每线程寄存器分配。更多关于寄存器分配节省的数据可以在第V-D节找到。

C. 寄存器文件队列

使用SMEM作为队列的开销。虽然当前GPU的warp专业化硬件支持全局内存和SMEM之间的粗粒度内存块数据移动，但一些对内存敏感的CUDA应用大量使用一次性数据。这种用例表现为细粒度的流式和收集内存访问模式，WASP通过在流水线阶段之间使用队列来支持。在WASP的队列设计中，我们研究了数据队列的数据存储和同步应该在哪里进行。一种选择是简单地使用与粗粒度内存块相同的软件机制：SMEM双缓冲和硬件同步单元。然而，管理进出SMEM的数据会导致执行开销。每个原始的全局加载指令在生产者阶段被转换为一个加载全局-存储共享（LDGSTS）指令来存入缓冲区，这需要额外的地址生成指令和控制指令来跟踪缓冲区是否已满。消费者阶段的代码类似，只是需要一个共享内存加载（LDS）指令来读取队列。当缓冲区满或空时，生产者和消费者必须相互通知，使用相应的到达/等待屏障。虽然对于粗粒度内存传输，这种开销可以通过数据重用得到摊销，但在一次性使用的流式和收集内存访问模式中没有摊销。经过评估，我们发现这些开销很高并影响性能（第V-C节）。

硬件队列（RFQ）的设计。考虑到这些因素，我们在WASP中设计了硬件队列，将它们映射为现有寄存器文件空间中的循环缓冲区。图6描述了寄存器文件队列（RFQs）如何集成到SM处理块中。与每warp寄存器分配类似，连接流水线阶段的命名队列在线程块规范中被明确定义，并在warp映射时在物理寄存器文件中分配。例如，TB0 W0 S0S1指定线程块0有一个连接原始warp 0的阶段0和阶段1的命名队列。每个命名队列都由一个附带的硬件表支持，以维护队列的状态，包括头指针、尾指针、队列分配起始索引和队列分配结束索引。

RFQ的ISA扩展和使用。我们扩展了GPU的ISA，以直接使用这些名称作为warp指令的操作数，类似于寄存器。例如，当TB0 W0 S0执行以下warp指令时：LDG.E.DECOUPLED Q0, [RX]，一个解耦的全局内存加载被发出，使用TB0 W0 QS0S1队列将数据发送到TB0 W0 S1。类似地，TB0 W0 S1可以执行以下warp指令来从TB0 W0 S1队列中读取数据：MOV RY, Q0。位于warp指令调度器处的一个RFQ记分板支持队列中是否有数据（即读取时）或者队列中是否有空间容纳更多数据（即写入时）。这些记分板位is_empty、is_full在成功读取或写入RFQ时更新。

D. SM中的流水线感知Warp调度

引入新的调度策略。我们还增强了处理块内现有的硬件warp调度器，增加了一个新的流水线感知的优先级策略，以促进执行重叠。现有的硬件warp调度器存储有关每个warp状态的信息（例如，等待长/短记分板），用于确定其调度优先级。在WASP中，每个硬件warp调度器都被增强，以存储每个warp的流水线阶段ID和传入数据队列状态（例如，队列不为空，队列已满），这些作为调度决策算法的额外输入。利用这些新输入，我们实现了利用流水线阶段元信息进行优先级排序的新颖warp调度策略。

两种调度策略的结合。第一个策略使用流水线阶段ID输入来根据warp在流水线中的阶段来确定其优先级。我们的命名约定将pipe_stageId_0定义为流水线中的第一个阶段，pipe_stageId_N-1定义为一个N阶段流水线中的最后一个阶段。给定一组在SM处理块上执行的warp，WASP提供了根据其pipe_stageId值的升序来优先处理早期阶段的能力。我们发现，优先处理流水线的早期阶段通常是个好主意，因为这些阶段往往是内存访问阶段，我们希望重叠其长内存延迟。第二个策略建立在第一个策略之上，并使用RFQ状态输入来确定调度哪个warp的优先级。对于每个warp，我们保留以下位：is_empty、is_full。我们的硬件warp调度器使用is_empty、is_full RFQ状态位来在数据准备好时优先处理消费者warp。这两种策略在我们增强的硬件warp调度器中结合使用。在第V-E节中，我们发现，优先处理具有满队列的warp，其次是具有就绪缓冲区或非空队列的warp，再其次是流水线早期阶段的warp，这种策略表现最好。

E. WASP-TMA：加速地址生成

WASP-TMA支持多种数据移动模式。TMA是一个用于全局内存和SMEM之间粗粒度内存块传输的卸载加速器。TMA使用新的指令来指定要传输的全局内存块，使用元信息如基地址、步幅、偏移量和维度【23, Nvidia h100 tensor core gpu architecture, 2022】【24, Cuda binary utilities, 2023】。将内存访问指令流卸载到TMA有两个好处。首先，硬件加速减少了warp指令发射槽和寄存器，为SM处理块中的计算（如TensorCore操作）留出更多资源。其次，利用专用硬件更高效地生成访问，降低了能耗。WASP通过使用类似于先前ISA扩展提案【8, N. C. Crago, M. Stephenson, and S. W. Keckler, “Exposing memory access patterns to improve instruction and memory efficiency in gpus,” ACM Trans. Archit. Code Optim., vol. 15, no. 4, oct 2018.】的硬件，增强了TMA处理线程粒度细粒度内存访问的能力。图8展示了WASP支持的模式，包括全局内存和SMEM之间的内存块传输（图8a）、全局内存和RFQ之间的流式数据传输（图8b）以及全局内存和SMEM/RFQ之间的收集操作（图8c）。

WASP-TMA的指令和操作机制。一个类似TMA的global-SMEM指令用于在全局内存和SMEM之间移动粗粒度数据，使用到达-等待屏障进行同步。增加了一个新的global-RFQ配置，以命名RFQ为目标而不是SMEM。细粒度同步通过RFQ的状态表来完成。典型的解耦LDG指令写入寄存器文件队列时会获取队列中的一个条目。WASP-TMA的global-RFQ指令会获取多个条目，并在它们可用之前延迟发射。收集访问模式 (C[i] = B[A[i]]) 专注于最小化数据移动并有效地将两个操作融合在一起。WASP-TMA的 (gather-SMEM) 和 (gather-RFQ) 指令首先生成一个到全局内存中索引数组的收集内存访问流。当收集到的索引到达全局内存时，它们被WASP-TMA消耗并在第二阶段进行处理。第二阶段将收集到的索引与基地址结合，生成一个内存请求流，目标是SMEM (gather-SMEM) 或RFQ (gather-RFQ)。传入的索引保存在一个具有两个条目的乒乓缓冲中，一个用于TMA当前正在处理的索引集，另一个用于在同一周期内接收新的索引集。通过不将收集到的索引写回SMEM，WASP-TMA消除了否则会需要的额外RFQ和SMEM流量。

IV. WASP编译器设计

编译器自动化的目标与方法。Warp专业化是一种非平凡的转换，目前由程序员手动实现或使用像CUTLASS这样的优化库来处理GEMM【12, A. Kerr, D. Merrill, J. Demouth, and J. Tran. (2017) Cutlass: Fast linear algebra in cuda c++. [Online]. Available: https://devblogs.nvidia. com/cutlass-linear-algebra-cuda/】。WASP通过一个由程序员指导的编译器转换来消除这一限制，该转换自动生成在支持WASP的GPU上运行的warp专业化CUDA内核。我们的编译器执行二进制重编译，使用Nvidia的nvdisasm【24, Cuda binary utilities, 2023】生成的SASS汇编。我们还做出简化假设，以积极利用GPU的弱一致性模型：线程块同步仅用于SMEM传输，且不使用内存栅栏。程序员应在启用warp专业化转换之前确保其内核满足这两个假设。虽然我们的实现是一个源码到源码的编译器，但其他针对解耦访问执行处理器的工作表明，这类思想可以嵌入到主流的GPU编译基础设施中【21, Q. M. Nguyen and D. Sanchez, “Phloem: Automatic acceleration of irregular applications with fine-grain pipeline parallelism,” in 2023 IEEE International Symposium on High-Performance Computer Architecture (HPCA), 2023, pp. 1262–1274.】【43, N. Topham, A. Rawsthorne, C. McLean, M. Mewissen, and P. Bird, “Compiling and optimizing for decoupled architectures,” in Supercomputing ’95:Proceedings of the 1995 ACM/IEEE Conference on Supercomputing, 1995, pp. 40–40.】。编译器的目标输出是在我们的WASP GPU性能模型（第V-A节）上运行的SASS程序。我们首先使用程序依赖图（PDG）提取潜在的流水线阶段（第IV-A节），然后将warp专业化流水线最终确定为一个SASS程序（第IV-B节）。每种内存访问模式如何转换的详细示例见第IV-C节。

A. 流水线阶段提取

基于PDG和回溯切片（backslice）的阶段提取。与过去的工作类似，WASP的分区过程专注于通过在全局内存加载-使用边界处将原始程序分解为流水线阶段，来重叠长延迟的内存访问与计算【9, N. C. Crago and S. J. Patel, “Outrider: Efficient memory latency tolerance with decoupled strands,” SIGARCH Comput. Archit. News, vol. 39, no. 3, p. 117–128, jun 2011.】【10, T. J. Ham, J. L. Aragon, and M. Martonosi, “Desc: Decoupled supply-´ compute communication management for heterogeneous architectures,” ser. MICRO-48. New York, NY, USA: Association for Computing Machinery, 2015, p. 191–203.】【20, Q. M. Nguyen and D. Sanchez, “Pipette: Improving core utilization on irregular applications through intra-core pipeline parallelism,” in 2020 53rd Annual IEEE/ACM International Symposium on Microarchitecture (MICRO), 2020, pp. 596–608.】【40, J. E. Smith, “Decoupled access/execute computer architectures,” in Proceedings of the 9th Annual Symposium on Computer Architecture, ser. ISCA ’82. Washington, DC, USA: IEEE Computer Society Press, 1982, p. 112–119.】【47, Z. Wang and T. Nowatzki, “Stream-based memory access specialization for general purpose processors,” in Proceedings of the 46th International Symposium on Computer Architecture, ser. ISCA ’19. New York, NY, USA: Association for Computing Machinery, 2019, p. 736–749.】。首先，我们构建原始SASS程序的程序依赖图（PDG）。我们的PDG维护程序指令之间的控制流边和数据依赖关系。接下来，我们使用类似于OUTRIDER【9, N. C. Crago and S. J. Patel, “Outrider: Efficient memory latency tolerance with decoupled strands,” SIGARCH Comput. Archit. News, vol. 39, no. 3, p. 117–128, jun 2011.】的流水线阶段提取方案将PDG划分为warp流水线阶段。我们识别内核中所有的全局内存加载指令（LDG），并使用指令的回溯切片来确定哪些指令适合进行流水线阶段提取。回溯切片包括LDG直接依赖的原始程序中的所有指令，以及LDG间接依赖的、对程序控制流有贡献的指令。包含SMEM加载指令（LDS）的回溯切片表明LDG可能与SMEM存储（STS）指令存在无法追踪的内存依赖关系，因此不适合进行流水线阶段提取。类似地，包含LDG自身依赖循环的回溯切片也被排除。我们发现在实践中，这两种情况在GPU代码中都不常见。每个符合条件的LDG指令在加载-使用边界处被拆分为两个新指令：用于地址的LDG_PRODUCER和用于加载结果的LDG_CONSUMER。作为一种特殊情况，我们识别仅连接到SMEM存储的LDG指令，并将其合并为单个加载全局-存储共享（LDGSTS）指令。这些LDGSTS指令需要稍后描述的特定同步。

两阶段提取过程。在第一个提取阶段，为每个LDG_PRODUCER和LDGSTS创建一个初始的新阶段，使用所需指令和基本块的集合。通过遍历PDG，识别每个LDG_PRODUCER和LDGSTS的地址生成指令的回溯切片，并将其添加到集合中。回溯切片的深度优先搜索在依赖链的端点或上游的LDG_PRODUCER指令处终止。图9描绘了LDG B的地址生成回溯切片在LDG A处终止，并包含两条指令（①）。第二个提取阶段通过添加最小的基本块和相关的控制流指令来完成新阶段的PDG子图。对于阶段集合中的每条指令，将流入该指令所在基本块的父基本块添加到搜索列表中。对于列表中的每个未访问过的基本块，将该块标记为已访问并检查其是否包含分支指令。如果存在分支，则为该分支生成一个指令回溯切片并添加到集合中。搜索过程持续进行，直到集合中的所有指令都已检查完毕，表示该集合代表了执行该阶段所需的所有指令。使用这些指令生成一个PDG子图，保持它们在原始程序中的顺序。图9显示，①中的地址生成指令被分配到基本块B（BB_B），其父基本块BB_A和BB_B被添加到搜索列表中。BB_A没有分支，而BB_B有一个分支需要回溯。②中的指令被添加到阶段的集合中，过程继续。基本块BB_A和BB_B已经被访问过，所以第二个提取阶段完成。

WASP-TMA卸载优化。在第二阶段结束时会执行一个优化遍，以检测LDG或LDGSTS指令是否包含在可以卸载到WASP-TMA的循环中。如果循环和地址生成模式合适，循环中识别出的控制流和地址生成指令将被替换为WASP-TMA配置指令。

B. 流水线最终化与配置

流水线的生成与配置。当所有符合条件的LDG_PRODUCER指令都处理完毕并且所有阶段都已提取后，会生成一个代表warp专业化流水线的最终程序。首先，有一组原始程序指令仍未被分配。这组指令代表属于最终流水线阶段的计算，包括所有共享内存和全局存储指令。这些指令被分配到一个新阶段，并保持程序顺序。接下来，根据SM是否能支持提取出的阶段数量，可以选择性地合并阶段。我们发现，为CUDA内核优化的SASS可能有数十到数百个静态LDG指令被提取到阶段中。考虑到我们研究的CUDA应用的原始线程块维度，这样的warp专业化内核将无法容纳在SM上。为了减少大量提取的阶段，我们使用先前工作中的一种方案，该方案合并具有相似内存间接级别的内存访问阶段【9, N. C. Crago and S. J. Patel, “Outrider: Efficient memory latency tolerance with decoupled strands,” SIGARCH Comput. Archit. News, vol. 39, no. 3, p. 117–128, jun 2011.】。我们发现这种方案在适应SM资源的同时表现良好。

阶段间通信与同步的最终化。在这个可选的合并步骤完成后，阶段根据其自然顺序被赋予显式的流水线阶段名称。对于一个大小为N的流水线，阶段零总是最早的阶段，阶段N-1是最后的阶段。然后用流水线阶段的数量更新线程块规范。表I描绘了在这个最终过程中填充的线程块规范表。阶段间的通信和同步也被最终确定。我们对带有LDG和LDGSTS指令的阶段处理方式不同。对于LDG_PRODUCER和相应的LDG_CONSUMER指令，使用RFQ，并在线程块规范中添加一个连接两个阶段的、具有固定大小的新命名队列。LDG_PRODUCER和相应的LDG_CONSUMER指令被更新以使用该命名队列，并且如果只有一个依赖指令，LDG_CONSUMER指令可以选择性地合并到该依赖指令中。

SMEM缓冲与同步的自动化。带有LDGSTS指令的阶段必须与计算流水线阶段同步以确保正确行为。如图1b所示，这意味着插入协调的到达/等待屏障。我们的编译器自动化了CudaDMA中用于单缓冲和双缓冲的过程。对于这两种情况，我们首先通过检查PDG来识别包围LDGSTS的一对BAR.SYNC屏障指令。对于单缓冲，每个BAR.SYNC指令在两个warp中使用BAR.WAIT和BAR.ARRIVE指令在相同的程序位置被替换为到达/等待屏障（示例见第IV-C节）。双缓冲更为复杂，需要将SMEM缓冲区和到达/等待屏障加倍【1, M. Bauer, H. Cook, and B. Khailany, “Cudadma: Optimizing gpu memory bandwidth via warp specialization,” in Proceedings of 2011 International Conference for High Performance Computing, Networking, Storage and Analysis, ser. SC ’11. New York, NY, USA: Association for Computing Machinery, 2011.】。我们使用LDGSTS指令的地址回溯切片和nvdisasm提供的SMEM分配信息来识别使用的是哪个共享缓冲区。会进行检查以确定是否有足够的SMEM容量用于双缓冲，然后调整大小并应用转换。图10展示了程序如何使用这个更大的缓冲区来实现图1b中的结果。在两个包围屏障之间的全局到SMEM传输的子程序被复制，以创建两个访问SMEM缓冲区不同半区的新子程序，每个子程序使用独特的到达/等待屏障集（图10中的A和B）。在子程序前后添加了新的基本块（swap_buffer, completed），其中包含通过交替基地址来在两个缓冲区之间切换执行的指令。对于后一个流水线阶段（图10中的阶段1），屏障A最初被设置为已到达。

最终SASS程序的生成。最后，将流水线阶段的程序组合起来，创建一个完整的SASS程序。首先执行寄存器重新分配；WASP编译器通过将寄存器压缩到连续空间来进行简单的重新分配。线程块规范被更新，以分配每个阶段的寄存器大小（表I）。然后将流水线阶段顺序写入最终的SASS程序。在程序的顶部添加一个跳转表，使用WASP显式命名的特殊寄存器（第III-A节）将每个warp引导到相应的代码段。

C. 内存访问模式示例

不同模式的转换实例。图11、12和13描绘了为流式全局到RFQ、收集全局到RFQ和全局到SMEM模式生成warp专业化流水线的示例。每个示例都关注CUDA内核的内层循环，并描绘了该内层循环基本块的指令级数据流图。

流式全局到RFQ。图11显示了一个用于将输入向量复制到输出向量的流式全局到RFQ示例。在内层循环中，warp执行一个全局加载指令（LDG），将数据转发到一个全局存储指令（STG）。LDG和STG指令都有自己的地址生成指令（ADD），内层循环还有管理控制流的指令（黑色）。考虑到内层循环中的单个LDG指令，WASP内核由两个流水线阶段（warp0 stage0和warp0 stage1）组成。对LDG地址有贡献的指令的向后切片被分配到较早的流水线阶段，其余指令分配到较晚的流水线阶段。基本块的控制流指令在两个warp中被复制，以维持两个阶段之间的一致执行。最后，LDG的目的操作数通过QS0S1命名队列连接到STG的源操作数。

int* in; //global memory   
int* out; //global memory   
...   
int start= (…) + threadId.x;   
int end = start + warp_transfer_size; 
for (int i=start; i<end; i+=WARP_WIDTH) 
    out[i] = in[i];

收集全局到RFQ。图12显示了一个收集全局到RFQ操作如何导致一个额外的流水线阶段。CUDA内层循环从全局内存数组index中读取，并使用结果作为data数组的索引。这种间接性在DFG中表现为额外的LDG和地址生成。与之前的全局到RFQ示例类似，原始内层循环的两个LDG指令都被用作分区的点。总共生成了三个阶段，并使用复制的控制流指令和队列来连接流水线阶段。

int* index; //global memory 
int* data; //global memory 
int* out; //global memory 
... 
int start= (…) + threadId.x; 
int end = start + warp_transfer_size; 
for (int i=start; i<end; i+=WARP_WIDTH) 
    out[i] = data[ index[i] ];

全局到SMEM。图13显示了WASP如何使用单缓冲为全局到SMEM的内存块传输生成代码。在原始CUDA程序中，为缓冲区数组分配了共享内存，该数组在线程块中的所有线程之间共享。__syncthreads()函数用于在传输前后同步缓冲区的状态。数据流图在其他方面与流式复制示例相似，唯一的例外是SMEM存储指令（STS）替换了STG。首先，LDG和STS指令被融合成一个单一的LDGSTS指令，适当的地址生成和控制流指令被分配到较早的流水线阶段warp0 stage0。由于内层循环中没有计算，warp0 stage1是空的。最后，插入BAR.WAIT和BAR.ARRIVE指令来同步线程，类似于图1b中的示例。

int* in; //global memory   
__shared__ int buffer[]; 
..   
__syncthreads(); //sync buffer ready   
int start= (…) + threadId.x;   
int end = start + warp_transfer_size;   
for (int i=start; i<end; i+=WARP_WIDTH) 
    buffer[i] = in[i];   
__syncthreads(); //sync transfer done

A4 实验环境

• 数据集/基准测试: 实验使用了20个来自多个领域的基准测试程序，包括MLPerf、cuSPARSE、机器学习、机器人和高性能计算。这些程序经过了充分优化，具体列表见下表II。
  表II：基准测试及WASP带来的中位数/最大内核加速比
  

  

• 模型架构: 本文未评估特定模型，而是评估在各种应用内核上的性能。
• 硬件配置: 实验基于模拟的NVIDIA Ampere A100 GPU架构进行。具体配置如下表III所示，该模型增加了对快速arrive/wait屏障和类似TMA的加速器支持，以模拟现代GPU特性。WASP的模拟在此基础上增加了新的线程块规范、RFQs、流水线感知的映射、分配、调度和WASP-TMA加速器。
  表III：NVArchSim A100+ 配置
  

  

• 软件配置:
  • 模拟器: 使用了NVIDIA的NVArchSim (NVAS) 的修改版，这是一个经过验证的混合跟踪与执行驱动的GPU模拟器。
  • 编译器: 自研的WASP编译器工具，用于自动生成warp专业化的SASS程序和线程块配置。编译器根据性能提升情况，逐个内核决定是否应用warp专业化转换。
  • 基线模型: 基线GPU模型在GEMM和cuBLAS内核上模拟了CUTLASS的warp专业化，利用WASP编译器进行粗粒度块传输和理想化的warp映射。

A5 实验结果

总体性能提升: 图14显示，与采用CUTLASS warp专业化的现代GPU基线（BASELINE）相比，完整的WASP系统（WASP GPU+WASP COMPILER ALL）在所有20个基准测试中均实现了超过10%的性能提升，平均性能提升达到47%。仅使用WASP编译器（WASP COMPILER ALL）而不依赖WASP硬件，通过软件方式（SMEM队列）支持细粒度模式，也能带来23%的平均性能提升，证明了自动化编译器的价值。

WASP硬件特性贡献分析: 图15逐步展示了各项WASP硬件特性带来的性能增益。在编译器支持的基础上：
* 流水线感知寄存器分配: 减少了寄存器文件压力，提高了SM占用率，带来了4%的几何平均性能提升。
* WASP-TMA: 通过卸载地址生成，减少了动态指令数，为bfs (+29%) 和mst (+30%) 等指令发射受限的应用带来显著好处，总体上额外提供了2%的几何平均性能提升。
* 寄存器文件队列 (RFQs): 替代了高开销的SMEM队列，显著降低了SMEM带宽消耗和指令开销。对SMEM带宽敏感的应用如3d_unet (+32%) 和spmm2_web (+32%) 提升超过15%，其他应用平均提升4%。
* 流水线感知Warp调度: 优先调度早期内存访问阶段的warp，促进了执行重叠，在所有基准上额外带来4%的平均性能提升。

寄存器分配效率: 图16表明，传统的统一寄存器分配方案会导致warp专业化后线程块的寄存器占用量急剧增加（在某些内核中接近翻倍）。而WASP的每阶段分配策略平均将寄存器占用量减少了44%，使得相对于非warp专业化基线的开销仅为26%。

Warp调度策略影响: 图17显示，优先处理流水线早期阶段（producer-first）的调度策略通常性能最佳，因为它能促进内存级并行。结合队列状态（full/ready queue）的策略在计算和内存访问不平衡的应用（如3d_unet, pointnet）上表现更优。

RFQ队列大小敏感性: 图18表明，RFQ队列大小对性能有显著影响。平均而言，每个通道32个条目是最佳平衡点，相比8个条目提升了15%的性能。更大的队列会因寄存器压力增加而降低SM上的线程块并发度，从而影响性能。

动态指令开销: 图19显示，WASP编译器因引入额外warp管理控制流，平均增加了18%的动态指令。然而，当使用WASP-TMA硬件卸载后，内存、地址生成和控制指令被大幅减少，使得在hpcg、pointnet等应用中，总动态指令数甚至可以降低到基线的46%。

内存带宽利用率和敏感性: 图21显示WASP普遍提高了L2带宽利用率。图20的敏感性分析表明，在内存带宽减半的情况下，WASP仍能达到甚至超过基线在正常带宽下的性能，因为它能更有效地利用有限的带宽。在带宽加倍时，WASP能更好地利用增加的带宽，而基线则受限于地址生成开销。

硬件复杂性: 表IV分析了WASP的硬件开销，主要是元数据存储。增加的总存储需求小于162KB，预计总额外硬件面积开销低于整个GPU芯片面积的1%。

表IV：WASP面积开销（存储需求）

A7 补充细节 (相关工作)

重叠内存访问与计算的现有工作。在GPU上重叠内存访问和计算方面已有大量先前工作。Warp专业化最早由CudaDMA【1, M. Bauer, H. Cook, and B. Khailany, “Cudadma: Optimizing gpu memory bandwidth via warp specialization,” in Proceedings of 2011 International Conference for High Performance Computing, Networking, Storage and Analysis, ser. SC ’11. New York, NY, USA: Association for Computing Machinery, 2011.】引入，用于到SMEM的粗粒度内存块传输，并已在许多应用中用于重叠计算和内存【2, M. Bauer, S. Treichler, and A. Aiken, “Singe: Leveraging warp specialization for high performance on gpus,” in Proceedings of the 19th ACM SIGPLAN symposium on Principles and practice of parallel programming, 2014, pp. 119–130.】【19, N. Maruyama and T. Aoki, “Optimizing stencil computations for nvidia kepler gpus,” 2014.】【48, H. Wei, E. Liu, Y. Zhao, and H. Yu, “Efficient non-fused winograd on gpus,” in Advances in Computer Graphics: 37th Computer Graphics International Conference, CGI 2020, Geneva, Switzerland, October 20–23, 2020, Proceedings. Berlin, Heidelberg: Springer-Verlag, 2020, p. 411–418.】。类似的工作将规则的仿射计算从GPU内核的其余部分解耦出来【46, K. Wang and C. Lin, “Decoupled affine computation for simt gpus,” in Proceedings of the 44th Annual International Symposium on Computer Architecture, ser. ISCA ’17. New York, NY, USA: Association for Computing Machinery, 2017, p. 295–306.】。WASP通过新的硬件和编译器支持扩展了warp专业化的最新技术。

调度与融合技术。先前的工作也专注于改进warp和线程块调度以更好地重叠内存访问和计算【39, A. Sethia, D. A. Jamshidi, and S. Mahlke, “Mascar: Speeding up gpu warps by reducing memory pitstops,” in 2015 IEEE 21st International Symposium on High Performance Computer Architecture (HPCA), 2015, pp. 174–185.】以及通过改进内存访问调度来提高缓存局部性【16, S.-Y. Lee, A. Arunkumar, and C.-J. Wu, “Cawa: Coordinated warp scheduling and cache prioritization for critical warp acceleration of gpgpu workloads,” in Proceedings of the 42nd Annual International Symposium on Computer Architecture, ser. ISCA ’15. New York, NY, USA: Association for Computing Machinery, 2015, p. 515–527.】【36, T. G. Rogers, M. O’Connor, and T. M. Aamodt, “Cache-conscious wavefront scheduling,” in 2012 45th Annual IEEE/ACM International Symposium on Microarchitecture, 2012, pp. 72–83.】。一项特别的工作研究了水平融合作为一种并发执行具有相反内存密集型和计算密集型行为的内核的方法【17, A. Li, B. Zheng, G. Pekhimenko, and F. Long, “Automatic horizontal fusion for gpu kernels,” in 2022 IEEE/ACM International Symposium on Code Generation and Optimization (CGO), 2022, pp. 14–27.】。WASP使用warp专业化来显式提取单个内核内的并行性，并重叠内存访问和计算。

硬件加速器与预取技术。先前关于粗粒度传输的工作包括D2MA和Hopper，它们采用了暴露给程序员的硬件加速器【11, D. A. Jamshidi, M. Samadi, and S. Mahlke, “D2ma: Accelerating coarsegrained data transfer for gpus,” in Proceedings of the 23rd International Conference on Parallel Architectures and Compilation, ser. PACT ’14. New York, NY, USA: Association for Computing Machinery, 2014, p. 431–442.】【23, Nvidia h100 tensor core gpu architecture, 2022】，以及为收集/跨步模式添加到GPU ISA中的向量指令【8, N. C. Crago, M. Stephenson, and S. W. Keckler, “Exposing memory access patterns to improve instruction and memory efficiency in gpus,” ACM Trans. Archit. Code Optim., vol. 15, no. 4, oct 2018.】。WASP使用WASP-TMA，这是一种新的用于细粒度和粗粒度数据移动的卸载加速器，它与GPU上的warp专业化流水线集成。预取技术在GPU的广泛背景下已有探索【15, J. Lee, N. B. Lakshminarayana, H. Kim, and R. Vuduc, “Manythread aware prefetching mechanisms for gpgpu applications,” in 2010 43rd Annual IEEE/ACM International Symposium on Microarchitecture, 2010, pp. 213–224.】【38, A. Sethia, G. Dasika, M. Samadi, and S. Mahlke, “Apogee: Adaptive prefetching on gpus for energy efficiency,” in Proceedings of the 22nd International Conference on Parallel Architectures and Compilation Techniques, 2013, pp. 73–82.】，用于改进线程块调度【13, G. Koo, H. Jeon, Z. Liu, N. S. Kim, and M. Annavaram, “Cta-aware prefetching and scheduling for gpu,” in 2018 IEEE International Parallel and Distributed Processing Symposium (IPDPS), 2018, pp. 137–148.】【29, Y. Oh, K. Kim, K. Y. Myung, H. P. Jong, Y. Park, W. R. Won, and M. Annavaram, “Apres: Improving cache efficiency by exploiting load characteristics on gpus,” in 2016 ACM/IEEE 43rd Annual International Symposium on Computer Architecture (ISCA), 2016, pp. 191–203.】，以及用于间接内存收集【49, X. Yu, C. J. Hughes, N. Satish, and S. Devadas, “Imp: Indirect memory prefetcher,” in Proceedings of the 48th International Symposium on Microarchitecture, ser. MICRO-48. New York, NY, USA: Association for Computing Machinery, 2015, p. 178–190.】。WASP利用warp专业化仅向内存子系统发出按需请求，避免了推测性内存访问。

其他流水线并行方案。现代GPU还支持在不同SM上执行的线程块之间的流水线并行。Nvidia的Hopper H100允许通过分布式共享内存实现SM之间的直接数据通信【23, Nvidia h100 tensor core gpu architecture, 2022】。VersaPipe使线程块能够使用存在于全局内存中的任务队列动态地充当不同的流水线阶段【50, Z. Zheng, C. Oh, J. Zhai, X. Shen, Y. Yi, and W. Chen, “Versapipe: a versatile programming framework for pipelined computing on gpu,” in Proceedings of the 50th Annual IEEE/ACM International Symposium on Microarchitecture, 2017, pp. 587–599.】。与WASP不同，这两种技术都没有在硬件或软件中显式利用流水线阶段的异构性，也不在SM级别提供执行重叠。Symphony是一种新的类似GPU的加速器和编程语言设计，它支持显式编程连接在一起的流水线阶段【31, M. Pellauer, J. Clemons, V. Balaji, N. Crago, A. Jaleel, D. Lee, M. O’Connor, A. Parashar, S. Treichler, P.-A. Tsai, S. W. Keckler, and J. S. Emer, “Symphony: Orchestrating sparse and dense tensors with hierarchical heterogeneous processing,” ACM Trans. Comput. Syst., vol. 41, no. 1–4, dec 2023.】。WASP为流水线并行提供了对程序员更透明的编译器支持，同时在现有GPU上集成了支持性的硬件特性。

解耦访问-执行处理器 (DAE)。解耦访问-执行处理器（DAE）显式地将内存和计算拆分为两个并发执行的独立程序【40, J. E. Smith, “Decoupled access/execute computer architectures,” in Proceedings of the 9th Annual Symposium on Computer Architecture, ser. ISCA ’82. Washington, DC, USA: IEEE Computer Society Press, 1982, p. 112–119.】。DAE已被引入多核处理器【10, T. J. Ham, J. L. Aragon, and M. Martonosi, “Desc: Decoupled supply-´ compute communication management for heterogeneous architectures,” ser. MICRO-48. New York, NY, USA: Association for Computing Machinery, 2015, p. 191–203.】【47, Z. Wang and T. Nowatzki, “Stream-based memory access specialization for general purpose processors,” in Proceedings of the 46th International Symposium on Computer Architecture, ser. ISCA ’19. New York, NY, USA: Association for Computing Machinery, 2019, p. 736–749.】，增强以在多线程多核处理器上处理内存间接【9, N. C. Crago and S. J. Patel, “Outrider: Efficient memory latency tolerance with decoupled strands,” SIGARCH Comput. Archit. News, vol. 39, no. 3, p. 117–128, jun 2011.】，并集成到乱序核心中以支持图算法【20, Q. M. Nguyen and D. Sanchez, “Pipette: Improving core utilization on irregular applications through intra-core pipeline parallelism,” in 2020 53rd Annual IEEE/ACM International Symposium on Microarchitecture (MICRO), 2020, pp. 596–608.】。其他相关工作开发了语言和编译器，将程序划分为子上下文，在通过队列通信的并行处理器上执行【30, G. Ottoni, R. Rangan, A. Stoler, and D. August, “Automatic thread extraction with decoupled software pipelining,” in 38th Annual IEEE/ACM International Symposium on Microarchitecture (MICRO’05), 2005, pp. 12 pp.–118.】【37, K. Sankaralingam, R. Nagarajan, H. Liu, C. Kim, J. Huh, D. Burger, S. W. Keckler, and C. R. Moore, “Exploiting ilp, tlp, and dlp with the polymorphous trips architecture,” SIGARCH Comput. Archit. News, vol. 31, no. 2, p. 422–433, may 2003.】【42, W. Thies, M. Karczmarek, and S. Amarasinghe, “Streamit: A language for streaming applications,” in Compiler Construction: 11th International Conference, CC 2002 Held as Part of the Joint European Conferences on Theory and Practice of Software, ETAPS 2002 Grenoble, France, April 8–12, 2002 Proceedings 11. Springer, 2002, pp. 179–196.】【44, F. Tseng and Y. N. Patt, “Achieving out-of-order performance with almost in-order complexity,” in Proceedings of the 35th Annual International Symposium on Computer Architecture, ser. ISCA ’08. USA: IEEE Computer Society, 2008, p. 3–12.】。WASP专注于为GPU设计新的硬件和软件机制，以实现流水线感知的执行。

A6 结论

本文提出了WASP，一套用于warp专业化的硬件和编译器支持。Warp专业化是一种通过重叠内存访问和计算操作来提升GPU性能的强大技术。WASP通过引入以下创新改进了现有技术：处理细粒度内存访问模式的新硬件、流水线感知的warp映射和调度，以及一个减轻程序员负担的编译器。研究发现，仅WASP编译器就能比现有顶尖GPU提升23%的运行时性能，当与新的WASP硬件结合时，性能提升可达47%。