Optimizing Cache Bypassing and Warp Scheduling for GPUs

文章标题：优化GPU的缓存绕过和Warp调度
作者/机构：Yun Liang, Xiaolong Xie, Yu Wang, Guangyu Sun, and Tao Wang

A1 主要贡献

本文旨在解决图形处理单元（GPU）中因大规模线程并行性导致的缓存资源争用问题。该问题通常会导致缓存命中率低、内存子系统性能受限以及流水线停顿。为了缓解这些问题，作者提出了一种协同的静态（编译时）和动态（运行时）缓存绕过框架，并设计了一种与之配合的旁路感知warp调度器（Bypass-Aware Warp Scheduler, BAWS）。

核心问题：
1. 高缓存争用：大规模并发线程争抢有限的片上缓存资源，导致缓存命中率极低（例如，对于16-kB L1数据缓存，平均命中率仅为27.1%）。
2. 流水线停顿：低缓存命中率导致大量的缓存未命中，进而引发未命中状态保持寄存器（MSHRs）等相关资源的拥塞，造成流水线停顿。在许多情况下，使用GPU缓存甚至会损害性能。

研究目标与创新点：
本文的主要目标是通过智能的缓存绕过机制，在不牺牲大规模线程并行性的前提下，提升GPU应用的性能。其核心贡献如下：
1. 协同的静态与动态缓存绕过框架：
* 静态（编译时）：通过分析内存访问模式或性能剖析，将全局加载指令（global loads）分为三类：具有良好局部性的（标记为ca）、具有差局部性的（标记为cg）和具有中等局部性的（标记为cm）。这种分类信息被编码到应用程序二进制文件中，为硬件提供全局局部性提示。
* 动态（运行时）：硬件根据编译时提供的提示，对ca和cg指令执行固定的缓存或绕过策略。对于cm指令，硬件在线程块（thread block）级别动态调整使用缓存的线程块比例，以平衡缓存利用和资源争用。
2. 内存访问模式分析与性能剖析：提出了两种方法来对全局加载指令进行分类：一种是基于静态分析的内存访问模式识别，另一种是基于性能剖析的方法。
3. 运行时管理技术：开发了运行时技术，通过在线学习动态调整使用或绕过缓存的线程块比例，以适应程序运行时的行为变化。
4. 旁路感知Warp调度器（BAWS）：设计了一种新的warp调度器，它能够根据实时的缓存性能指标（如缓存命中和流水线停顿情况），自适应地调整缓存warp（caching warps）和绕过warp（bypassing warps）的调度优先级，从而进一步优化性能。

关键成果：
* 协同的静态和动态缓存绕过技术在多种GPU应用中取得了最高2.28倍（几何平均1.32倍）的性能加速。
* 与最先进的缓存绕过技术相比，该协同方法将性能加速从1.11倍和1.17倍提升至1.32倍。
* 结合旁路感知warp调度器（BAWS），平均性能加速进一步提升至1.38倍。


图 1. (a) GPU架构。(b) CPU和(c) GPU上的缓存绕过。

A3 背景知识/关键Observation/设计原则

II. 背景

A. GPU架构

GPU的层次化结构。GPU由多个流式多处理器（Streaming Multiprocessors, SMs）构成。每个SM内部集成了多个流处理器（Streaming Processors）、寄存器、暂存板内存（scratchpad memory，也称共享内存）和L1数据缓存。SM以单指令多数据（SIMD）的方式协调执行。这些SM通过互联网络连接到L2缓存。图1(a)展示了本文的目标GPU架构。GPU应用程序的线程被组织成线程块（thread blocks），同一线程块内的线程可以通过共享内存共享数据并使用屏障指令进行同步。一个SM上能同时执行的线程块数量受限于可用资源。每个线程块进一步组织成warps，每个warp包含32个线程，warp内的线程以SIMD方式执行。

图 2. 将L1缓存大小和共享内存大小从16kB增加到48kB时的性能结果。

B. GPU缓存特性

GPU缓存的特点与挑战。GPU采用两级缓存层次结构，L1缓存位于片上，L2缓存位于片外。与软件管理的共享内存不同，缓存由硬件管理。通用应用程序天然倾向于使用缓存。如图2所示，将L1数据缓存的大小从16kB增加到48kB，可带来1.17倍的性能提升；而将共享内存从16kB增加到48kB，性能仅提升1.02倍。因此，本文专注于缓存优化。由于大规模并行性，缓存争用是GPU上的一个严重问题。对于一个典型的16-kB L1数据缓存，当一个SM上并发运行1024个线程时，每个线程的缓存容量仅为16字节。如此小的每线程缓存容量导致了低缓存命中率——对于Rodinia【3, Rodinia: A benchmark suite for heterogeneous computing, 2009, IISWC】和Parboil【4, The IMPACT Research Group, Parboil Benchmark Suite, http://impact.crhc.illinois.edu/Parboil/parboil.aspx】基准套件中的应用 ，16-kB L1数据缓存的平均命中率仅为27.1%。L2缓存通常足够大，可以捕获大部分数据局部性。当我们将同一组应用的L2缓存大小加倍时，平均性能提升不到1%。因此，本文仅关注L1缓存。除了缓存争用，大规模并行性带来的另一个问题是流水线停顿。当发生缓存未命中时，缓存首先在MSHR（Miss Status Holding Registers）表中寻找一个空闲条目，并在相应的缓存组中寻找一个空闲的缓存行【11, W. Jia, K. A. Shaw, and M. Martonosi, “MRPB: Memory request prioritization for massively parallel processors,” in Proc. 20th Int. Symp. High Perform. Comput. Archit. (HPCA), Orlando, FL, USA, 2014, pp. 272–283.】。该缓存行会为这次未命中保留，直到服务完成。如果没有空闲的MSHR条目或相应缓存组中的所有缓存行都已被预留，流水线就必须停顿，直到有空闲资源可用。如果缓存未命中频繁发生，由于流水线停顿，访问缓存甚至可能损害性能。

