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 实验环境

  • 硬件配置：

    • GPU模拟器：基于GPGPU-Sim（版本3.2.1）进行性能评估，使用GPUWattch进行功耗测量。
    • GPU架构（如表II所示）：15个计算单元（SM），每个SM配置32个核心，频率为700MHz。
    • SM资源：每个SM最多支持1536个线程、8个线程块、48KB共享内存和32K寄存器。
    • 缓存配置：
      • L1数据缓存：大小为16/32/48KB，采用LRU替换策略，设有32个MSHR条目。
      • L2统一缓存：6个bank，每个bank 128KB，16路组相联。
    • 调度器：每个SM有两个调度器，默认采用LRR（松散轮询）策略。

  • 软件配置：

    • 代码实现：本文提出的协同缓存绕过和warp调度策略在GPGPU-Sim模拟器中实现。
    • 数据集/基准：使用了来自Parboil、Rodinia、PolyBench和MapReduce的多个应用程序。实验重点关注对缓存敏感的应用，这些应用在之前最先进的GPU缓存管理技术研究中也被广泛使用。
    • 实验设置：如表III所示，实验中使用了不同的问题规模作为性能剖析（Profiling）和最终评估（Evaluation）的输入，以验证方法的普适性。默认的静态绕过方法采用基于性能剖析的分类。

TABLE II GPGPU-SIM 配置

TABLE III 评估的缓存敏感应用

A4 实验结果

A. 性能和能效结果
  • 实验内容：在16-kB L1缓存配置下，评估了仅静态绕过、仅动态绕过和协同缓存绕过的性能提升，基线为所有全局加载都使用L1缓存的默认设置。
  • 实验结果：
    • 性能（图9）：仅静态绕过、仅动态绕过和协同绕过分别实现了平均1.13倍、1.28倍和1.32倍的性能加速。协同绕过在应用PVC上取得了最高的2.28倍加速。
    • 性能提升来源（图10, 图11）：协同绕过通过两种方式提升性能：1) 缓解缓存争用，平均提升23.6%的缓存命中率；2) 减少流水线停顿，平均减少36.5%的L1缓存拥塞导致的停顿。对于本身局部性差的应用（如HIS, CVR），性能提升主要来自停顿减少；对于局部性好但受争用影响的应用（如SPV, PVC），性能提升来自命中率提高。
    • 能效（图13, 图14）：协同绕过平均提升了9.6%的能效（Perf/Watt）。虽然该技术导致平均功耗增加了21.3%（主要由于L2缓存和NoC活动更频繁），但执行时间的大幅缩短带来了整体能效的提升。

      图 9. 协同静态和动态缓存绕过对缓存敏感应用在16-kB缓存上的性能结果。

      图 10. 协同绕过的缓存命中率提升（16 kB）。

      图 11. 协同缓存绕过对L1缓存拥塞导致的流水线停顿的相对减少（16 kB）。

      图 13. 协同缓存绕过的归一化能效（Perf /Watt）（16-kB L1数据缓存）。

      图 14. 应用PVC的功耗分解。y轴为对数刻度。
B. 旁路感知Warp调度器（BAWS）
  • 实验内容：将协同缓存绕过与不同的warp调度策略（LRR, GTO, Two-level, BAWS）相结合，评估其性能。
  • 实验结果（图12）：协同绕过分别与LRR, GTO, Two-level, 和BAWS结合，平均性能加速为1.32倍, 1.31倍, 1.30倍, 和1.38倍。
  • 分析结论：传统的GTO和Two-level调度器在协同绕过框架下性能提升不明显，因为缓存争用问题已被缓解。而BAWS作为一种补充技术，通过在运行时感知缓存性能（命中和停顿）来动态调整调度策略，能够进一步释放性能潜力，将平均加速提升至1.38倍。

    图 12. 协同缓存绕过与LRR、GTO、两级调度和BAWS在16-kB缓存下的性能结果。
C. 静态缓存绕过方法的比较
  • 实验内容：比较基于性能剖析（profile-based）和基于内存模式分析（pattern analysis）的两种静态分类方法在协同绕过框架下的性能。
  • 实验结果（图15）：使用内存模式分析和性能剖析的协同绕过分别取得了1.30倍和1.32倍的平均性能加速。
  • 分析结论：内存模式分析方法的性能与更耗时的性能剖析方法非常接近，且无需任何剖析开销。这使得它在无法进行剖析的动态系统中具有很强的实用性。

    图 15. 采用不同静态绕过技术的协同缓存绕过性能结果。
