文章标题：面向现代GPU的局部性感知CTA聚类
作者/机构：
- Ang Li, Pacific Northwest National Lab
- Shuaiwen Leon Song, Pacific Northwest National Lab
- Weifeng Liu, University of Copenhagen
- Xu Liu, College of William and Mary
- Akash Kumar, Technische Universität Dresden
- Henk Corporaal, Eindhoven University of Technology

A1 主要贡献

本文探讨了现代GPU中一种长期被忽视但具有性能提升潜力的数据局部性——CTA间局部性（inter-CTA locality）。利用这种局部性面临三大挑战：硬件可行性不明确、底层的CTA调度器未知且无法访问、以及核心内缓存容量小。

为解决这些问题，本文做出了以下主要贡献：
1. 实证证明了CTA间局部性的可利用性：通过在多种现代GPU（Fermi、Kepler、Maxwell、Pascal）上进行的实证探索，证明了在L1或L1/Tex统一缓存上，可以利用空间和时间上的CTA间局部性。（第3.1节）
2. 量化并分类了CTA间局部性：发现CTA间重用占全局数据重用的重要部分，并根据其来源将其分为算法相关、缓存行相关、数据相关、写相关和流式五类，讨论了其可利用性。（第3.2节）
3. 提出了CTA聚类（CTA-Clustering）的概念和方法：提出了一种旨在将具有潜在重用的CTA分组到同一个SM上并发或连续执行的软件技术。详细阐述了其理论和包含分区（Partitioning）、反转（Inverting）和绑定（Binding）三个步骤的设计方案。（第4.2-4.3节）
4. 构建了自动优化框架：将所提出的技术整合到一个自动化框架中，用于自动利用通用应用程序的CTA间局部性。该框架可作为编译器的一部分，直接部署在现有商用GPU上，无需修改硬件。（第4.4节）
5. 在多代GPU上验证了有效性：在所有现代NVIDIA GPU架构上进行了广泛评估。实验结果表明，对于具有算法相关CTA间重用的应用，该技术通过平均减少55%（Fermi）、65%（Kepler）、29%（Maxwell）和28%（Pascal）的L2缓存事务，分别带来了平均1.46倍、1.48倍、1.45倍和1.41倍的性能加速，最高可达3.8倍、3.6倍、3.1倍和3.3倍。（第5节）

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

第2节背景知识

总体架构。GPU处理器由多个称为流式多处理器（Streaming Multiprocessors, SMs）的SIMD核心组成。图1展示了现代NVIDIA GPU的通用架构图。通常，一个SM包含标量处理器、寄存器文件、特殊功能单元、加载/存储单元、共享内存和不同类型的缓存。表1列出了四种代表性产品的基本架构规格。Fermi和Kepler GPU的每个SM都有一个48KB可配置、128B缓存行、写驱逐（write-evict）的L1数据缓存，用于通用片外DRAM访问，并与共享内存共享同一片上存储。然而，近期的Maxwell和Pascal GPU将该存储完全用于共享内存，而依赖于纹理缓存（Texture cache，32B缓存行，非相干）提供L1缓存能力，称为L1/Tex统一缓存【28，NVIDIA. GTX980 Whitepaper: Featuring Maxwell, the Most Advanced GPU Ever Made, 2014】。同时，GPU中所有的SM通过一个片上网络（NoC）连接到一个共享的L2缓存。L2缓存是分块（banked）且可写的（写回和写分配【29，Inderpreet Singh, et al. Cache coherence for GPU architectures. 2013, HPCA】），缓存行大小为32B。传统上GPU不启用硬件缓存预取【30，Sara S. Baghsorkhi, et al. Efficient Performance Evaluation of Memory Hierarchy for Highly Multithreaded Graphics Processors. 2012, PPoPP】。但用户可以通过特定的编译器选项【7，NVIDIA. CUDA Programming Guide, 2015】启用/禁用L1缓存。值得注意的是，L1缓存行大小大于或等于L2，这对后续讨论很重要。

图1：现代NVIDIA GPU架构图：Fermi, Kepler, Maxwell 和 Pascal。箭头表示L1和L2数据缓存的不同全局内存读取数据路径。

执行模型。GPU SM遵循单指令多线程（Single-Instruction-Multiple-Threads, SIMT）执行模型【6，Erik Lindholm, et al. NVIDIA Tesla: A unified graphics and computing architecture. 2008, IEEE Micro】。一组32个线程构成一个执行向量，称为warp。Warp是SM指令获取和解码的基本单位，一个warp内的所有线程以锁步方式前进。Warp在SM的warp插槽中注册（见表1）。多个warp构成一个称为线程块（thread block）或协作线程阵列（Cooperative-Thread-Array, CTA）的块，它封装了所有的线程同步和屏障操作。CTA是向SM交付作业的基本单位。从硬件角度看，CTA之间不应有依赖关系——一个内核在任意CTA调度/执行顺序下都应获得正确结果。这一特性确保了当CTA调度策略和/或SM架构被修改时，应用程序保持不变。这种“无序”特性对本文也至关重要：可以操纵CTA调度而不会危及一致性。

CTA调度。GPU上的默认CTA调度策略被认为是轮询（Round-Robin, RR）【11，Adwait Jog, et al. OWL: Cooperative Thread Array Aware Scheduling Techniques for Improving GPGPU Performance. 2013, ASPLOS】、【27，Minseok Lee, et al. Improving GPGPU resource utilization through alternative thread block scheduling. 2014, HPCA】、【31，Onur Kayıran, et al. Neither more nor less: optimizing thread-level parallelism for GPGPUs. 2013, PACT】、【32，Hyeran Jeon, et al. CTA-aware Prefetching for GPGPU. 2014, CENG-2014-08】、【33，Jin Wang, et al. LaPerm: Locality Aware Scheduler for Dynamic Parallelism on GPUs. 2016, ISCA】。首先，CTA调度器（即GigaThread引擎）为每个SM分配至少一个CTA。如果一个SM仍有足够的资源（如寄存器、共享内存、warp插槽等）来支持额外的CTA，则会进行第二轮分配。这个轮询分配过程会重复进行，直到所有SM都饱和，要么受资源限制，要么受硬件限制【32，Hyeran Jeon, et al. CTA-aware Prefetching for GPGPU. 2014, CENG-2014-08】。之后，每当一个现有的CTA完成退出，一个新的CTA就会被分配给一个SM。需要注意的是，CTA调度完全由硬件实现的GigaThread引擎管理，没有任何明确的软件策略可以影响默认的CTA调度，也无法修改CTA如何分派到SM。此外，一旦一个CTA被分配给一个SM，它就不能被抢占或重新分配到另一个SM【7，NVIDIA. CUDA Programming Guide, 2015】。这种对CTA调度器缺乏控制，为利用调度提升GPU性能带来了主要障碍【25，Bo Wu, et al. Enabling and exploiting flexible task assignment on GPU through SM-centric program transformations. 2015, ICS】、【26，Kunal Gupta, et al. A study of persistent threads style GPU programming for GPGPU workloads. 2012, InPar】。最后，RR调度终究是一个假设，因为GigaThread引擎中的实际调度算法从未被公开过【25，Bo Wu, et al. Enabling and exploiting flexible task assignment on GPU through SM-centric program transformations. 2015, ICS】。