C. 缓存绕过

缓存绕过的机制与平台差异。缓存绕过已被证明能有效缓解CPU的缓存争用和颠簸（thrashing）问题【12, Y. Wu et al., “Compiler managed micro-cache bypassing for high performance EPIC processors,” in Proc. 35th Annu. ACM/IEEE Int. Symp. Microarchit. (MICRO), Istanbul, Turkey, 2002, pp. 134–145.】。近期的GPU架构，如NVIDIA的Fermi、Kepler和Maxwell，也开始引入缓存绕过。然而，由于架构特性不同，两个平台上的缓存绕过接口实现方式有所不同。图1(b)展示了CPU上的缓存绕过实现。在CPU上，L1缓存命中率通常很高，因此每个内存请求总是会查询L1缓存。缓存绕过主要用于末级缓存（LLC）【13, S. M. Khan, Y. Tian, and D. A. Jimenez, “Sampling dead block prediction for last-level caches,” in Proc. 43rd Annu. IEEE/ACM Int. Symp. Microarchit. (MICRO), Atlanta, GA, USA, 2010, pp. 175–186.】、【14, C.-J. Wu et al., “Ship: Signature-based hit predictor for high performance caching,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 430–441.】。根据数据局部性，缓存绕过可以选择性地在分配阶段将数据分配到LLC或绕过LLC。如图1(b)所示，从内存请求的数据只分配到L1缓存，而L2缓存被绕过。在GPU上，内存请求可以在访问时选择绕过L1缓存，如图1(c)所示。之后，请求的数据直接转发给处理器，而不填充L1缓存。这种缓存绕过已在NVIDIA Fermi（GTX480）和Kepler（GTX680）架构上实现。与CPU上的缓存绕过不同，GPU上的缓存绕过不仅缓解了缓存争用，还减少了资源拥塞和流水线停顿。当前几代GPU上的缓存绕过支持仍较初级。例如，在NVIDIA Fermi和Kepler上，程序员可以通过控制编译标志（-dlcm=ca或-dlcm=cg）来为整个程序选择使用或绕过缓存。然而，这种粗粒度的解决方案未能充分利用缓存和绕过所提供的全部机会。目前，GPU不在L1缓存中缓存全局存储数据，因为L1缓存对于全局数据是不一致的。因此，全局存储通过使L1缓存中的副本失效来直接更新L2缓存。因此，与之前的工作【5, W. Jia, K. A. Shaw, and M. Martonosi, “Characterizing and improving the use of demand-fetched caches in GPUs,” in Proc. Int. Conf. Supercomput. (ICS), 2012, pp. 15–24.】、【11, W. Jia, K. A. Shaw, and M. Martonosi, “MRPB: Memory request prioritization for massively parallel processors,” in Proc. 20th Int. Symp. High Perform. Comput. Archit. (HPCA), Orlando, FL, USA, 2014, pp. 272–283.】类似，我们只关注全局内存加载。

III. 系统概述与动机

系统实现概览。本文提出的技术基于GPGPU-Sim【15, A. Bakhoda, G. L. Yuan, W. W. L. Fung, H. Wong, and T. M. Aamodt, “Analyzing CUDA workloads using a detailed GPU simulator,” in Proc. IEEE Int. Symp. Perform. Anal. Syst. Softw. (ISPASS), Boston, MA, USA, Apr. 2009, pp. 163–174.】（版本3.2.1）的编译和运行时系统实现，如图3所示。输入是PTX格式的GPU应用程序代码，这是CUDA代码的中间表示。静态缓存绕过组件在编译时分析PTX代码，根据内存请求的局部性对其进行分类，并通过指令集扩展将分类编码。在运行时，动态缓存绕过组件对于那些有强烈缓存或绕过偏好的内存请求，会遵循其决定，但对于其余的内存请求，则有灵活性在线程块级别调整其行为。之后，缓存旁路感知warp调度器（BAWS）管理并发的缓存和绕过warp的调度。


图 3. 系统概览。静态绕过使用标签ca、cg和cm将全局加载分为三类。标记为ba的线程块使用L1缓存，而标记为bg的线程块绕过L1缓存。BAWS管理并发warp的执行以进一步提高性能。

A. 静态缓存绕过概述

内存访问模式的多样性。通用应用程序中使用的内存访问模式多种多样。一些访问，如array[tid % 2]和array[bid]，可能天生具有良好的局部性。相反，一些访问可能天生具有较差的局部性，例如用于读写流数据的访问。最后，还有一些访问既没有好的也没有坏的局部性，它们主要受到大规模并发warp引起的严重缓存争用的影响。我们称之为中等局部性访问。
静态分类机制。静态缓存绕过组件首先检测所有的全局加载指令。在我们的目标平台上，PTX中的全局加载指令格式如下：ld.global.l1_cache_tag, ...，其中l1_cache_tag字段用于指定访问类型。原生存在两种选项：1) ca 和 2) cg。如果一个全局加载被标记为ca，那么内核的所有线程在执行该加载时都会使用缓存。如果一个全局加载被标记为cg，那么内核的所有线程在执行该加载时都会绕过缓存。这里，我们为具有中等局部性的加载指令扩展了第三个选项，称为cm。静态缓存绕过组件将所有全局加载分为三类：1) ca；2) cg；和 3) cm，并将分类编码到全局加载指令中，如图3所示。我们用ca表示具有良好局部性的指令；cg表示具有差局部性的指令；cm表示具有中等局部性的指令。第四节提供了静态缓存绕过组件的实现细节。

B. 动态缓存绕过概述

