Locality-Aware CTA Clustering for Modern GPUs

发表时间: 2017-04

文章标题：面向现代GPU的局部性感知CTA聚类
作者/机构：

速读

一句话结论 本文提出了一种纯软件的 CTA 聚类（CTA-Clustering）技术，通过绕开不可控的硬件调度器将具有数据重用的线程块绑定到同一个流式多处理器上，在多代 GPU 上显著减少了 L2 缓存事务并实现了平均 1.4 倍以上的性能加速。

要解决什么问题 现代 GPU 程序的全局数据重用中有 45% 发生在不同的协作线程阵列（CTA，即线程块）之间，但这种 CTA 间局部性长期被浪费。原有做法卡在 GPU 底层的 CTA 调度器（GigaThread 引擎）完全由硬件控制，对软件既不可见也无法修改。硬件通常采用轮询或随机策略将 CTA 分发给各个流式多处理器（SM），导致原本访问同一块数据的 CTA 被打散到不同的 SM 上，或者在时间上错开执行。由于 GPU 核心内的 L1 缓存容量极小，且采用写驱逐策略，这种无序调度会引发严重的缓存颠簸，导致 L1 命中率极低，迫使 SM 频繁向延迟极高且带宽有限的 L2 缓存和全局内存发起请求。如果不能在机制上接管调度权，把具有潜在重用关系的 CTA 强行聚集到同一个 SM 上并发或连续执行，就无法利用这部分巨大的局部性收益。

怎么做的 作者提出了 CTA 聚类技术，核心思路是废弃默认的内核启动方式，通过纯软件的代码转换构建一个映射机制，欺骗或完全规避硬件调度器，强制有局部性的 CTA 在同一个 SM 上执行。关键设计由分区、反转和绑定三个步骤构成。首先是分区（Partitioning），负责将原始内核中的所有 CTA 划分为 $M$ 个平衡的簇，使得簇内保留最多的 CTA 间局部性。作者通过分析数组引用对网格维度的依赖性来决定是按 X 还是 Y 维度划分，其启发式分区函数定义为：$$f(v) = \left( \lfloor v/M \rfloor - \left( \lfloor (v\%M)/M \rfloor * (\lfloor |V|/M \rfloor + 1) + \lfloor v/M \rfloor * \lfloor |V|/M \rfloor \right), v\%M \right)$$ 其中 $v$ 是原始 CTA ID，$|V|$ 是总 CTA 数，$M$ 是簇的数量，输出结果为簇 ID $i$ 和该 CTA 在簇内的位置 $w$。其次是反转（Inverting），负责在已知簇 ID 和簇内位置时，反向推导出原始的 CTA ID，其反转函数定义为：$$v = f^{-1}((w, i)) = i * (\lfloor |V|/M \rfloor + 1) + w + \min(i - (|V|\%M), 0)$$ 最后是绑定（Binding），负责将实际执行的计算资源与划分好的簇对应起来。为了彻底绕开黑盒的硬件调度器，作者设计了基于代理的聚类（Agent-based Clustering）。该机制在每个 SM 上只启动少量在内核执行期间持久驻留的 CTA（称为代理），代理通过读取特殊寄存器 `%smid` 获取当前物理 SM 的 ID，从而知道自己该负责哪个簇。代理内部运行一个任务循环，不断获取簇内位置 $w$，结合所在簇 $i$，利用反转函数计算出原始 CTA ID $v$，进而执行原始内核的计算逻辑。作为补充，该方法还引入了 CTA 节流机制，通过限制每个 SM 上活跃的代理数量来缓解缓存争用，并配合缓存旁路和重塑顺序后的 CTA 预取来进一步掩盖内存延迟。

效果如何 实验在 Fermi、Kepler、Maxwell 和 Pascal 四代 NVIDIA GPU 上搭建，测试了来自 Rodinia、Parboil 等基准测试集的 23 个应用。对比基线包括默认硬件调度的基线（BSL）、基于轮询假设的重定向聚类（RD）、基于代理的聚类（CLU）、结合 CTA 节流的代理聚类（CLU+TOT）、缓存旁路（BPS）以及 CTA 预取（PFH）。量化结果显示，对于具有算法相关局部性的应用，结合节流的代理聚类在四代架构上分别实现了平均 1.46 倍、1.48 倍、1.45 倍和 1.42 倍的性能加速，最高加速达 3.83 倍，L2 缓存事务数分别锐减了 55%、65%、29% 和 28%。该方法存在明确的适用边界：它对算法相关和缓存行相关的应用极为有效，但对数据相关、写相关或纯流式应用几乎没有收益。此外，架构差异会带来显著的代价波动，在较新的 Maxwell 和 Pascal 架构上，由于 L1 缓存行较小（32B）且被划分为多个扇区，阻碍了跨扇区的数据重用，加上动态 warp 绑定机制引入了额外的原子操作和同步开销，其聚类效果略逊于拥有 128B 大缓存行的老架构。作者也承认，对于像矩阵乘法这种数据重用距离远超 L1 缓存容量、且每个 CTA 包含大量 warp 导致并发代理数受限的应用，时间与空间局部性均难以被有效捕获，性能提升十分有限。

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数据缓存的不同全局内存读取数据路径。
图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()）用于禁止多发射。

清单3：用于验证空间和时间CTA间局部性的微基准测试。
清单3：用于验证空间和时间CTA间局部性的微基准测试。

实验观察与分析。图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间局部性）。
图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】。

图3：常见GPU应用程序的CTA内和CTA间重用百分比。
图3：常见GPU应用程序的CTA内和CTA间重用百分比。