第3节理解GPU上的CTA间重用

3.1 在L1上利用CTA间重用的可行性

设计微基准测试验证CTA间局部性。我们设计了一个微基准测试（如清单3所示），以验证GPU L1缓存是否有能力利用被分派到同一个SM上的CTA之间的数据重用，无论是同时执行（空间局部性）还是相继执行（时间局部性）。评估的GPU平台配置列于表1，包括Fermi、Kepler、Maxwell和Pascal。如代码18-21行所述，我们分别为这四个平台启动了480、960、1024、1280个CTA，对应每个SM 4、4、2、2个轮次（turnaround）。每个CTA包含一个warp，且基本上只利用主线程以避免CTA边界内的warp内合并和warp间冲突效应。由于每个CTA只有一个warp，一个SM的所有硬件CTA插槽（见表1）实际上都被占满。因此，在每个轮次中，一个SM可以在这四个平台上分别同时执行8、16、32、32个CTA。为了更准确地测量内存访问延迟，在Maxwell和Pascal上运行时，我们在第12行前后增加了两条同步指令，因为在Maxwell之后，计时指令（第10和14行）和内存访问指令（第12行）可以同时发出，导致计时结果倾斜。CTA屏障（_syncthreads()）用于禁止多发射。

实验观察与分析。图2显示了在两种设置下，分派到容纳第一个CTA（即CTA-0）的特定SM上的CTA的平均内存访问延迟：（A）默认CTA调度，和（B）交错执行。我们将这个SM表示为“SM 0”。x轴代表分派到SM 0的CTA-ID。我们还使用CUDA分析器分析了另外两个缓存指标（即L1读取事务数和L1到L2读取事务数）。总结观察结果如下：
1. 时间局部性：图2-(A)表明，L1级别可以利用后续轮次的CTA（即在一个CTA退出后）的时间局部性。这一点通过以下观察得到证实：只有第一轮的CTA会经历由全局内存加载引起的高访问延迟（例如，在Pascal上约为750个周期）。所有后续的CTA都受益于第一轮在L1中产生的局部性，因此它们的访问延迟与L1访问延迟相似（例如，在Pascal上约为132个周期）。请注意，在第一轮的多个CTA中，实际上只有一个或两个从片外DRAM取数据；其他CTA尽管在L1缓存中命中，但实际上是“命中保留”（hit reserved）【21，Ali Bakhoda, et al. Analyzing CUDA workloads using a detailed GPU simulator. 2009, ISPASS】（命中但请求的数据仍在传输中），因此表现出同样长的延迟。分析器报告的L1读取事务数和L1到L2读取事务数也证实了这一点：对于Fermi和Kepler，一个128B的L1未命中等同于四个32B的L2读取事务；而对于Maxwell和Pascal，L1/Tex统一缓存被划分为两个扇区（sector），每个扇区产生一个32B的L1未命中，导致两个L2读取事务。据推测，这些扇区是私有于特定的CTA插槽，遵循某种映射机制。
2. 空间局部性：为了验证L1缓存是否可以利用同时分派到同一个SM的CTA（即在同一轮次中）的空间局部性，我们启用了交错执行（清单3中的第7-9行）来错开它们的内存访问，这样并发CTA的同步内存请求就不会像之前那样在L1中被聚合。第8行的DELAY变量控制交错的程度，设置得足够长（例如1200个周期），以便前一个CTA请求的数据可以在同一轮次中其他CTA获取数据之前到达L1。这也使所有并发的CTA保持存活和活动状态。图2-(B)展示了同时启动的CTA之间的这种空间重用。例如，只有第一轮中的第一个CTA表现出比其余CTA长得多的延迟。
3. 观察到的硬件CTA调度策略：通过在表1中的四种GPU以及GTX750Ti（第一代Maxwell GPU，CC-5.0）上的实验，我们发现默认的底层CTA调度器在不同架构间有所不同。我们基于微基准测试的运行观察到两种通用的调度模式：（1）对于表1中的架构，第一轮的CTA通常遵循轮询（RR），但其余轮次的CTA则不然，它们遵循一种更“需求驱动”的策略。（2）在GTX750Ti上，每个轮次内的CTA被随机分配给SM 0，而不是遵循任何特定规则。对真实世界GPU应用程序的进一步调查证实，默认的硬件CTA调度器实际上接近模式（2），即第一轮也不一定遵循RR。我们还发现，即使SM数量能整除CTA数量，工作负载在SM间的分布也不均衡。例如，在图2-(A)的Kepler GPU上，SM 0只执行了60个CTA，而不是预期的64个。

图2：在持有CTA-0的SM上利用CTA间重用：（A）默认场景（L1上的时间CTA间局部性）；（B）交错场景（L1上的空间CTA间局部性）。

3.2 CTA间局部性的量化与来源

CTA间局部性的量化。我们量化了CTA间重用在总全局数据重用中所占的百分比，这个计算基于SM产生的所有内存请求在进入L1缓存之前的数据重用情况。由于当前的GPU分析器无法追踪特定内存请求的地址和来源（如CTA-ID），我们仅为演示目的，使用GPGPU-Sim【21，Ali Bakhoda, et al. Analyzing CUDA workloads using a detailed GPU simulator. 2009, ISPASS】来追踪所有内存访问请求的数据重用，并估算CTA间重用在整体数据重用中的百分比。需要注意的是，这个估算是数据驱动的，与缓存设计或CTA调度策略无关。每次运行的结果都是一致的。图3量化了33个常见GPU应用程序的CTA间重用情况，清楚地表明这些应用程序中的CTA间重用占整体数据重用的非常重要的部分（平均为45%）。这表明旨在改善CTA间重用的策略可能会带来显著的整体性能效益，超越了仅试图促进warp内和warp间数据重用的传统方法【16，R. Ausavarungnirun, et al. Exploiting Inter-Warp Heterogeneity to Improve GPGPU Performance. 2015, PACT】、【34，Veynu Narasiman, et al. Improving GPU performance via large warps and two-level warp scheduling. 2011, MICRO】、【35，A. Sethia, et al. Mascar: Speeding up GPU warps by reducing memory pitstops. 2015, HPCA】、【36，Adwait Jog, et al. Orchestrated Scheduling and Prefetching for GPGPUs. 2013, ISCA】。