运行时决策的灵活性。动态缓存绕过组件首先解码由静态缓存绕过组件编码的全局加载分类。对于标记为ca或cg的全局加载，它遵循编译时做出的决定。然而，对于标记为cm的全局加载，我们将缓存或绕过的决定留给运行时的动态缓存绕过。动态缓存绕过的目标是在缓存和绕过之间取得平衡——让一部分线程绕过缓存以减轻缓存争用和资源拥塞，而其他线程仍然使用缓存以利用数据局部性。
线程块级别的决策。如果一个warp中的线程获取数据的延迟不同（例如，缓存和绕过的行为不同），那么该warp必须停顿直到最长的内存请求完成。此外，GPU应用程序通常使用屏障指令来同步同一线程块中的线程。如果由于不同的缓存行为，线程到达屏障的时间不同，先到的线程必须等待后到的线程才能一起前进。在这两种情况下，性能都会因为等待线程占用的资源闲置而下降。因此，我们选择在线程块级别上调节缓存或绕过的线程，以确保同一线程块中的线程具有相同的缓存或绕过行为。
线程块标签与动态调节。为了区分缓存和绕过的线程块，每个线程块在分派时被分配一个标签。我们使用标签ba或bg。如果一个线程块被标记为ba，那么其中的所有线程在执行标记为cm的全局加载时都将使用缓存。如果一个线程块被标记为bg，那么其中的所有线程在执行标记为cm的全局加载时都将绕过缓存。一旦GPU核心取到一个全局加载指令，它会根据全局加载的标签和当前线程块的标签来决定其行为。我们用$TB_{max}$表示一个SM上可以同时执行的最大线程块数。$TB_{max}$由每个线程块的资源使用量和SM的可用资源决定。我们用$TB_{ba}$和$TB_{bg}$分别表示运行时每个SM上标记为ba和bg的活动线程块数。根据定义，我们有$TB_{ba} + TB_{bg} = TB_{max}$。动态缓存绕过组件在运行时调整$TB_{bg}$并确定每个线程块的标签。例如，在图3中，线程块2-5被标记为bg，而线程块0和1被标记为ba。第五节提供了动态缓存绕过组件的实现细节。

C. 缓存旁路感知Warp调度器

Warp调度的重要性。Warp调度器对GPU系统性能有相当大的影响【7, V. Narasiman et al., “Improving GPU performance via large warps and two-level warp scheduling,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 308–317.】，因为它决定了活动warp被调度的顺序。在每个周期，warp调度器从就绪的warp中选择一个。一个warp可能因为数据依赖或长内存延迟而停顿。默认的松散轮询（Loosely Round-Robin, LRR）调度器平等对待所有warp，并轮流调度它们。然而，这已被证明是低效的，因为它可能加剧缓存争用问题【7, V. Narasiman et al., “Improving GPU performance via large warps and two-level warp scheduling,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 308–317.】、【16, A. Jog et al., “OWL: Cooperative thread array aware scheduling techniques for improving GPGPU performance,” in Proc. Int. Conf. Archit. Support Program. Lang. Oper. Syst. (ASPLOS), Houston, TX, USA, 2013, pp. 395–406.】。我们用k表示一个线程块中的warp数。在动态缓存绕过组件的控制下，在所有的$TB_{max} \times k$个warp中，$TB_{ba} \times k$个warp使用缓存，而$TB_{bg} \times k$个warp绕过缓存。现有的warp调度器不区分这些缓存和绕过的warp。此外，它们只关注提高缓存命中率。然而，由过多缓存未命中引起的流水线停顿也是GPU缓存的一个重要性能因素。因此，我们提出了一种缓存旁路感知warp调度器（BAWS），通过warp调度来缓解缓存争用和流水线停顿。具体细节在第六节。


图 4. 使用SPV应用的协同缓存绕过和旁路感知warp调度策略的动机。(a) 指令特性。(b) TBbg的性能影响。(c) Warp调度策略。

D. 动机

SPV应用案例。我们使用SPV应用来说明协同静态和缓存绕过的好处。图4(a)通过剖析显示了SPV中所有全局加载的L1缓存命中率。总共有十个全局加载。其中，ld8的缓存命中率非常高（> 99%），因此我们将其归类为良好局部性加载（标签ca）。然而，ld1, ld2, ld4, ld5, ld7, 和 ld10 的缓存命中率非常低（< 1%），因此我们将它们归类为差局部性加载（标签cg）。对于剩下的加载（ld3, ld6, 和 ld9），我们将其归类为中等局部性加载（标签cm）。图4(b)显示了将$TB_{bg}$从0变化到8时的性能（SPV的$TB_{max}$为8）。性能相对于默认设置（所有全局加载都使用缓存）进行了归一化。显然，我们的协同缓存绕过具有很高的性能提升潜力。在最左边的点，除了标记为cg（差局部性）的加载外，所有全局加载都使用缓存，这带来了1.1倍的加速。通过改变$TB_{bg}$，性能可以进一步加速到1.19倍（当$TB_{bg}=6$时）。请注意，在这个实验中，我们在整个内核执行期间固定了$TB_{bg}$。但在我们的动态缓存绕过中，我们可以在内核执行期间调整$TB_{bg}$。
BAWS的动机。我们使用SPV应用对不同的warp调度器进行了评估，以说明BAWS的好处。将三种warp调度器与默认的LRR策略进行比较：1) Greedy-then-oldest (GTO)，总是优先调度年龄最大的warp。2) Bypass-first，总是优先调度绕过缓存的warp。当所有就绪的warp都是绕过warp或缓存warp时，使用GTO策略选择一个。3) Cache-first，与bypass-first策略相同，但优先调度使用缓存的warp。在运行时，我们固定$TB_{bg}$为4。图4(c)显示了相对于LRR调度器的归一化性能结果。这些策略的性能加速分别为1.12倍、1.18倍和1.12倍。SPV应用因缓存资源拥塞而遭受缓存争用和流水线停顿。通过优先调度绕过缓存的warp，由于缓存压力得到缓解，流水线停顿的次数减少了。因此，bypass-first调度器优于其他调度器。不同的应用有不同的缓存偏好，因此需要一个自适应的warp调度器。我们的BAWS根据运行时信息自适应地应对不同场景（第六节）。