CTA间局部性的来源分类。然而，这些显著的CTA间局部性是否能在GPU硬件上被利用以提升性能？为了回答这个问题，我们需要理解CTA间局部性的来源。通过对广泛的GPU应用程序（详见表2）进行特征分析，我们将GPU CTA间局部性的来源分为以下五个类别，其模式如图4所示：

图4：基于CTA间局部性来源的五个应用类别。
图4：基于CTA间局部性来源的五个应用类别。

A2 方法细节

4.1 总体策略概述

基于可利用性的两种优化策略。如前所述，一些应用程序（如算法相关的）具有丰富的CTA间局部性，而另一些（如流式）则没有。根据应用程序是否具有可利用的CTA间局部性，我们提出了一种操纵GPU上CTA调度的软件方法，即CTA聚类（CTA-Clustering）。我们将应用程序是否具有“可利用的”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)节）。

图5：优化策略。O代表原始内核。N代表新内核。
图5：优化策略。O代表原始内核。N代表新内核。

4.2 CTA聚类

CTA聚类的核心思想与步骤。CTA聚类的基本思想是启动一个新内核来替换原始内核，从而建立一个预定义的聚类规则。目标是将具有CTA间局部性的CTA聚集在一起，并在同一个SM上并发或连续执行。我们使用O、N、C分别表示原始内核、新内核和簇。CTA聚类本质上是找到一个映射：$N \rightarrow O$，该映射受制于C，如图6所示。这个过程包括三个步骤：

简而言之，寻找映射 $N \rightarrow O$ 的过程如下：

公式1
公式1

以矩阵乘法（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))$。
图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节）的一部分来确保聚类的平衡性。我们将划分函数定义为：

公式2
公式2

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

公式3
公式3

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

公式4
公式4

否则：

公式5
公式5

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间局部性的应用程序的划分方法：

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

图7：2D网格的四种主要CTA索引方法。
图7：2D网格的四种主要CTA索引方法。

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

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

公式6
公式6

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

公式7
公式7

这两个方程可以统一为：

公式8
公式8

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调度器：

公式9
公式9

在MM中，对于一个ID为4的N中的CTA，其关联的(w, i)可以通过公式8计算：$(w,i)=(4/2, 4\%2)=(2,0)$。

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。
图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：基于重定向的聚类头文件
清单4：基于重定向的聚类头文件

图9：通过基于重定向的聚类进行的内核转换
图9：通过基于重定向的聚类进行的内核转换

清单5：基于代理的聚类头文件
清单5：基于代理的聚类头文件

图10：通过基于代理的聚类进行的内核转换
图10：通过基于代理的聚类进行的内核转换

4.3 CTA聚类的补充优化

  1. (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投票方案。
  2. (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间重用的访问争夺资源。
  3. (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重用，表明是写相关的。

图11：CTA间局部性感知优化框架
图11：CTA间局部性感知优化框架

A3 实验环境

表1：实验平台。 “CC.”是计算能力。“Dri/Rtm”是CUDA驱动和运行时版本。“Warp slots”和“CTA slots”是每个SM的最大warp数和CTA数。“Regs”是每个SM的寄存器数。“SMem”是每个SM的共享内存大小。对于Fermi和Kepler GPU，L1缓存和共享内存是可配置的。
表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”指来源。
表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数的比率。它不是理论占用率。所有结果都归一化到基线，并由多次运行的平均值测量。
图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缓存命中率。
图13：在Fermi、Kepler、Maxwell和Pascal GPU上，CTA聚类及优化的L2缓存事务（或L1/Tex-L2事务）和L1缓存命中率。

5.2 观察、分析与局限性

  1. 两种聚类方法对比：基于重定向的聚类（RD）由于依赖于不准确的RR调度假设，通用性不强，尽管对某些应用（如NN, IMD）有效。相比之下，基于代理的技术（CLU, CLU+TOT）效果更好，尤其是在Fermi和Kepler上，其静态CTA到warp-slot的绑定使得基于SM的绑定成本更低。

  2. 跨架构效果分析：CTA聚类对算法相关的应用更有效，因为其CTA间重用是算法固有的。对于缓存行相关的应用，该技术主要对Fermi和Kepler有效，因为它们的L1缓存行尺寸（128B）远大于Maxwell和Pascal（32B），大缓存行有利于跨CTA边界的空间重用。

  3. 补充优化效果分析：

    • 节流（Throttling）：对执行资源竞争激烈的应用（如KMN, SYK）非常有效。
    • 旁路（Bypassing）：通常效果不佳，因为CTA节流已经缓解了大部分资源竞争。
    • 预取（Prefetching）：对于没有可利用CTA间局部性的应用，性能提升不显著。原因有二：(a) 对于缓存友好但容量有限的应用，预取可能破坏原有局部性；(b) 预取需要复杂的地址计算，开销较大。

  4. 节流的必要性：多数算法相关应用（KMN除外）无需节流即可从聚类中直接获益。仅通过控制共享内存来节流的传统方法有时性能甚至更差，表明节流是聚类的补充，仅在容量冲突成为瓶颈时使用。

  5. 性能评估指标：对比图12和图13，证实了L2缓存事务数是比L1命中率更好的GPU性能评估指标。L2事务数下降通常对应性能提升。

  6. MM应用的局限性分析：尽管MM被用作示例，但其性能提升不显著。原因有三：(1) CTA间数据重用距离远超L1缓存容量，导致时间局部性难以利用；(2) 每个CTA的warp数多（32个），限制了SM上并发代理的数量，使得空间局部性难以利用；(3) Maxwell和Pascal的L1/Tex缓存是分扇区的，阻碍了跨扇区的数据重用。

  7. 平台差异总结：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节）中，引用了以下文献来支持其设计和论证：