CTA间局部性的来源分类。然而，这些显著的CTA间局部性是否能在GPU硬件上被利用以提升性能？为了回答这个问题，我们需要理解CTA间局部性的来源。通过对广泛的GPU应用程序（详见表2）进行特征分析，我们将GPU CTA间局部性的来源分为以下五个类别，其模式如图4所示：
- (A) 算法相关（Algorithm Related）：其CTA间局部性源于特定的算法设计，其中某些数据被来自不同CTA的线程多次使用（图4-(A)）。对于这些应用，算法设计者的提示对性能优化很重要。此外，这些应用通常为CTA间重用提供了很好的机会。典型例子包括MM、KMN、DCT。
- (B) 缓存行相关（Cache-line Related）：其CTA间局部性由GPU缓存设计引入，更具体地说是长缓存行尺寸【17，Lingda Li, et al. Tag-Split Cache for Efficient GPGPU Cache Utilization. 2016, ICS】、【27，Minseok Lee, et al. Improving GPGPU resource utilization through alternative thread block scheduling. 2014, HPCA】。如表1所示，在Fermi和Kepler上，一个线程的单个整数访问（4B）导致的缓存未命中必须将整个128B的缓存行取入L1。因此，该缓存行中的其他31个整数可能被来自不同CTA的线程访问（图4-(B)）。当内存访问行为不完全合并或未与缓存行边界完全对齐时（例如，线程访问光环区域或用户定义的对象数组），这种情况就会发生，并且在具有大L1缓存行尺寸的架构中（如Fermi和Kepler）尤其普遍。典型例子包括SYK、NBO、ATX。
- (C) 数据相关（Data Related）：此类应用主要处理不规则数据结构，如图、树、哈希、指针列表等。CTA间局部性来自数据在存储中的组织方式，或数据在内存中的访问方式（图4-(C)）。由于数据组织的非规律性，这种局部性通常是偶然实现的。典型的数据相关应用包括BFS、HST、BTR。
- (D) 写相关（Write Related）：此类应用可能具有CTA间局部性。然而，由于GPU L1缓存采用写驱逐（write-evict）策略【7，NVIDIA. CUDA Programming Guide, 2015】、【29，Inderpreet Singh, et al. Cache coherence for GPU architectures. 2013, HPCA】，可能被重用的数据可能会被另一个CTA对同一缓存行的不相关写入提前驱逐（图4-(D)）。当一个内核读写同一个数组，但访问距离小于缓存行大小时（例如，包含a[i + 1]的缓存行在另一个CTA写入a[i]时被驱逐），就会发生这种情况。NW是一个典型的写相关应用。
- (E) 流式（Streaming）：流式应用的内存访问大多是合并和对齐的（图4-(E)），而数据只使用一次或仅在CTA范围内重用（例如，通过共享内存）。这些应用很少有CTA间重用。典型的流式应用包括BS、SAD、DXT。

A2 方法细节

4.1 总体策略概述

基于可利用性的两种优化策略。如前所述，一些应用程序（如算法相关的）具有丰富的CTA间局部性，而另一些（如流式）则没有。根据应用程序是否具有可利用的CTA间局部性，我们提出了一种操纵GPU上CTA调度的软件方法，即CTA聚类（CTA-Clustering）。我们将应用程序是否具有“可利用的”CTA间局部性定义如下：
* 算法相关（程序定义）和缓存行相关（架构定义）的局部性可以在运行时之前识别，因此是可利用的。
* 数据相关（数据定义）、写相关（有局部性但无法利用）和流式（无局部性）的局部性要么不显著，要么只能在运行时确定。因此，利用它们的可能性很小或没有好处。我们认为它们没有可利用的CTA间局部性。
需要注意的是，对于某些具有非常特定运行时访问模式的数据相关应用，它们在内存中的数据组织可能可以在运行时之前预测【37，Lingda Li, et al. A Graph-based Model for GPU Caching Problems. 2016, arXiv:1605.02043】，从而使其CTA间局部性变得可利用。例如，先前的工作【38，Eddy Z. Zhang, et al. On-the-fly elimination of dynamic irregularities for gpu computing. 2011, ASPLOS】、【39，Jianqiao Liu, et al. Hybrid CPU-GPU scheduling and execution of tree traversals. 2016, ICS】建议在运行时前使用一个轻量级的检查器内核（inspector kernel）来分析某些图处理应用（如BFS的前几层）的局部数据访问，以预测全局数据组织，从而优化运行时访问。图5概述了在GPU上利用CTA间局部性的总体流程。对于具有可利用CTA间局部性的应用程序，我们应用CTA聚类来最大化它们在L1上的CTA间数据重用（第4.2节和4.3-(I)(II)节）。对于没有可利用CTA间局部性的应用程序，我们首先应用CTA聚类来重塑默认的CTA调度模式（不是为了利用CTA间局部性，而是为了施加一个特定的CTA执行顺序），然后使用CTA预取（CTA-Prefetching）在当前CTA结束前预加载后续CTA所需的数据（第4.3-(III)节）。

4.2 CTA聚类

CTA聚类的核心思想与步骤。CTA聚类的基本思想是启动一个新内核来替换原始内核，从而建立一个预定义的聚类规则。目标是将具有CTA间局部性的CTA聚集在一起，并在同一个SM上并发或连续执行。我们使用O、N、C分别表示原始内核、新内核和簇。CTA聚类本质上是找到一个映射：$N \rightarrow O$，该映射受制于C，如图6所示。这个过程包括三个步骤：
- 步骤1. 分区（Partitioning）：将O的CTA划分为M个平衡的簇C，使得每个簇内保留最多的CTA间局部性。这由图6中的函数f定义。
- 步骤2. 反转（Inverting）：从C重建O。换句话说，给定某个簇中的一个CTA，我们可以检索其在原始内核O中的CTA ID。反转由图6中的函数$f^{-1}$定义。
- 步骤3. 绑定（Binding）：将新内核N的CTA绑定到簇C。假设C的大小等于SM的数量，那么从C到SM的映射是一对一的映射。绑定由图6中的函数g定义。
简而言之，寻找映射 $N \rightarrow O$ 的过程如下：

以矩阵乘法（MM）为例。为了清晰地展示理论，我们使用一个著名的应用——来自CUDA SDK【40，NVIDIA. CUDA SDK Code Samples, 2015】的矩阵乘法（MM）来展示聚类过程。根据MM的源代码，其CTA内数据重用完全由共享内存处理（即，矩阵A和B的局部缓冲区在共享内存中声明，用于线程间的CTA内数据共享）——CTA间局部性没有被显式探索。然而，如图8-(A)所示，MM固有地具有算法相关的CTA间局部性，这可能被利用来提升性能。由于篇幅限制，我们不在此展示MM的源代码，本节中将称其为“MM内核体”。

图6：CTA聚类是找到 N → O (红色虚线) 或 $f^{-1}(g(N))$。

4.2.1 步骤1. CTA分区：$f = O \rightarrow C$