A2 方法细节

IV. 静态绕过的实现

本节介绍静态缓存绕过组件的实现细节，其目标是将全局加载指令分为好、差和中等三类局部性。我们提出了两种不同的方法：内存访问模式分析和基于性能剖析的方法。

A. 内存访问模式分析

基于模式的静态分类。我们发现，为了最大化内存吞吐量，大多数内存指令遵循几种常见的内存访问模式，其中一些模式表现出明确的缓存或绕过偏好，这为通过静态模式分析分类全局加载提供了机会。我们提出的模式分析依赖于分析内存地址的计算方式。通常，一个全局加载可以表示为一个函数，输入包括线程块标识符（bid）、线程标识符（tid）、循环索引（i）、常量（C）以及一些未知值（如链表中的指针）。我们不分析包含未知值的模式，因为它们依赖于输入且缓存行为不可预测。对于其余的全局加载，我们观察到大多数属于四种常用模式之一，如图5所示，具体细节在表I中。
* 模式一：代表流式数据操作，内存地址是tid和一个常量偏移（C）的和。对于这种模式，内存合并组件能完美利用数据局部性，而缓存无法带来额外好处，因此倾向于绕过缓存（标记为cg）。
* 模式二：通常使用bid或tid mod C1与一个常量（C0）的和来计算内存地址。在这种模式下，多个线程可能访问同一个缓存块，通常具有良好的数据局部性且不会引起缓存争用，因此倾向于使用缓存（标记为ca）。
* 模式三：常出现在矩阵或模板（stencil）应用中，使用二维或三维线程索引（tid.x, tid.y, tid.z）计算内存地址。由于工作集较大和复杂的数据映射机制，它们通常会遭受缓存争用。
* 模式四：代表循环中的连续数据操作，大量数据在短时间内被加载到缓存。模式三和模式四的共同特点是，它们通常存在数据共享，但由于工作集较大而遭受缓存争用和资源拥塞。这种复杂的权衡使得模式分析难以判断其缓存偏好。因此，我们保守地将它们归类为中等局部性加载（cm）。

我们的内存访问模式分析流程如下：首先分析PTX代码，递归地展开内存地址中每个源值的表达式，从而获得全局加载的内存地址函数。如果得到的地址函数包含未知值，我们简单地将其分类为cm加载。对于其余的加载，我们检查它是否匹配上述四种模式中的任何一种。如果不匹配任何模式，我们将其分类为cm加载。否则，分类由其匹配的模式决定，如表I所示。注意，我们不分析指令间的局部性。在实验部分，我们将展示我们的模式分析对大量GPU应用表现良好。

图 5. GPU应用中常用的内存访问模式。(a) 模式1。(b) 模式2。(c) 模式3。(d) 模式4。
TABLE I 常用内存访问模式的详细信息


图 6. 带标注的LCFG。访问同一数组的加载使用相同颜色。(a) LCFG。(b) 分组后的全局加载。

B. 基于性能剖析的方法

通过剖析数据进行分类。基于性能剖析的方法依赖于收集缓存性能指标来对加载指令进行分类。对于一个包含N个全局加载的GPU应用，我们用$ld_i$表示程序顺序中的第i个全局加载。我们首先分析GPU内核的PTX代码并构建内核的控制流图（CFG）。然后，我们将CFG扩展为加载CFG（LCFG）。具体来说，在LCFG中，每个节点代表一个全局加载。如果CFG中的一个基本块包含多个加载，它将被拆分为LCFG中的多个节点，每个节点对应一个加载。然后，我们用控制流边将拆分后的节点逐一连接。之后，如果一个基本块不包含任何加载，我们用一个虚拟节点替换它。如果源CUDA程序中两个节点之间存在同步，我们就在LCFG中的这两个节点之间插入一个同步节点。最后，我们从最后一个节点到第一个节点添加一条后向边。图6(a)展示了一个LCFG的例子。接着，我们构建指令间局部性图$G=(V, E)$。节点$v_i \in V$代表$ld_i$。如果节点$v_i$和$v_j$有指令间局部性机会，它们就通过一条无向边连接。然后，我们在G中找到所有的连通分量。连通分量可以使用多项式时间的广度优先搜索得到。每个连通分量被视为一个组。图6(b)展示了图6(a)中LCFG的连通分量。在这个例子中，它将LCFG中的加载分成了七个组。我们使用以下指标来表征每个加载的指令内局部性和每个组（即G中的连通分量）的指令间局部性：
1. access(i): $ld_i$的访问次数。
2. hit(i): 当$ld_i$独占使用缓存时，$ld_i$的L1缓存命中次数。
3. hit(g): 当组g中的加载独占使用缓存时，组g的L1缓存命中次数。

我们通过性能剖析获得这些指标。具体来说，我们为$ld_i$使用ca标签，为其余全局加载使用cg标签来收集access(i)和hit(i)。类似地，我们为组g中的加载使用ca标签，为其余全局加载使用cg标签来收集hit(g)。这意味着我们需要对应用程序进行$N+N_G$次插桩和剖析，其中N是GPU应用程序的全局加载数，$N_G$是分组后LCFG中的连通分量数。显然，$N_G \le N$。我们在二进制级别实现插桩【10, X. Xie, Y. Liang, G. Sun, and D. Chen, “An efficient compiler framework for cache bypassing on GPUs,” in Proc. Int. Conf. Comput.-Aided Design (ICCAD), San Jose, CA, USA, 2013, pp. 516–523.】。与剖析开销相比，插桩的开销可以忽略不计。最后，我们计算缓存命中数，然后将全局加载分为三类。基于剖析的方法的详细信息可以在【9, X. Xie, Y. Liang, Y. Wang, G. Sun, and T. Wang, “Coordinated static and dynamic cache bypassing for GPUs,” in Proc. Int. Symp. High Perform. Comput. Archit. (HPCA), Burlingame, CA, USA, 2015, pp. 76–88.】中找到。