D. 与最先进技术的比较
  • 实验内容：将本文提出的协同绕过技术与多种最先进的GPU缓存优化技术进行比较，包括指令级绕过、访问级绕过、两级调度器、最优线程并行度（Optimal-TLP）和缓存感知波前调度（CCWS-Best-SWL）。
  • 实验结果（图16）：协同绕过的平均性能加速为1.32倍，显著优于其他技术（分别为1.11倍, 1.17倍, 1.04倍, 1.09倍, 1.15倍）。
  • 分析结论：协同绕过技术的优势在于它结合了静态全局分析和动态运行时调整，且不以牺牲线程级并行度（TLP）为代价，从而在性能上超越了仅依赖静态或动态信息、或通过限制并行度来缓解争用的方法。

    图 16. 与最先进优化技术的性能比较。
E. 与线程节流技术（Optimal-TLP）的结合
  • 实验内容：评估将协同缓存绕过与线程节流技术结合（称为“Throttling + Bypassing”，通过穷举搜索确定最优的线程块数和绕过比例）的性能潜力。
  • 实验结果（图17）：结合线程节流后，平均性能加速从1.32倍提升至1.39倍。
  • 分析结论：线程节流是协同绕过的一种补充技术，尤其对于那些除了L1缓存外还存在NoC和DRAM拥塞的应用。本文的协同绕过技术本身已能达到接近这种组合优化上限的性能。

    图 17. 结合缓存绕过与线程节流的性能结果。
F. 对缓存大小的敏感性分析
  • 实验内容：在16-kB、32-kB和48-kB三种不同大小的L1缓存上评估协同绕过的性能。
  • 实验结果（图18）：随着缓存容量的增加，性能加速逐渐减小，三种配置下的平均加速分别为1.32倍、1.21倍和1.11倍。
  • 分析结论：更大的缓存本身能缓解缓存争用，因此协同绕过的收益会降低。然而，大缓存会带来更高的访问延迟、功耗和面积成本。对于固有时间局部性差的应用，即使大缓存也无效，而缓存绕过依然能通过减少流水线停顿来提升性能。

    图 18. 不同缓存大小下的性能结果。
G. 集中式与分散式控制的比较
  • 实验内容：比较动态绕过中的两种控制机制：集中式控制（由一个SM决策并广播）和分散式控制（每个SM独立决策）。
  • 实验结果：平均而言，集中式控制比分散式控制性能略好3.2%。
  • 分析结论：集中式控制由于学习成本较低，表现更优。

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倍。