分区问题的形式化定义与启发式解法。分区问题的形式化定义是：给定一个无向图 $G(V, E_\mu)$，其中每个统一的顶点代表一个CTA，每条加权边代表度为 $\mu$ 的CTA间重用，将该图划分为M个平衡的簇，使得簇内边的权重总和最大化。我们将 $y_{v,i} = 1$ 定义为顶点v被分配到簇i（$i \in [M]$）；否则 $y_{v,i} = 0$。因此，我们可以将簇i定义为：$C_i = [(v, e)|v \in V, y_{v,i} = 1]$。分区问题可以表述为：对于所有顶点 $v \in V$，$\sum_{i=1}^{M} y_{v,i} = 1$（一个顶点分配到一个簇）；对于所有簇 $i \in [M]$，$\sum_{v \in V} y_{v,i} = |V|/M$（簇是平衡的）。一般的平衡图划分是NP完全问题【41，Konstantin Andreev and Harald Racke. Balanced graph partitioning. 2006, Theory of Computing Systems】。然而，应用开发者可以根据他们对算法的了解生成一个期望的划分函数。在这项工作中，我们提供一个启发式解决方案，作为自动化框架（第4.4节）的一部分来确保聚类的平衡性。我们将划分函数定义为：

其中v是O中的CTA ID，w是$C_i$中的ID。例如，如果一个簇包含3个CTA，则$w \in [0, 1, 2]$。由于V定义了O中CTA的顺序，而$|V|$是O中CTA的数量，我们的解决方案是通过以下方式将$|V|$分成M个块：

然而，$|V|$不一定是M的倍数。为了尽可能平衡地分配这些块，我们将公式3扩展为一个条件方程：如果 $v\%M \le |V|\%M$，则：

否则：

CTA索引方法与分区策略。与传统观念【27，Minseok Lee, et al. Improving GPGPU resource utilization through alternative thread block scheduling. 2014, HPCA】、【36，Adwait Jog, et al. Orchestrated Scheduling and Prefetching for GPGPUs. 2013, ISCA】不同，我们的划分方法不只是简单地将CTA分成连续的块，因为顶点v的顺序完全由CTA在O中的索引方式决定。图7展示了2D网格的四种主要CTA索引方法。对于默认的行主序索引，如果使用CUDA术语，v = blockIdx.y * gridDim.x + blockIdx.x，它会将连续的或行相邻的CTA聚类在一起，即沿着Y维度进行划分（Y-partitioning）。如果采用列主序索引（即v = blockIdx.x * gridDim.y + blockIdx.y），CTA将沿着X维度进行划分（X-partitioning）。此外，基于瓦片（Tile-based）的2D索引可以同时沿着X和Y维度进行划分，但可能因复杂的索引计算而产生更高的开销（第5.2节）。最后，也可以通过选择自定义索引方法来聚类任意的CTA。现在回顾第4.1节，我们描述了两种具有可利用CTA间局部性的应用程序的划分方法：
- (A) 算法相关：为了使划分过程通用和自动化，我们提出一个基于数组引用对不同网格维度坐标的依赖性分析的解决方案，这类似于对循环嵌套的依赖性分析【2，Randy Allen and Ken Kennedy. Optimizing compilers for modern architectures a dependence-based approach. 2001, Morgan Kaufmann】。如果内核网格（gridDim）是1D的，我们简单地执行X-partitioning。如果内核网格是2D的，且基于CTA X方向的变量（如blockIdx.x）是数组引用的最后一个或唯一的维度（如$A[\alpha(by) + bx + \epsilon(tx,ty)]$或$A[\beta(bx)]$），这表明它可能在X方向上具有CTA间局部性，我们则使用行主序索引进行Y-partitioning。否则，如果基于Y的变量（如blockIdx.y）是数组引用的最后一个或唯一的维度（如$A[\alpha(bx) + by + \epsilon(tx,ty)]$或$A[\beta(by)]$），我们则使用列主序索引进行X-partitioning。3D内核的划分过程与2D类似，但选项更多。然而，大多数常见的GPU应用程序只包含2D内核网格。
- (B) 缓存行相关：由于CUDA和C/C++通常采用行主序策略来组织和存储多维数组，与缓存行相关的CTA间局部性通常存在于行相邻的CTA之间。因此，我们使用行主序CTA索引（Y-partitioning）来计算O中的CTA顺序。

MM分区示例。如第4.1节所述，一些数据相关的应用如果用户提供了详细的自定义分区，也可以利用CTA间局部性，但这超出了本文的范围。以图8-(A)中的MM为例，A和B分别在S和T区域有CTA间重用。但是，是沿着Y（为了A的局部性）还是沿着X（为了B的局部性）划分CTA，取决于A的高度是否大于B的宽度（即方向性局部性强度）。在这种情况下，我们以A为目标，使用Y-partitioning（即行主序索引）。通过已知参数：M=2，网格宽度=3，网格高度=2，如图8-(A)所示，我们有$|V|=6$和$|V|\%M = 0$，这总是$\le v\%M$。因此，我们可以使用公式5来获得划分函数$f(v)$，并用它来定位O中任意CTA v应该被分派到哪个簇。例如，CTA-(0,1)可以被定位为簇C1中的第0个元素，通过计算$v = 1 * 3 + 0 = 3$，然后$f(3)=(3/2, 3\%2)=(1,1)$，再通过公式5计算w：$w = 3/2 - (1 * (6/2) + min(3\%2 - 1, 0)) = 1 - (3+0) = -2$ (原文此处似乎有误，应为 $w=\lfloor v/M \rfloor - (\lfloor (v\%M)/M \rfloor * (\lfloor |V|/M \rfloor + 1) + \lfloor v/M \rfloor * \lfloor |V|/M \rfloor ...)$，按论文公式5计算 $w = \lfloor 3/2 \rfloor - (1 * \lfloor 6/2 \rfloor + \min(3\%2 - 1, 0)) = 1 - (3+0) = -2$，这显然是错的，原文公式可能存在印刷错误或者我理解有偏差。根据意图，v=3，M=2，应该在cluster 1，w为1。我们遵循论文的意图，即CTA-(0,1) v=3，应映射到(w=1, i=1)。)

4.2.2 步骤2. CTA反转：$f^{-1} = C \rightarrow O$

反转函数的推导。由于公式3是一个一对一的映射函数，我们可以得到其反函数为：

类似地，公式4和5的反函数为：

这两个方程可以统一为：

MM反转示例。给定一对(w, i)（一个簇i和该簇中的位置w），它对应的CTA ID v可以通过公式7在O中定位。如果索引方法也已知，我们可以进一步获得它在O中的坐标。对于MM，一个在C1中标记为(w=2, i=1)的CTA可以通过以下方式找到其在O中对应的v：$v = f^{-1}((2, 1)) = 1 * (\lfloor6/2\rfloor + 0) + 2 + \min(1 - 0, 0) = 1*3+2+0=5$ (此处原文公式应用似乎也有问题，根据意图(w=2,i=1)应对应v=5，即坐标(1,2)，公式8: $v=1*(\lfloor 6/2 \rfloor+0)+2+min(1-0,0) = 5$，此计算正确)