C. 比较

两种静态方法的权衡。基于性能剖析的方法应该比内存访问模式分析更准确，而内存访问模式分析不产生任何剖析开销。内存访问模式分析可以应用于应用程序在编译时未知的动态系统，而基于剖析的方法可用于应用程序和输入在部署前已知的情况。

V. 动态绕过的实现

动态绕过的目标与方法。动态缓存绕过的目标是通过调整每个SM上使用缓存的活动线程块数（$TB_{ba}$）和绕过缓存的活动线程块数（$TB_{bg}$），在运行时平衡缓存和绕过行为。为实现此目标，它利用在线学习。学习过程包括以下三个步骤：
* 步骤1: 内核开始执行时，为$TB_{bg}$设置一个初始值。
* 步骤2: 当有$TB_{bg}$个标记为bg的活动线程块时启动计时器，并在一个采样周期P内保持$TB_{bg}$不变。
* 步骤3: 将当前采样周期的性能指标与历史记录进行比较，并更新$TB_{bg}$。
最初，我们设置$TB_{bg} = TB_{max}$。然后，我们迭代执行步骤2和3，直到所有线程块都执行完毕。在步骤2中，我们仅当标记为bg的活动线程块数达到目标值（$TB_{bg}$）时才启动计时器。我们使用一个线程块的生命周期作为采样周期，这是因为网格执行具有离散性。这也确保了一个线程块在执行期间不会改变其标签，从而避免了因不同缓存或绕过行为导致的内存分歧问题。

性能指标与更新策略。在步骤3中，我们首先定义缓存命中停顿分数（Cache Hit Stall Score, CHSS）作为性能指标，并在在线学习中使用它来调整$TB_{bg}$。CHSS的定义如下：
$$ \text{CHSS} = \frac{\text{Hits} \cdot \text{L2\_Latency}}{\text{Stall} \cdot \text{WarpCount}}. $$
具体来说，CHSS是L1缓存命中数（Hits）、由L1缓存拥塞引起的流水线停顿（Stall）和活动warp数（WarpCount）的函数。$L2\_Latency$是一个常量。缓存绕过通过缓解缓存争用和流水线停顿来提高性能。CHSS同时考虑了缓存命中节省的延迟（$Hits \cdot L2\_Latency$）和为流水线停顿付出的延迟（$Stall \cdot WarpCount$）。当CHSS > 1时，缓存命中能够补偿由大量缓存请求引起的流水线停顿，我们倾向于让更多线程块访问缓存。否则，使用缓存可能会损害性能，我们需要减少缓存请求的数量。理想情况下，我们希望同时最大化缓存命中并最小化停顿。因此，我们寻求一个能使CHSS尽可能大的$TB_{bg}$值。为此，我们维护一个CHSS表CHSS[]，其中CHSS[a]表示当$TB_{bg}=a$时采样周期内达到的CHSS值。我们在执行开始时将CHSS[]的值初始化为1。在步骤3结束时，使用当前周期收集的CHSS值更新CHSS[]表（条目CHSS[TB_{bg}]）。然后，我们通过比较CHSS[TB_{bg}]与其邻居CHSS[TB_{bg}+1]和CHSS[TB_{bg}-1]来确定下一个采样周期的$TB_{bg}$。
$$TB_{bg}=\begin{cases}TB_{bg}+1, & \text{if CHSS}[TB_{bg}+1] \text{ is the maximal} \\ TB_{bg}, & \text{if CHSS}[TB_{bg}] \text{ is the maximal} \\ TB_{bg}-1, & \text{if CHSS}[TB_{bg}-1] \text{ is the maximal.}\end{cases}$$

标签分配机制。接下来，我们说明如何为每个线程块确定标签（bg或ba）。当一个线程块被分派时，我们首先计算$Cur_{bg}$，即当前标记为bg的活动线程块数。注意，$Cur_{bg}$可能与目标$TB_{bg}$不同，因为线程块是动态分派和提交的。我们用tag[bid]表示分配给线程块bid（线程块标识符）的标签。我们通过比较$Cur_{bg}$与目标$TB_{bg}$来确定tag[bid]，如下所示：
$$ \text{tag}[\text{bid}] = \begin{cases} bg, & \text{if } \text{Cur}_{bg} < TB_{bg} \\ ba, & \text{otherwise.} \end{cases} $$
图7通过一个例子说明了采样周期、$TB_{bg}$更新和每个线程块的标签分配。在这个例子中，我们假设$TB_{max}=4$。最初，我们设置$TB_{bg}=4$。因此，前四个线程块都被标记为bg。我们在t1时刻为第一个周期启动计时器，因为在t1时正好有四个标记为bg的线程块。在一个采样周期（线程块3的生命周期）结束前，线程块0-2提交，线程块4-6被分派。当线程块4开始时，有三个标记为bg的块（线程块1-3）。因此，$Cur_{bg}=3$。然后，根据(2)，线程块4的标签被设置为bg。类似地，线程块5和6的标签也被设置为bg。第一个采样周期后，$TB_{bg}$变为3。然后，根据(2)，线程块7的标签被设置为ba。步骤2和3重复进行，直到所有线程块执行完毕。注意，第二个采样周期后，$TB_{bg}$变为2，我们在t4时刻为下一个采样周期启动计时器，因为在t4时有两个标记为bg的线程块（线程块11和12）。