引用文献汇总

  • 【1】 O. Kayıran 等人, "Neither more nor less: Optimizing thread-level parallelism for GPGPUs", PACT, 2013.
    • 引用位置：引言，实验章节。
    • 引用内容描述：指出GPU应用性能常受内存子系统限制，且最大化线程级并行度可能因缓存争用导致性能下降。
  • 【2】 M. Gebhart 等人, "Unifying primary cache, scratch, and register file memories in a throughput processor", MICRO, 2012.
    • 引用位置：引言，相关工作。
    • 引用内容描述：讨论片上内存设计，作为GPU优化的相关工作。
  • 【3】 S. Che 等人, "Rodinia: A benchmark suite for heterogeneous computing", IISWC, 2009.
    • 引用位置：背景知识，实验章节。
    • 引用内容描述：作为实验评估所使用的基准测试集之一。
  • 【4】 J. A. Stratton 等人, The IMPACT Research Group, Parboil Benchmark Suite.
    • 引用位置：背景知识，实验章节。
    • 引用内容描述：作为实验评估所使用的基准测试集之一。
  • 【5】 W. Jia 等人, "Characterizing and improving the use of demand-fetched caches in GPUs", ICS, 2012.
    • 引用位置：引言，背景知识，实验章节，相关工作。
    • 引用内容描述：指出在某些情况下GPU缓存会损害性能，并作为指令级绕过的代表性工作进行比较。
  • 【6】 S. Mu 等人, "Orchestrating cache management and memory scheduling for GPGPU applications", IEEE Trans. Very Large Scale Integr. (VLSI) Syst., 2014.
    • 引用位置：方法细节。
    • 引用内容描述：提及内存分歧问题，解释为何在线程块级别进行绕过决策。
  • 【7】 V. Narasiman 等人, "Improving GPU performance via large warps and two-level warp scheduling", MICRO, 2011.
    • 引用位置：引言，系统概述，方法细节，实验章节，相关工作。
    • 引用内容描述：指出warp调度对GPU性能有重要影响，并作为两级调度器的代表性工作进行比较。
  • 【8】 T. G. Rogers 等人, "Cache-conscious wavefront scheduling", MICRO, 2012.
    • 引用位置：引言，方法细节，实验章节，相关工作。
    • 引用内容描述：提出缓存感知的warp调度来提升块内缓存局部性，作为CCWS的代表性工作进行比较。
  • 【9】 X. Xie 等人, "Coordinated static and dynamic cache bypassing for GPUs", HPCA, 2015.
    • 引用位置：引言，方法细节。
    • 引用内容描述：本文的会议初版。
  • 【10】 X. Xie 等人, "An efficient compiler framework for cache bypassing on GPUs", ICCAD, 2013.
    • 引用位置：引言，方法细节，实验章节，相关工作。
    • 引用内容描述：作为指令级绕过的代表性工作进行比较。
  • 【11】 W. Jia 等人, "MRPB: Memory request prioritization for massively parallel processors", HPCA, 2014.
    • 引用位置：引言，背景知识，实验章节，相关工作。
    • 引用内容描述：描述了缓存未命中时对MSHR和缓存行的资源需求，并作为访问级绕过的代表性工作进行比较。
  • 【12】 Y. Wu 等人, "Compiler managed micro-cache bypassing for high performance EPIC processors", MICRO, 2002.
    • 引用位置：背景知识，相关工作。
    • 引用内容描述：指出缓存绕过对CPU缓存争用问题的有效性。
  • 【13】 S. M. Khan 等人, "Sampling dead block prediction for last-level caches", MICRO, 2010.
    • 引用位置：背景知识。
    • 引用内容描述：作为CPU上用于末级缓存（LLC）的缓存绕过的例子。
  • 【14】 C.-J. Wu 等人, "Ship: Signature-based hit predictor for high performance caching", MICRO, 2011.
    • 引用位置：背景知识。
    • 引用内容描述：作为CPU上用于末级缓存（LLC）的缓存绕过的例子。
  • 【15】 A. Bakhoda 等人, "Analyzing CUDA workloads using a detailed GPU simulator", ISPASS, 2009.
    • 引用位置：系统概述，实验章节。
    • 引用内容描述：作为本文实验所使用的GPGPU-Sim模拟器。
  • 【16】 A. Jog 等人, "OWL: Cooperative thread array aware scheduling techniques for improving GPGPU performance", ASPLOS, 2013.
    • 引用位置：系统概述。
    • 引用内容描述：指出默认的LRR调度器效率低下，可能加剧缓存争用。
  • 【17】 S. Grauer-Gray 等人, "Auto-tuning a high-level language targeted to GPU codes", InPar, 2012.
    • 引用位置：实验章节。
    • 引用内容描述：作为实验评估所使用的PolyBench基准测试集的来源。
  • 【18】 B. He 等人, "Mars: A mapreduce framework on graphics processors", PACT, 2008.
    • 引用位置：实验章节。
    • 引用内容描述：作为实验评估所使用的MapReduce应用的来源。
  • 【19】 J. Leng 等人, "GPUWattch: Enabling energy optimizations in GPGPUs", ISCA, 2013.
    • 引用位置：实验章节。
    • 引用内容描述：作为本文实验所使用的GPU功耗测量工具。
  • 【20】 X. Chen 等人, "Adaptive cache management for energy-efficient GPU computing", MICRO, 2014.
    • 引用位置：实验章节，相关工作。
    • 引用内容描述：提出结合缓存绕过和线程节流，并指出缓存绕过对TLP的惩罚较小。
  • ... (其余参考文献主要用于相关工作章节的背景介绍) ...