4.2.3 步骤3. CTA绑定：$g = N \rightarrow C$

绑定问题的定义与两种绑定方法。绑定问题的形式化定义是：给定新内核N，如何将N中的CTA与簇C关联起来，以便C中的所有项都能被完整而精确地执行？此步骤的目标是找到从$N \rightarrow C$的映射，或者本质上是$N \rightarrow (w, i)$。换句话说，给定新内核N中的任意CTA u，如何获知它负责的目标任务(w, i)。由于GigaThread引擎是未知且不可访问的，根据在N的上下文中如何获得(w, i)，我们提出了两种方法来欺骗或规避硬件CTA调度器：
- (A) 基于RR的绑定（RR-based Binding）：这种简单的方法是基于GigaThread引擎在新内核中遵循严格的RR策略这一假设来获得N中的(w, i)。有了这个假设，我们可以通过以下方式计算$N \rightarrow C$的映射：

在MM中，对于一个ID为4的N中的CTA，其关联的(w, i)可以通过公式8计算：$(w,i)=(4/2, 4\%2)=(2,0)$。
- (B) 基于SM的绑定（SM-based Binding）：与基于RR的绑定不同，这种方法对默认的CTA调度策略没有任何预设。假设簇i被绑定到SM i，我们需要通过获取它们对应的(w, i)来了解默认调度器在SM i上分配的CTA是如何映射到簇i的。为了在运行时识别相应的簇i，一个CTA可以从一个特殊寄存器中获取它当前所在的SM的物理ID【25，Bo Wu, et al. Enabling and exploiting flexible task assignment on GPU through SM-centric program transformations. 2015, ICS】：asm("mov.u32 %0,%%smid;":"=r"(sm_id));。为了识别其在簇中的位置w，一个CTA必须与同一个SM上的其他CTA同步以避免冲突（即两个CTA获得相同的w）。我们发现在Fermi和Kepler上，CTA绑定到硬件warp插槽的方式基本上是连续和固定的。因此，一个CTA可以根据其硬件warp插槽ID除以WARPS_PER_CTA来决定其在簇中的位置w，如清单5的第5-6行所示。然而，在Maxwell和Pascal上，来自不同CTA的warp是动态绑定到硬件warp插槽的【42，Ang Li, et al. SFU-Driven Transparent Approximation Acceleration on GPUs. 2016, ICS】。因此，我们依赖于一个全局原子操作和共享内存广播（清单5的第16-19行）。

4.2.4 整合所有步骤

两种CTA聚类方法的提出。结合所有三个步骤，我们基于两种绑定方案提出了两种CTA聚类方法：
1. 基于重定向的聚类（Redirection-based Clustering）：如图8-(C)所示，该设计建立在基于RR的绑定之上。新内核中的CTA数量与旧内核相同（即$|N|=|O|$，从N到O的一对一映射）。“重定向”意味着N中的每个CTA u被重定向到O中的一个CTA v。我们将这个想法（公式7和8）实现为一个头文件（清单4），基于此，图9展示了使用MM的简单代码转换。尽管该方案易于实现且成本低，但它建立在硬件CTA调度器严格遵循RR的假设之上，而这在真实GPU硬件上已被证明是不正确的（第3.1-(3)节）。之前的一些工作也假设CTA调度为RR【11，Adwait Jog, et al. OWL: Cooperative Thread Array Aware Scheduling Techniques for Improving GPGPU Performance. 2013, ASPLOS】、【27，Minseok Lee, et al. Improving GPGPU resource utilization through alternative thread block scheduling. 2014, HPCA】、【31，Onur Kayıran, et al. Neither more nor less: optimizing thread-level parallelism for GPGPUs. 2013, PACT】、【32，Hyeran Jeon, et al. CTA-aware Prefetching for GPGPU. 2014, CENG-2014-08】、【33，Jin Wang, et al. LaPerm: Locality Aware Scheduler for Dynamic Parallelism on GPUs. 2016, ISCA】。
2. 基于代理的聚类（Agent-based Clustering）：该设计（图8-(D)）建立在基于SM的绑定方案之上，且N中CTA的数量与O不同。与基于重定向的聚类中欺骗GigaThread引擎不同，基于代理的聚类完全规避了硬件CTA调度器。如图8-(D)所示，我们在每个SM上分配几个在内核执行期间持久驻留的CTA（类似于持久线程的概念【26，Kunal Gupta, et al. A study of persistent threads style GPU programming for GPGPU workloads. 2012, InPar】、【43，Sreepathi Pai, et al. Improving gpgpu concurrency with elastic kernels. 2013, ASPLOS】），称为“代理”（agents），通过一个任务循环服务于属于该SM的簇中的所有(w, i)（或任务）。例如，SM-0上的代理A0和A1将执行在分区步骤中被聚类到C0的O中的所有CTA。请注意，一个SM上的代理数量通常远小于簇的大小，例如16个代理对2K个CTA。当一个SM的代理并行工作时（空间局部性），以及在连续任务之间（时间局部性），CTA间局部性得到利用。我们也为该设计实现了一个头文件，如清单5所示，基于此，图10展示了使用MM的非常简单的代码转换。尽管该设计无论默认CTA调度器如何都能工作，但它确实在基于SM的绑定中带来了额外成本（即SM ID获取、代理ID计算和同步）。我们将在第5节讨论这些开销的性能影响。

图8：使用MM作为例子的CTA聚类过程。(A): CTA(1,1)与CTA(0,1)和CTA(2,1)在矩阵A的S区域有CTA间重用，与CTA(1,0)在矩阵B的T区域有CTA间重用。(B): 先前工作假设的默认RR调度策略无法保留CTA间局部性。(C): 基于重定向的聚类使用RR假设将新内核的CTA u绑定到(w, i)，然后通过反转定位旧内核的CTA v。(D): 基于代理的聚类在每个SM上启用持久性CTA（活动代理）来处理(w, i)任务列表，通过寄存器获取sm_id和agent_id来绑定u，然后计算(w, i)并通过反转定位v。

4.3 CTA聚类的补充优化