图 7. 采样周期、TBbg更新和标签分配。

控制机制与硬件开销。最后，我们提出了两种在学习成本和灵活性上有所权衡的方法：集中式控制和分散式控制。在集中式控制中，只有SM0执行在线学习并将其决策（$TB_{bg}$）广播给其他SM。所有SM以相同的方式调节线程块。在分散式控制中，每个SM独立地执行学习。对于一个给定的GPU应用，其线程块以轮询方式分派到各个SM。SM可能会有不同的工作负载，因为线程块可能有不同的计算和内存行为。与集中式控制相比，分散式控制允许对不同的SM进行灵活调整，但学习成本更高。我们将在实验部分比较这两种方法。动态缓存绕过需要少量的硬件改动。首先，在GPU上，每个SM上同时活动的线程块数$TB_{max}$是有限的（例如8）。为了收集CHSS指标，需要两个32位硬件计数器来收集Hits和Stall，而$L2\_Latency$和$WarpCount$在执行期间是常量。CHSS[]表的大小是$TB_{max}+1$。因此，CHSS[]表只需要72字节。其次，对于每个线程块，我们只需要一个额外的位来表示其标签（ba或bg）。实际上，我们只需要为每个SM分配额外的$TB_{max}$个位，每个位代表一个活动线程块的标签。最后，需要两个额外的4位寄存器：一个用于$Cur_{bg}$，另一个用于$TB_{bg}$。$Cur_{bg}$在标记为bg的线程块退役或分派时更新。

VI. 旁路感知WARP调度策略的实现

本节讨论GPU的warp调度器，并提出我们的旁路感知warp调度器（Bypass-Aware Warp Scheduler, BAWS）。BAWS根据运行时的缓存性能，决定使用或绕过缓存的warp的调度顺序，以进一步缓解缓存争用并减少流水线停顿。

现有Warp调度器的回顾

现有方法的局限性。我们从回顾现有的warp调度器开始。LRR（松散轮询）【7, V. Narasiman et al., “Improving GPU performance via large warps and two-level warp scheduling,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 308–317.】调度器被用作本文的基线调度器。LRR平等对待所有活动warp并轮流调度它们。LRR有两个主要缺点。首先，warp很可能同时遇到长内存延迟，导致整个流水线停顿。其次，它可能损害缓存性能，因为所有warp可能在短时间内访问缓存。GTO（Greedy-then-oldest）和两级（two-level）策略被提出来改进LRR【7, V. Narasiman et al., “Improving GPU performance via large warps and two-level warp scheduling,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 308–317.】。GTO根据warp在GPU上存在的时间对所有warp进行优先级排序，最老的warp总是被优先调度。这已被证明能有效地保护数据局部性【8, T. G. Rogers, M. O’Connor, and T. M. Aamodt, “Cache-conscious wavefront scheduling,” in Proc. Int. Symp. Microarchitect. (MICRO), Vancouver, BC, Canada, 2012, pp. 72–83.】。两级调度【7, V. Narasiman et al., “Improving GPU performance via large warps and two-level warp scheduling,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 308–317.】将所有warp分成多个取指组，每个取指组的大小是经验性设置的。两级调度只从一个取指组中调度warp。当取指组中的所有warp都停顿时，它会移动到下一个取指组。因此，取指组内部的数据局部性得以保留。然而，正如我们在本文中展示的，缓存争用和流水线停顿都会影响GPU性能，而现有的warp调度器只关注缓存争用问题。为了达到最优性能，有必要设计一个同时考虑这两个因素的调度器。

图 8. 提出的旁路感知warp调度策略。(a) 概览。(b) 算法。

BAWS调度策略

BAWS的设计与算法。动态缓存绕过组件决定了线程块的缓存或绕过行为，从而决定了warp的缓存或绕过行为。我们将使用（绕过）缓存的warp称为ba（bg）warp。在图8(a)所示的情况下，我们有四个ba warp和六个bg warp。尽管动态缓存绕过组件调整了绕过缓存的线程块数量以缓解缓存争用和流水线停顿问题，但缓存争用仍可能发生。如前所述，现有的warp调度器不区分ba或bg warp。如果warp调度器在短时间内调度了许多ba warp，可能会导致缓存拥塞。另一方面，当缓存性能良好时，现有的调度器可能会调度许多bg warp，从而未能利用数据局部性。我们在图8(b)中展示了我们的BAWS算法。直观地说，当缓存性能良好时，BAWS应该优先调度那些ba warp；当缓存拥塞时，应该优先调度那些bg warp。为实现此目标，我们使用动态缓存绕过组件的CHSS指标（第五节）作为性能指示器。CHSS是缓存命中和流水线停顿的一个良好监视器。回顾CHSS的定义，它被定义为因缓存命中而节省的内存延迟与因缓存拥塞而产生的流水线停顿惩罚的商。理想情况下，当CHSS约等于1时，缓存命中和流水线停顿惩罚达到平衡。当CHSS过高（$\ge HThres$）时，我们认为缓存没有拥塞，并优先调度那些ba warp以利用缓存的好处，使用GTO策略在所有ba warp中调度一个。当CHSS过低（$\le LThres$）时，我们认为缓存拥塞，访问缓存会引起过多的流水线停顿。因此，为避免流水线停顿，使用GTO策略在所有bg warp中选择一个。否则，我们简单地使用基础的GTO策略，因为两种类型的warp都没有明显优势。
阈值设定与开销。当CHSS约为1时，缓存和绕过达到平衡。我们测试了$H_{thres}$在1.6-2.5范围内的一些值，步长为0.1。类似地，我们测试了$L_{thres}$在0.1-0.6范围内的一些值，步长为0.1。最后，我们根据经验将$H_{thres}$和$L_{thres}$设置为2和0.5，因为它们返回了最好的结果。与原始的GTO和两级策略相比，BAWS能够根据运行时的缓存争用和流水线停顿信息调整调度顺序。它帮助我们的协同缓存绕过在warp调度时微调缓存性能。在实验部分，我们将展示当与动态缓存绕过组件结合时，它可以进一步增加性能收益。与现有的LRR和GTO warp调度策略相比，我们提出的BAWS不产生额外的调度和硬件开销。CHSS指标以及warp的ba和bg标签由动态绕过组件维护。