(I) CTA节流（CTA Throttling）。CTA节流限制了一个SM上并发CTA的数量，以减少对执行资源（如缓存和带宽）的争用。先前的工作【27，Minseok Lee, et al. Improving GPGPU resource utilization through alternative thread block scheduling. 2014, HPCA】、【31，Onur Kayıran, et al. Neither more nor less: optimizing thread-level parallelism for GPGPUs. 2013, PACT】已经观察到，每个SM使用最大数量的CTA并不总是最优的。我们通过控制每个SM的活动代理数量来实现基于软件的CTA节流。这是通过当代理ID大于指定值时直接退出CTA来实现的（清单5第7行）。然而，仅仅通过减少内核网格配置中的代理总数来调整节流程度可能导致不正确的执行，因为硬件调度器在SM间的代理分派不均衡（第3.1-(3)节）。因此，我们总是在内核网格配置中为每个SM分配应用程序允许的最大代理数（例如，受寄存器使用和共享内存限制）（图10中的MAX_AGENTS），以占据CTA插槽，从而强制代理均衡分布，但在运行时只激活其中的一部分（图10中的ACTIVE_AGENTS），这取决于其代理ID（清单5第7和20行）。通过这种方式，我们可以控制节流程度，同时保证代理均匀分布到SM。此外，我们利用__launch_bounds__【7，NVIDIA. CUDA Programming Guide, 2015】在CTA数量少于允许数量时，在编译期间增加寄存器使用，从而利用寄存器重用并隐藏指令延迟。注意，节流并非对所有应用程序都是必需的，它仅在聚类后性能未改善甚至下降时应用（即减少容量未命中）。为了在运行时决定活动代理的数量，我们参考了一个类似于【12，Ang Li, et al. Adaptive and transparent cache bypassing for GPUs. 2015, SC】中使用的动态CTA投票方案。
(II) 缓存旁路（Cache Bypassing）。大量工作【12，Ang Li, et al. Adaptive and transparent cache bypassing for GPUs. 2015, SC】、【13，Chao Li, et al. Localitydriven dynamic GPU cache bypassing. 2015, ICS】、【14，Xuhao Chen, et al. Adaptive Cache Management for Energy-Efficient GPU Computing. 2014, MICRO】、【15，Xiaolong Xie, et al. An Efficient Compiler Framework for Cache Bypassing on GPUs. 2013, ICCAD】、【16，R. Ausavarungnirun, et al. Exploiting Inter-Warp Heterogeneity to Improve GPGPU Performance. 2015, PACT】、【19，Wenhao Jia, et al. Mrpb: Memory request prioritization for massively parallel processors. 2014, HPCA】、【20，Xiaolong Xie, et al. Coordinated static and dynamic cache bypassing for GPUs. 2015, HPCA】已经为GPU提出了缓存旁路技术，旨在避免不必要的缓存污染（如严重的颠簸），减少由于极其有限的缓存容量和资源（如MSHR和未命中队列槽）导致的容量和冲突未命中。作为进一步增强CTA聚类的补充技术，我们旁路流式访问到L1或L1/Tex统一缓存，以防止它们与具有CTA间重用的访问争夺资源。
(III) 使用重塑顺序的CTA预取（CTA Prefetching using Reshaped Order）。如第4.1节所述，CTA预取适用于没有可利用CTA间局部性的应用程序（即数据相关、写相关或流式）。对于这些应用，我们的核心技术——CTA聚类预计不会直接从CTA间局部性的角度带来性能益处。但是CTA聚类（例如，Y-partitioning）可以施加一个特定的CTA调度顺序（即重塑默认顺序），这使得当前CTA能够提前为其后继者预加载所需的数据，从而隐藏到L2或DRAM的长内存访问延迟。这是可行的，因为GPU L1缓存在一个CTA退出后仍保留数据。没有这个重塑的顺序，由于CTA分派的“无序”特性，预取只能限制在CTA范围内【32，Hyeran Jeon, et al. CTA-aware Prefetching for GPGPU. 2014, CENG-2014-08】、【36，Adwait Jog, et al. Orchestrated Scheduling and Prefetching for GPGPUs. 2013, ISCA】、【44，Jaekyu Lee, et al. Many-thread aware prefetching mechanisms for GPGPU applications. 2010, MICRO】、【45，Nagesh B Lakshminarayana and Hyesoon Kim. Spare register aware prefetching for graph algorithms on GPUs. 2014, HPCA】。用于GPU软件预取的宏在清单5的第34-37行列出。

4.4 感知CTA间优化的框架

框架流程与局部性来源估计。图11展示了我们基于软件的利用CTA间局部性的优化框架。前面的小节已经解释了，一旦我们对目标内核有了某些了解（例如，它是否具有可利用的CTA间局部性），如何进行相应的优化。本节主要讨论框架如何估计应用程序的CTA间局部性来源。这个过程在图11中以蓝色突出显示。该框架应用几种简单的粗粒度技术来估计应用程序的CTA间局部性来源，然后才对其进行进一步的优化分析。由于CTA内局部性通常在CTA范围内捕获（例如，通过共享内存），改变CTA调度只会影响CTA间局部性。因此，我们首先应用简单的基于重定向的聚类（图9）来施加一个新的CTA执行顺序（X-或Y-聚类），并检查性能和/或L1命中率是否发生变化（这是CTA间局部性潜力的一个指标）。潜力较高的内核（例如，L1命中率有显著变化）可能属于算法相关或缓存行相关。注意，为了这次验证，最好减小问题规模（总CTA数），因为每个SM的CTA数量过多会引发严重的颠簸，通常使L1命中率接近于零。此外，我们可以打开/关闭L1缓存，看是否有任何性能变化。如果关闭L1后L2缓存访问显著减少，这很可能是由大的L1缓存行引起的，暗示了缓存行相关。如果L2事务没有变化，但CUDA分析器报告的合并度很高，这个内核可能是流式的。然而，如果合并度低，内存访问行为可能更随机，表明是数据相关的。最后，框架评估内核内部对数组的引用。如果内核读写同一个数组，但引用有移位或偏斜，那么写的结果很可能稍后会被其他CTA重用，表明是写相关的。

A3 实验环境

硬件配置: 实验在四代NVIDIA GPU上进行，具体平台参数如表1所示。
- Fermi: Tesla C2050 (CC 2.0), 14 SMs, 3GB GDDR5.
- Kepler: Tesla K40c (CC 3.5), 15 SMs, 12GB GDDR5.
- Maxwell: GTX Titan X (CC 5.2), 24 SMs, 12GB GDDR5.
- Pascal: Tesla P100 (CC 6.0), 56 SMs, 16GB HBM2.
- 各平台每个SM的Warp插槽数、CTA插槽数、寄存器数、共享内存大小以及L1/L2缓存配置详见表1。
软件配置:
- CUDA版本: 使用了与各平台兼容的CUDA Driver & Runtime版本，具体为Fermi (7.0/7.0), Kepler (8.0/8.0), Maxwell (8.0/8.0), Pascal (8.0/8.0)。
- 编译器: 使用NVCC进行编译。
- 操作系统: 未明确说明，但通常为Linux。
基准测试集:
- 名称与来源: 选取了23个有代表性的应用，来自多个流行的GPU基准测试套件，包括Rodinia【46】, Parboil【47】, CUDA SDK【40】和Mars【48】。
- 特征与用途: 这些应用覆盖了第3.2节中定义的五种CTA间局部性来源类别（算法相关、缓存行相关、数据相关、写相关、流式）。每个应用的具体特征，如每个CTA的warp数、基线CTA数、寄存器使用、共享内存使用等，详见表2。这些应用被用来全面评估所提出的CTA聚类技术及其优化方法的性能。