A4 实验环境

TABLE II GPGPU-SIM 配置

TABLE III 评估的缓存敏感应用

A4 实验结果

A. 性能和能效结果
B. 旁路感知Warp调度器（BAWS）
C. 静态缓存绕过方法的比较
D. 与最先进技术的比较
E. 与线程节流技术（Optimal-TLP）的结合
F. 对缓存大小的敏感性分析
G. 集中式与分散式控制的比较

A7 补充细节

VIII. 相关工作

GPU优化技术概览。最先进的GPU优化技术包括性能建模【21, S. Hong and H. Kim, “An analytical model for a GPU architecture with memory-level and thread-level parallelism awareness,” in Proc. Int. Symp. Comput. Archit. (ISCA), Austin, TX, USA, 2009, pp. 152–163.】-【23, S. S. Baghsorkhi, M. Delahaye, S. J. Patel, W. D. Gropp, and W.-M. W. Hwu, “An adaptive performance modeling tool for GPU architectures,” in Proc. ACM Symp. Principles Pract. Parallel Program. (PPoPP), Bengaluru, India, 2010, pp. 105–113.】，控制流分歧优化【24, W. W. L. Fung, I. Sham, G. Yuan, and T. M. Aamodt, “Dynamic warp formation and scheduling for efficient GPU control flow,” in Proc. Int. Symp. Microarchit. (MICRO), Chicago, IL, USA, 2007, pp. 407–420.】-【26, Y. Liang, M. T. Satria, K. Rupnow, and D. Chen, “An accurate GPU performance model for effective control flow divergence optimization,” IEEE Trans. Comput.-Aided Design Integr. Circuits Syst., vol. 35, no. 7, pp. 1165–1178, Jul. 2016.】，多任务处理【27, J. T. Adriaens, K. Compton, N. S. Kim, and M. J. Schulte, “The case for GPGPU spatial multitasking,” in Proc. IEEE Int. Symp. High Perform. Comput. Archit. (HPCA), New Orleans, LA, USA, Feb. 2012, pp. 1–12.】-【30, X. Li and Y. Liang, “Efficient kernel management on GPUs,” in Proc. Design Autom. Test Europe (DATE), Dresden, Germany, 2016, pp. 85–90.】，数据布局转换【31, Y. Liang et al., “Real-time implementation and performance optimization of 3D sound localization on GPUs,” in Proc. Design Autom. Test Europe (DATE), Dresden, Germany, 2012, pp. 832–835.】，以及片上内存设计【2, M. Gebhart, S. W. Keckler, B. Khailany, R. Krashinsky, and W. J. Dally, “Unifying primary cache, scratch, and register file memories in a throughput processor,” in Proc. 45th Annu. IEEE/ACM Int. Symp. Microarchit. (MICRO), Vancouver, BC, Canada, 2012, pp. 96–106.】、【32, M. Gebhart, S. W. Keckler, and W. J. Dally, “A compile-time managed multi-level register file hierarchy,” in Proc. IEEE/ACM Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 465–476.】-【34, S. Wang et al., “Performance-centric register file design for GPUs using racetrack memory,” in Proc. 21st Asia South Pac. Design Autom. Conf. (ASP DAC), Macau, China, 2016, pp. 85–90.】。随着越来越多具有不规则内存访问模式的应用被移植到GPU，内存系统优化变得日益重要。线程调度策略【7, V. Narasiman et al., “Improving GPU performance via large warps and two-level warp scheduling,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 308–317.】、【35, M. Gebhart et al., “Energy-efficient mechanisms for managing thread context in throughput processors,” in Proc. 38th Annu. Int. Symp. Comput. Archit. (ISCA), San Jose, CA, USA, 2011, pp. 235–246.】、【36, T. G. Rogers, M. O’Connor, and T. M. Aamodt, “Divergence-aware warp scheduling,” in Proc. 46th Annu. IEEE/ACM Int. Symp. Microarchit. (MICRO), Davis, CA, USA, 2013, pp. 99–110.】和并行性管理技术【1, O. Kayıran, A. Jog, M. T. Kandemir, and C. R. Das, “Neither more nor less: Optimizing thread-level parallelism for GPGPUs,” in Proc. Int. Conf. Parallel Archit. Compilation Techn. (PACT), Edinburgh, U.K., 2013, pp. 157–166.】、【8, T. G. Rogers, M. O’Connor, and T. M. Aamodt, “Cache-conscious wavefront scheduling,” in Proc. Int. Symp. Microarchitect. (MICRO), Vancouver, BC, Canada, 2012, pp. 72–83.】已被讨论。内存调度策略【7, V. Narasiman et al., “Improving GPU performance via large warps and two-level warp scheduling,” in Proc. Int. Symp. Microarchit. (MICRO), Porto Alegre, Brazil, 2011, pp. 308–317.】和数据预取【37, A. Jog et al., “Orchestrated scheduling and prefetching for GPGPUs,” in Proc. 40th Annu. Int. Symp. Comput. Archit. (ISCA), Tel Aviv, Israel, 2013, pp. 332–343.】也得到了研究。