表1：实验平台。 “CC.”是计算能力。“Dri/Rtm”是CUDA驱动和运行时版本。“Warp slots”和“CTA slots”是每个SM的最大warp数和CTA数。“Regs”是每个SM的寄存器数。“SMem”是每个SM的共享内存大小。对于Fermi和Kepler GPU，L1缓存和共享内存是可配置的。

表2：基准测试特征。“WP”代表每个CTA的warp数。“CTAs”表示基线中每个SM的默认CTA数。“6/8/8/8”中的四个项分别是Fermi/Kepler/Maxwell/Pascal的值。此表中的其他列也适用类似的含义。“Registers”表示基线中每个线程的寄存器成本。“SMem”是基线中每个CTA的共享内存成本。“Partition”是在CTA聚类期间采用的划分方法。“Opt Agents”是CTA节流的最优代理数。“Ref”指来源。

A4 实验结果

5.1 总体结果概述

性能提升：如图12所示，对于算法相关的应用，CTA聚类及其相关优化在Fermi、Kepler、Maxwell和Pascal GPU上分别取得了平均1.46倍、1.48倍、1.45倍和1.42倍的性能加速，最高分别达到3.83倍、3.63倍、3.1倍和3.32倍。对于缓存行相关的应用，在L1缓存行较大的Fermi和Kepler GPU上，也分别取得了平均1.47倍和1.29倍的加速，最高达到2.57倍和1.75倍。

缓存性能改善：如图13所示，该技术显著改善了缓存性能。对于算法相关的应用，在四种架构上L2事务数分别减少了55%、65%、29%和28%。对于缓存行相关的应用，在Fermi、Kepler和Maxwell上，L2缓存事务数分别减少了81%、71%和34%。

图12：在四种GPU架构上，通过CTA聚类及其相关优化获得的归一化性能加速和已实现占用率（AC OCP）。每行中的三个子图包括：（左）算法相关，（中）缓存行相关，（右）数据、写和流式相关。“BSL”是基线。“RD”代表基于重定向的聚类。“CLU”代表使用每SM最大允许代理数作为活动代理的基于代理的聚类。“TOT”代表CTA节流。“CLU+TOT”通过节流使用每SM最优的活动代理数。“BPS”是缓存旁路。“PFH”代表CTA预取。已实现占用率定义为每个活动周期平均活动warp数与SM支持的最大warp数的比率。它不是理论占用率。所有结果都归一化到基线，并由多次运行的平均值测量。

图13：在Fermi、Kepler、Maxwell和Pascal GPU上，CTA聚类及优化的L2缓存事务（或L1/Tex-L2事务）和L1缓存命中率。

5.2 观察、分析与局限性

两种聚类方法对比：基于重定向的聚类（RD）由于依赖于不准确的RR调度假设，通用性不强，尽管对某些应用（如NN, IMD）有效。相比之下，基于代理的技术（CLU, CLU+TOT）效果更好，尤其是在Fermi和Kepler上，其静态CTA到warp-slot的绑定使得基于SM的绑定成本更低。
跨架构效果分析：CTA聚类对算法相关的应用更有效，因为其CTA间重用是算法固有的。对于缓存行相关的应用，该技术主要对Fermi和Kepler有效，因为它们的L1缓存行尺寸（128B）远大于Maxwell和Pascal（32B），大缓存行有利于跨CTA边界的空间重用。
补充优化效果分析：
- 节流（Throttling）：对执行资源竞争激烈的应用（如KMN, SYK）非常有效。
- 旁路（Bypassing）：通常效果不佳，因为CTA节流已经缓解了大部分资源竞争。
- 预取（Prefetching）：对于没有可利用CTA间局部性的应用，性能提升不显著。原因有二：(a) 对于缓存友好但容量有限的应用，预取可能破坏原有局部性；(b) 预取需要复杂的地址计算，开销较大。
节流的必要性：多数算法相关应用（KMN除外）无需节流即可从聚类中直接获益。仅通过控制共享内存来节流的传统方法有时性能甚至更差，表明节流是聚类的补充，仅在容量冲突成为瓶颈时使用。
性能评估指标：对比图12和图13，证实了L2缓存事务数是比L1命中率更好的GPU性能评估指标。L2事务数下降通常对应性能提升。
MM应用的局限性分析：尽管MM被用作示例，但其性能提升不显著。原因有三：(1) CTA间数据重用距离远超L1缓存容量，导致时间局部性难以利用；(2) 每个CTA的warp数多（32个），限制了SM上并发代理的数量，使得空间局部性难以利用；(3) Maxwell和Pascal的L1/Tex缓存是分扇区的，阻碍了跨扇区的数据重用。
平台差异总结：CTA聚类在Fermi和Kepler上表现略好于Maxwell和Pascal，因为：(i) 前者有更大的L1缓存行，利于空间重用；(ii) 后者由于动态warp-slot映射，需要承担SM绑定带来的原子和同步开销；(iii) 后者的L1/Tex缓存分扇区，阻碍了代理间的数据重用。

A5 结论

本文提出了一种新颖的聚类技术，旨在发掘并利用一种长期被忽视的局部性——CTA间局部性。首先，通过实验证明了现有GPU硬件有能力在L1或L1/Tex统一缓存上利用空间和时间上的CTA间局部性。接着，通过量化其在广泛应用中的存在并分析其来源，验证了这种局部性的潜力。基于这些洞察，本文提出了CTA聚类的概念及其相关的软件技术。最后，在所有现代NVIDIA GPU架构上对这些技术进行了评估。实验结果表明，所提出的聚类技术能够显著提升片上缓存性能，从而带来可观的整体性能改进。

方法细节中的引用汇总

在论文的核心方法章节（第4节）中，引用了以下文献来支持其设计和论证：