CPU上的缓存绕过。缓存绕过是缓解缓存争用问题的有效技术。针对CPU，静态【12, Y. Wu et al., “Compiler managed micro-cache bypassing for high performance EPIC processors,” in Proc. 35th Annu. ACM/IEEE Int. Symp. Microarchit. (MICRO), Istanbul, Turkey, 2002, pp. 134–145.】和动态方法【38, J. Gaur, M. Chaudhuri, and S. Subramoney, “Bypass and insertion algorithms for exclusive last-level caches,” in Proc. Int. Symp. Comput. Archit. (ISCA), San Jose, CA, USA, 2011, pp. 81–92.】-【41, H. Liu, M. Ferdman, J. Huh, and D. Burger, “Cache bursts: A new approach for eliminating dead blocks and increasing cache efficiency,” in Proc. 41st Annu. IEEE/ACM Int. Symp. Microarchit. (MICRO), 2008, pp. 222–233.】都已被提出。CPU的技术主要使用缓存命中率作为缓存绕过的指导模型。然而，在GPU上，由于其独特的架构特性，包括大规模并行性、资源拥塞和内存分歧，基于缓存命中率的模型并不总能表现良好【5, W. Jia, K. A. Shaw, and M. Martonosi, “Characterizing and improving the use of demand-fetched caches in GPUs,” in Proc. Int. Conf. Supercomput. (ICS), 2012, pp. 15–24.】、【20, X. Chen et al., “Adaptive cache management for energy-efficient GPU computing,” in Proc. 47th Annu. IEEE/ACM Int. Symp. Microarchit. (MICRO), Cambridge, U.K., 2014, pp. 343–355.】。

GPU上的缓存绕过。最近的一些研究已经探索了GPU的缓存绕过。静态缓存绕过技术在【5, W. Jia, K. A. Shaw, and M. Martonosi, “Characterizing and improving the use of demand-fetched caches in GPUs,” in Proc. Int. Conf. Supercomput. (ICS), 2012, pp. 15–24.】、【10, X. Xie, Y. Liang, G. Sun, and D. Chen, “An efficient compiler framework for cache bypassing on GPUs,” in Proc. Int. Conf. Comput.-Aided Design (ICCAD), San Jose, CA, USA, 2013, pp. 516–523.】和【42, Y. Liang, X. Xie, G. Sun, and D. Chen, “An efficient compiler framework for cache bypassing on GPUs,” IEEE Trans. Comput.-Aided Design Integr. Circuits Syst., vol. 34, no. 10, pp. 1677–1690, Oct. 2015.】中被提出，它们在编译时分析加载的数据访问模式。动态方法也得到了讨论。Jia等人【11, W. Jia, K. A. Shaw, and M. Martonosi, “MRPB: Memory request prioritization for massively parallel processors,” in Proc. 20th Int. Symp. High Perform. Comput. Archit. (HPCA), Orlando, FL, USA, 2014, pp. 272–283.】利用绕过来改善缓存性能，绕过由关联度停顿触发。Li等人【43, C. Li et al., “Locality-driven dynamic GPU cache bypassing,” in Proc. 29th ACM Int. Conf. Supercomput. (ICS), Irvine, CA, USA, 2015, pp. 67–77.】提出了一种由局部性驱动的动态缓存绕过技术，利用局部性信息在运行时调整缓存行为。Ausavarungnirun等人【44, R. Ausavarungnirun et al., “Exploiting inter-warp heterogeneity to improve GPGPU performance,” in Proc. Int. Conf. Parallel Archit. Compilation Techn. (PACT), San Francisco, CA, USA, 2015, pp. 25–38.】提出利用缓存绕过来调解内存分歧问题。Li等人【45, D. Li et al., “Priority-based cache allocation in throughput processors,” in Proc. IEEE 21st Int. Symp. High Perform. Comput. Archit. (HPCA), Burlingame, CA, USA, Feb. 2015, pp. 89–100.】提出了一种缓存管理策略，只允许一部分warp向缓存中插入数据。这些技术要么只采用静态缓存绕过，要么只采用动态缓存绕过，而我们的协同绕过则结合了两者。最近，Chen等人【20, X. Chen et al., “Adaptive cache management for energy-efficient GPU computing,” in Proc. 47th Annu. IEEE/ACM Int. Symp. Microarchit. (MICRO), Cambridge, U.K., 2014, pp. 343–355.】通过结合基于保护距离的缓存绕过和线程节流来提出一种自适应缓存管理技术以改善缓存性能。与我们的发现类似，他们也证明了与纯线程节流技术【8, T. G. Rogers, M. O’Connor, and T. M. Aamodt, “Cache-conscious wavefront scheduling,” in Proc. Int. Symp. Microarchitect. (MICRO), Vancouver, BC, Canada, 2012, pp. 72–83.】相比，缓存绕过能有效缓解缓存压力，同时对TLP的惩罚较小。相比之下，我们的协同缓存绕过在不进行线程节流的情况下利用缓存绕过。此外，我们提出了一个BAWS，以最大化缓存绕过的性能收益。

A5 结论

由于其巨大的计算能力，GPU作为性能加速器变得越来越重要。现代GPU已经采用缓存来改善具有不规则内存访问模式的通用应用程序的内存性能。然而，由于大规模线程并行性引起的缓存争用和资源拥塞，GPU应用程序常常无法从缓存中受益。本文提出了一种协同的静态和动态缓存绕过方法以及一种旁路感知warp调度器（BAWS）。在编译时，协同绕过将加载指令分为三类。在运行时，协同绕过动态调整使用缓存的线程块比例，以减少缓存争用和流水线停顿。旁路感知warp调度器通过监控缓存性能来决定使用或绕过缓存的warp的调度。实验证明，协同缓存绕过为各种GPU应用带来了平均1.32倍的性能加速。当与提出的BAWS结合时，性能收益从1.32倍提高到1.38倍。

引用文献汇总