【2, Randy Allen and Ken Kennedy. Optimizing compilers for modern architectures a dependence-based approach. 2001, Morgan Kaufmann】: 在4.2.1节中被引用，用于说明其提出的基于依赖性分析的自动分区方法，类似于编译器中对循环嵌套的依赖性分析。
【7, NVIDIA. CUDA Programming Guide, 2015】: 在4.3节中被引用，用于说明__launch_bounds__这一CUDA特性，作者利用它在CTA数量较少时增加寄存器使用以优化性能。
【11, Adwait Jog, et al. OWL: Cooperative Thread Array Aware Scheduling Techniques for Improving GPGPU Performance. 2013, ASPLOS】: 在4.2.4节中被引用，作为先前假设硬件CTA调度为轮询（RR）的工作之一，为作者提出的基于重定向的聚类方法提供了背景，并指出了该假设的局限性。
【12, Ang Li, et al. Adaptive and transparent cache bypassing for GPUs. 2015, SC】: 在4.3节中被引用，其动态CTA投票方案被本文参考，用于在运行时决定节流（throttling）所需的活动代理数量。
【13, Chao Li, et al. Localitydriven dynamic GPU cache bypassing. 2015, ICS】, 【14, Xuhao Chen, et al. Adaptive Cache Management for Energy-Efficient GPU Computing. 2014, MICRO】, 【15, Xiaolong Xie, et al. An Efficient Compiler Framework for Cache Bypassing on GPUs. 2013, ICCAD】, 【16, R. Ausavarungnirun, et al. Exploiting Inter-Warp Heterogeneity to Improve GPGPU Performance. 2015, PACT】, 【19, Wenhao Jia, et al. Mrpb: Memory request prioritization for massively parallel processors. 2014, HPCA】, 【20, Xiaolong Xie, et al. Coordinated static and dynamic cache bypassing for GPUs. 2015, HPCA】: 这一系列文献在4.3节被引用，作为GPU缓存旁路技术研究的背景，本文将缓存旁路作为CTA聚类的补充优化手段。
【25, Bo Wu, et al. Enabling and exploiting flexible task assignment on GPU through SM-centric program transformations. 2015, ICS】: 在4.2.3节中被引用，用于证实可以通过特殊寄存器（%smid）在运行时获取CTA所在的物理SM ID，这是作者提出的“基于SM的绑定”方法的关键技术基础。
【26, Kunal Gupta, et al. A study of persistent threads style GPU programming for GPGPU workloads. 2012, InPar】: 在4.2.4节中被引用，其“持久线程”概念为作者提出的“基于代理的聚类”中的“代理”（agents）设计提供了思想来源。
【27, Minseok Lee, et al. Improving GPGPU resource utilization through alternative thread block scheduling. 2014, HPCA】: 在4.2.1节和4.2.4节中被引用，作为先前研究CTA调度的代表工作。作者指出其分区方法过于简单（仅划分连续块）且依赖RR假设，与本文更通用的分区和不依赖RR的代理方法形成对比。同时，在4.3节中引用它说明“使用最大CTA数并非最优”的观点，为CTA节流提供了动机。
【31, Onur Kayıran, et al. Neither more nor less: optimizing thread-level parallelism for GPGPUs. 2013, PACT】: 在4.2.4节和4.3节中被引用，同样作为先前假设RR调度以及提出CTA节流思想的工作。
【32, Hyeran Jeon, et al. CTA-aware Prefetching for GPGPU. 2014, CENG-2014-08】 与 【44, Jaekyu Lee, et al. Many-thread aware prefetching mechanisms for GPGPU applications. 2010, MICRO】 及 【45, Nagesh B Lakshminarayana and Hyesoon Kim. Spare register aware prefetching for graph algorithms on GPUs. 2014, HPCA】: 在4.3节被引用，用于说明现有的预取工作主要局限于CTA内部（intra-CTA），而本文通过重塑CTA执行顺序，将预取扩展到了CTA之间。
【33, Jin Wang, et al. LaPerm: Locality Aware Scheduler for Dynamic Parallelism on GPUs. 2016, ISCA】: 在4.2.4节中被引用，作为假设RR调度的先前工作之一。
【36, Adwait Jog, et al. Orchestrated Scheduling and Prefetching for GPGPUs. 2013, ISCA】: 在4.2.1节中被引用，其分区方法被认为过于简单。同时在4.3节中被引用，用于说明现有预取工作的局限性。
【37, Lingda Li, et al. A Graph-based Model for GPU Caching Problems. 2016, arXiv:1605.02043】, 【38, Eddy Z. Zhang, et al. On-the-fly elimination of dynamic irregularities for gpu computing. 2011, ASPLOS】, 【39, Jianqiao Liu, et al. Hybrid CPU-GPU scheduling and execution of tree traversals. 2016, ICS】: 在4.1节被引用，用于说明对于某些数据相关应用，可以通过运行时前的分析（如inspector kernel）来预测数据访问模式，从而使其CTA间局部性变得可利用。
【40, NVIDIA. CUDA SDK Code Samples, 2015】: 在4.2节被引用，作为矩阵乘法（MM）示例代码的来源。
【41, Konstantin Andreev and Harald Racke. Balanced graph partitioning. 2006, Theory of Computing Systems】: 在4.2.1节被引用，用于说明通用的平衡图划分问题是NP完全的，从而论证了本文采用启发式解法的合理性。
【42, Ang Li, et al. SFU-Driven Transparent Approximation Acceleration on GPUs. 2016, ICS】: 在4.2.3节被引用，用于说明Maxwell和Pascal架构上，warps是动态绑定到硬件warp-slots的，这解释了为何在这些架构上需要使用原子操作和共享内存广播来实现基于SM的绑定。
【43, Sreepathi Pai, et al. Improving gpgpu concurrency with elastic kernels. 2013, ASPLOS】: 在4.2.4节中被引用，其“弹性内核”的工作与“持久线程”类似，为“代理”设计提供了参考。

Locality-Aware CTA Clustering for Modern GPUs

A1 主要贡献

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

第2节背景知识

第3节理解GPU上的CTA间重用

3.1 在L1上利用CTA间重用的可行性

3.2 CTA间局部性的量化与来源

A2 方法细节

4.1 总体策略概述

4.2 CTA聚类

4.2.1 步骤1. CTA分区：$f = O \rightarrow C$

4.2.2 步骤2. CTA反转：$f^{-1} = C \rightarrow O$

4.2.3 步骤3. CTA绑定：$g = N \rightarrow C$

4.2.4 整合所有步骤

4.3 CTA聚类的补充优化

4.4 感知CTA间优化的框架

A3 实验环境

A4 实验结果

5.1 总体结果概述

5.2 观察、分析与局限性

A5 结论

方法细节中的引用汇总

💬 评论讨论

A1 主要贡献

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

第2节 背景知识

第3节 理解GPU上的CTA间重用

3.1 在L1上利用CTA间重用的可行性

3.2 CTA间局部性的量化与来源

A2 方法细节

4.1 总体策略概述

4.2 CTA聚类

4.2.1 步骤1. CTA分区：$f = O \rightarrow C$

4.2.2 步骤2. CTA反转：$f^{-1} = C \rightarrow O$

4.2.3 步骤3. CTA绑定：$g = N \rightarrow C$

4.2.4 整合所有步骤

4.3 CTA聚类的补充优化

4.4 感知CTA间优化的框架

A3 实验环境

A4 实验结果

5.1 总体结果概述

5.2 观察、分析与局限性

A5 结论

方法细节中的引用汇总

💬 评论讨论

第2节背景知识

第3节理解GPU上的CTA间重用