A Compiler Framework for Optimizing Dynamic parallelism on GPUs

文章标题：一个用于优化GPU动态并行的编译器框架
作者/机构：Mhd Ghaith Olabi¹, Juan Gomez Luna², Onur Mutlu², Wen-mei Hwu³,⁴, Izzat El Hajj¹
¹贝鲁特美国大学，黎巴嫩
²苏黎世联邦理工学院，瑞士
³英伟达，美国
⁴伊利诺伊大学厄巴纳-香槟分校，美国

A1 主要贡献

核心问题：GPU上的动态并行允许GPU线程动态启动其他GPU线程，这对于具有嵌套并行性的应用，特别是嵌套并行量不规则且无法预先预测的应用非常有用。然而，先前的工作表明，当启动大量小规模网格（grid）时，动态并行可能会带来巨大的性能损失。大量的启动会导致高昂的启动延迟（由于拥塞），而小规模的网格会导致硬件利用率不足。

研究目标：为了解决上述问题，本文提出了一个用于优化嵌套并行应用中动态并行使用的编译器框架。该框架旨在通过自动化和增强现有优化技术来降低动态并行的开销，从而提高性能。

创新点/主要贡献：
本文提出了一个集成了三种关键优化技术的编译器框架，主要贡献如下：
* 自动化阈值（Thresholding）优化：提出了一种编译器变换，能够自动实现阈值优化。该优化仅在子线程数量超过某个阈值时才动态启动网格，否则在父线程中串行执行子线程的工作。这减少了启动的网格数量，从而减轻了拥塞，并确保只有能有效利用硬件的大规模网格才被启动。
* 在动态并行中应用粗化（Coarsening）优化：提出将粗化优化应用于动态并行的上下文中，并提供了一个编译器变换来实现它。粗化通过将多个子线程块的工作合并到一个线程块中来执行，从而减少了需要调度的子线程块数量，并能与聚合优化相结合，摊销聚合开销。
* 提出多块粒度的聚合（Aggregation）技术：对先前工作中的聚合技术进行了增强，提出了一种新的使用多块粒度（multi-block granularity）的聚合技术。该技术在一组线程块中的父线程之间聚合启动，填补了先前工作中单线程块和整个网格这两种极端粒度之间的空白。
* 集成化的开源编译器框架：将阈值、粗化和聚合这三种优化技术整合到一个开源的编译器框架中，简化了优化动态并行代码的过程。

成果总结：评估表明，该编译器框架相比于仅使用动态并行的应用，性能提升了43.0倍（几何平均值）；相比于不使用动态并行的应用，性能提升了8.7倍；相比于先前工作中仅使用聚合优化的应用，性能提升了3.6倍。

A3 背景知识

A. 动态并行

动态并行的基本用法与效率问题。在实践中，GPU上执行的父线程可能会各自发现一些可以并行化的嵌套工作。如图1(a)所示，每个父线程启动一个子网格来并行执行这些嵌套工作，并且每个父线程可以为其子网格提供不同的启动配置和参数。由于嵌套工作的数量可能因线程而异，因此子网格的大小也不同。以这种方式使用动态并行可能产生的一个低效来源是，可能会启动大量的子网格，其中许多网格的规模可能很小【索引40, Characterization and analysis of dynamic parallelism in unstructured GPU applications, 2014, IISWC】。在这种情况下，大量的子网格启动会导致拥塞，而子网格的小规模则导致设备利用率低下。

图1. 动态并行背景
图1. 动态并行背景

B. 聚合

聚合优化的概念。聚合是一种将由多个父线程启动的子网格合并或聚合为单个网格的优化。许多工作【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】【索引24, Exploiting dynamic parallelism to efficiently support irregular nested loops on GPUs, 2015, CoSMiC】【索引25, Nested parallelism on GPU: Exploring parallelization templates for irregular loops and recursive computations, 2015, ICPP】【索引41, Compiler-assisted workload consolidation for efficient dynamic parallelism on GPU, 2016, IPDPS】通过手动或在编译器中应用此优化。图1(b)展示了如何将聚合应用于图1(a)中的示例。在此示例中，父线程们协调计算它们所有子网格的累积大小，以便启动一个单一的聚合网格。在先前的工作中，协调的范围（称为聚合粒度）是同一warp、同一线程块或整个网格内的父线程。如果范围是warp或块，则由其中一个参与的父线程代表其他线程启动聚合网格。如果范围是整个网格，则聚合网格从主机端启动。

聚合与反聚合逻辑。在原始代码中，每个父线程可能为其子网格提供不同的参数和启动配置。然而，在转换后的代码中，只能提供一套参数和启动配置。因此，在启动聚合网格之前，每个父线程将其原始参数和启动配置存储在内存中，并将指向该内存的指针传递给聚合网格。子线程随后必须识别其原始父线程，以便从内存中加载正确的参数和配置。为此，每个子线程块执行一个搜索操作。父线程为确定聚合网格大小并将各自参数和启动配置存入内存所做的工作称为聚合逻辑（aggregation logic）。子线程为识别其原始父线程并从内存中加载其原始参数和配置所做的工作称为反聚合逻辑（disaggregation logic）。

聚合的优缺点与粒度权衡。聚合的优点在于它减少了子网格的启动数量，从而减轻了拥塞，并增大了子网格的规模以确保更好的硬件利用率。其缺点是会因聚合和反聚合逻辑而产生开销，并且会延迟子网格的启动，直到所有父线程都准备好启动。聚合粒度（warp、块或网格）的选择涉及在这些优缺点之间进行权衡。使用更大的粒度可以减少拥塞并提高利用率，但会产生更高的聚合和反聚合逻辑开销，并使子网格的启动延迟更长。

A2 方法细节

III. 阈值（THRESHOLDING）

A. 优化概述

阈值优化的核心思想。阈值（Thresholding）是一种优化方法，即仅当子线程数量超过特定阈值时才动态启动子网格。否则，子线程将由父线程顺序执行。图2展示了阈值如何应用于图1(a)的示例中。在该示例中，两个父线程的子线程数量很少，并行化这些子线程所获得的收益不太可能值得付出启动开销。因此，这些父线程转而顺序执行子线程的工作。

图2. 带有阈值优化的动态并行示例
图2. 带有阈值优化的动态并行示例

自动化阈值的动机。阈值通常由程序员手动应用【索引24, Exploiting dynamic parallelism to efficiently support irregular nested loops on GPUs, 2015, CoSMiC】【索引25, Nested parallelism on GPU: Exploring parallelization templates for irregular loops and recursive computations, 2015, ICPP】【索引34, Controlled kernel launch for dynamic parallelism in GPUs, 2017, HPCA】【索引41, Compiler-assisted workload consolidation for efficient dynamic parallelism on GPU, 2016, IPDPS】。然而，手动应用阈值会使启动代码复杂化，需要代码复制，并损害代码的可读性。因此，我们建议通过编译器变换来自动化阈值优化（第III-B节），并讨论实现自动化所面临的挑战（第III-C和III-D节）。

B. 代码转换

阈值转换示例。图3展示了我们的编译器如何应用阈值转换。图3(a)中的原始代码包含一个父内核（04-08行）和一个子内核（01-03行）。父内核使用动态并行调用子内核（06行），并为其配置网格维度gDim和块维度bDim。

转换后的代码结构。图3(b)展示了应用阈值转换后的代码。此转换包含两个关键部分：构建一个由父线程执行的子内核的串行版本（09-15行），以及应用一个阈值来决定是执行启动还是调用串行版本（21-26行）。

构造串行版本。为了构造串行版本，我们复制子内核并将其属性更改为 __device__，使其成为一个设备函数（09行）。参数列表中追加了两个参数：_gDim，代表并行版本的原始网格维度；_bDim，代表原始块维度。在子内核主体周围插入循环（10-11行）以串行化子线程。第一个循环（10行）遍历子线程块，而第二个循环（11行）使用作为参数传递的边界来遍历每个子块中的子线程。最后，所有对保留索引和维度变量的使用都被替换为相应的循环索引和边界。为简单起见，该示例展示了一个一维子内核，但如果子内核是多维的，则会为每个维度插入循环。

应用阈值判断。为了应用阈值，首先识别子线程的数量并将其存储在 _threads 变量中（21行），第III-D节将讨论如何识别子线程数量。接下来，在子内核调用周围插入一个if语句（22行），以确保仅当子线程数大于或等于 _THRESHOLD 时才执行动态启动。这里的 _THRESHOLD 是一个宏变量，可以在编译时覆盖以进行调优。如果不满足阈值，则调用实现串行版本的设备函数（25行），从而在父线程中串行化子工作。

图3. 阈值代码转换示例
图3. 阈值代码转换示例
01 __global__ child(params) { 
02   child body 
03 } 
04 __global__ parent(...) { 
05   ... 
06   child <<< gDim, bDim >>> (args); 
07   ... 
08 }
(a) 原始代码
09 _device__ child_serial(params, dim3 _gDim, dim3 _bDim) { 
10   for(_bx = 0; _bx < _gDim.x; ++_bx) { 
11     for(_tx = 0; _tx < _bDim.x; ++_tx) { 
12       child body // 用 _bx 替换 blockIdx.x, 
13     }            // 用 _tx 替换 threadIdx.x, 用 _gDim
14   }              // 替换 gridDim, 用 _bDim 替换 blockDim
15 } 
16 __global__ child(params) { 
17   child body 
18 } 
19 _global__ parent(...) { 
20   
21   _threads = ...; // 从 gDim 表达式中提取
22   if(_threads >= _THRESHOLD) { 
23     child <<< gDim, bDim >>> (args); 
24   } else { 
25     child_serial (args, gDim, bDim); 
26   } 
27   
28 }
(b) 经过阈值转换的代码

C. 不可转换的内核

不可转换内核的类型。并非所有子内核都适合进行第III-B节中描述的转换。特别地，我们不转换两种类型的子内核：(1) 通过 syncthreads() 或warp级原语执行线程间屏障同步的子内核，以及 (2) 使用共享内存的子内核。

含屏障同步的内核。对于执行屏障同步的子内核，虽然文献中已有支持此类同步的GPU线程串行化方法【索引12, Dynamic loop vectorization for executing OpenCL kernels on CPUs, 2014, Master’s thesis】【索引18, Improving performance of OpenCL on CPUs, 2012, CC】【索引20, Localitycentric thread scheduling for bulk-synchronous programming models on CPU architectures, 2015, CGO】【索引21, Multi-tier dynamic vectorization for translating GPU optimizations into CPU performance, 2014, Tech. Rep.】【索引33, MCUDA: An efficient implementation of CUDA kernels for multi-core CPUs, 2008, LCPC】，但将它们扩展到在单个GPU线程中串行化多个GPU线程是不切实际的。原因有二：首先，这些技术对所有局部变量进行标量扩展以跨屏障保留所有线程的状态，在GPU上这种标量扩展会将所有寄存器访问转换为内存访问，代价极高。其次，包含屏障同步的代码通常实现的是一种在串行化时效率不高的并行算法（例如，并行归约），在这种情况下，最好由程序员手动应用阈值优化，因为最优的串行和并行算法是不同的。

使用共享内存的内核。对于使用共享内存的子内核，我们不构造其串行版本，因为每个父线程将需要与整个子块一样多的共享内存，这将使父块的共享内存需求过高。此外，使用共享内存的内核通常会使用 __syncthreads() 来协调线程间对共享内存的访问，因此这些内核很可能因为之前提到的屏障同步原因而无法转换。

D. 识别子线程数量

识别子线程数的挑战。第III-B节中描述的转换需要识别所需的子线程数量，以便与阈值进行比较。识别所需子线程数量具有挑战性，因为程序员在内核调用中提供的是网格维度（块数）和块维度（每块线程数），而不是直接提供所需子线程数。程序员使用所需的子线程数来计算网格维度。

不准确的识别方法。一种识别所需子线程数的方法是将网格维度与块维度相乘。然而，这种方法给出了子网格中的总线程数，包括未被使用的线程，这个值不能作为与阈值比较的代表性值。例如，如果一个父线程需要2个单位的嵌套并行工作，并将子块维度配置为1024，它将配置子内核为1个块。此时，网格维度与块维度相乘得到1024，远大于实际所需的2个线程。理想情况下，应与阈值比较的是2，而不是1024。

基于模式匹配的启发式方法。我们使用的方法基于一个观察：程序员通常将网格维度计算为所需线程数与块维度的向上取整除法（ceiling-division）。由于向上取整除法有多种表达方式，无法通过静态分析百分之百确定所需线程数。因此，我们识别程序员最常用的向上取整除法模式，并基于这些模式采用简单的静态分析。

常见的向上取整除法模式。图4显示了程序员编写的用于计算网格维度的常见向上取整除法表达式。选项(a)-(c)使用整数运算，而(d)-(e)转换为浮点数并使用ceil函数。选项(f)用于多维块。在所有选项中，表达式可以是整体的，也可以是分部分的。请注意，$N$ 和 $b$ 可以是任意表达式。

(a) $(N – 1)/b + 1$ $N$: 所需线程数
(b) $(N + b – 1)/b$ $b$: 块维度
(c) $N/b + (N\%b == 0)?0:1$
(d) ceil((float)N/b)
(e) ceil(N/(float)b)
(f) dim3(..., ..., ...)// dim3的参数可能是上述表达式之一

启发式分析过程。我们观察到，$N$ 通常位于除法左侧的子表达式中，该子表达式可能还包含像1或 $b$ 这样的常量。基于此观察，我们的分析过程寻找除法操作，取其左侧的子表达式，移除常量的加减法，并将剩余的子表达式视为所需的线程数。这种分析本质上是启发式的，不保证能找到真正的所需线程数。然而，在这种情况下使用启发式方法是可以接受的，因为结果仅用于选择串行化还是并行化工作，不会影响程序的正确性。

代码实现。找到的子表达式被赋给图3中第21行的 _threads 变量。然后，gDim 中该子表达式的出现位置被替换为 _threads，以确保在表达式有副作用的情况下不会在代码中重复计算。

IV. 粗化（COARSENING）

A. 优化概述

粗化优化的定义与优势。粗化是一种将原始代码中多个线程块的工作分配给单个线程块的优化。分配给每个粗化线程块的原始线程块数量称为粗化因子。当GPU硬件的线程块数量超负荷时，硬件会串行化这些线程块。在代码中而非由硬件来串行化这些线程块有多个优点：首先，它减少了需要调度的线程块数量，并允许粗化块中的一些warp在其他warp完成前一个原始块的工作之前，继续执行下一个原始块的工作。其次，如果原始线程块之间有共同的工作，这些工作可以被提取出来由粗化线程块执行一次，从而摊销其成本。粗化的缺点是它会降低并行度，如果粗化因子过高，可能会导致设备利用率不足。

在动态并行中应用粗化。粗化作为一种优化经常被程序员在许多不同上下文中使用【索引22, Programming Massively Parallel Processors: A Hands-On Approach, 2016, Book】。先前的工作也在编译器中应用了粗化【索引26, Automatic optimization of thread-coarsening for graphics processors, 2014, PACT】【索引32, Predictable thread coarsening, 2018, TACO】，但并未在动态并行的背景下进行。我们建议在动态并行的背景下应用粗化，并通过编译器转换来自动化其应用。

粗化与动态并行结合的优势。图5展示了如何将粗化应用于图1(a)中示例的子线程块。在此示例中，转换后代码中的每个粗化子线程块执行原始代码中两个子线程块的工作。在动态并行的背景下应用粗化的优势在于它减少了需要调度的子线程块数量。更重要的是，当在聚合之前应用粗化时，粗化还提供了每个块能完成更多工作的子线程块，这使它们更有能力摊销反聚合逻辑的开销。

图5. 带有粗化优化的动态并行示例
图5. 带有粗化优化的动态并行示例

B. 代码转换

粗化转换示例。图6展示了我们的编译器如何将粗化转换应用于图3(a)的原始代码。此转换包含两个关键部分：粗化子内核（01-05行）和修改启动配置以启动粗化后的子内核（08-10行）。

粗化子内核。为了粗化子内核，一个参数 _gDim 被附加到参数列表中（01行），它代表未经粗化的原始网格维度。插入一个粗化循环（02行），该循环遍历分配给粗化块的原始子线程块的工作。对保留索引和维度变量的使用被替换为相应的循环索引和边界。为简单起见，该示例仅显示了一维粗化，但如果子网格是多维的，则会为每个维度插入循环。

修改启动配置。为了修改启动配置以启动粗化后的子内核，原始网格维度 gDim 被存储在一个变量 _gDim 中（08行）。该值也被复制到代表粗化后网格维度的 _cgDim。然后，粗化网格维度 _cgDim 的x维度通过向上取整除法除以粗化因子 _CFACTOR（09行）。这里的 _CFACTOR 是一个宏变量，可以在编译时覆盖以进行调优。同样，为简单起见，该示例显示了一维粗化。最后，子内核被配置为使用粗化后的网格维度，而原始网格维度作为参数传递（10行）。

图6. 粗化代码转换示例
图6. 粗化代码转换示例
01 _global child(params, _gDim) { 
02   for(_bx = blockIdx.x; _bx < _gDim.x; _bx += gridDim.x) { 
03     child body // 用 _bx 替换 blockIdx.x
04   }            // 用 _gDim 替换 gridDim
05 } 
06 __global__ parent(...) { 
07   
08   _cgDim = _gDim = gDim ; 
09   _cgDim.x = (_gDim.x + _CFACTOR – 1)/_CFACTOR; 
10   child <<< _cgDim, bDim >>>(args, _gDim); 
11   
12 }

V. 聚合（AGGREGATION）

聚合已由先前的工作提出【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】【索引24, Exploiting dynamic parallelism to efficiently support irregular nested loops on GPUs, 2015, CoSMiC】【索引25, Nested parallelism on GPU: Exploring parallelization templates for irregular loops and recursive computations, 2015, ICPP】【索引41, Compiler-assisted workload consolidation for efficient dynamic parallelism on GPU, 2016, IPDPS】，并在第II-B节中进行了描述。在本文中，我们提出了一种新的聚合粒度，即多块粒度（第V-A节）。我们还建议应用聚合阈值来优化warp和块粒度的聚合（第V-B节）。

A. 多块粒度聚合

提出多块粒度的动机。回顾第II-B节，先前的工作已在warp、块和网格粒度上执行聚合。使用更大的粒度可以减少拥塞并提高硬件利用率，但会产生更高的聚合和反聚合逻辑开销，并使子网格的启动延迟更长。聚合粒度的选择在这些性能因素之间提供了一种权衡。然而，块和网格粒度聚合方案之间存在巨大差距，留下了大量未被探索的权衡空间。为了更好地探索权衡空间，我们提出了一种介于块和网格粒度之间的中间粒度，即多块粒度聚合。

多块粒度聚合的实现。在多块粒度聚合中，我们将父网格划分为多个块组，每个组具有固定数量的块。同一组块内的线程协作将其子网格聚合为单个聚合网格。在先前的工作中【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】，聚合逻辑涉及对原始网格维度进行扫描操作和对块维度进行最大值操作以确定聚合网格配置，并使用屏障同步等待所有参与线程。由于我们无法跨多个线程块进行同步，我们使用原子操作来执行多块粒度的扫描和最大值操作，类似于先前工作中网格粒度的做法。至于屏障同步，我们用一个组范围的计数器替换它，每个线程块完成时原子地增加该计数器。组中最后一个增加计数器的线程块执行聚合启动。

多块粒度聚合的代码转换。图7展示了我们的编译器如何将多块粒度聚合转换应用于图3(a)的原始代码。该转换包括两部分：父内核中的聚合逻辑（14-35行）和子内核中的反聚合逻辑（01-11行）。

聚合逻辑（父内核）。在父内核中，我们首先保存gDim和bDim表达式。然后，识别线程块所属的组（16行），并在预分配的内存缓冲区中找到该组的内存段（17行）。接着，每个启动子网格的线程同时原子地增加两个全局计数器（19-20行）：(1) _numParents，为父线程分配索引；(2) _sumGDim，计算先前父线程的子块总数。然后，每个线程将其参数、扫描后的网格维度和块维度存储到内存中（21-23行），并执行原子操作以找到最大块维度（24行）。

同步与启动。每个线程将其配置和参数写入全局内存后，执行一次fence操作（26行）以确保其子块可见这些数据。这在先前的工作中是不必要的，但在这里是必须的，因为启动可能由不同的线程块执行。还需要一个局部屏障（27行）确保块内所有线程完成数据存储。最后，块中的一个线程（28行）原子地增加组范围的计数器（29行），并检查其块是否是组中最后一个完成的（30-31行）。如果是，该线程就启动聚合网格（32-33行）。

反聚合逻辑（子内核）。在子内核中（01-11行），代码与先前的工作【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】基本相同。每个子块通过对扫描后的网格维度数组进行二分搜索来识别其原始父线程（02行），然后加载其参数和配置（04-06行），并基于这些信息执行子内核的工作（07-10行）。

01 _global child(_paramsArray, _gDimScannedArray, _bDimArray) { 
02   _parentIdx = binary search in _gDimScannedArray 
03   params = _paramsArray[_parentIdx] 
04   _gDim = _gDimScannedArray[_parentIdx] - _gDimScannedArray[_parentIdx - 1] 
05   _bx = blockIdx.x - _gDimScannedArray[_parentIdx - 1] 
06   _bDim = _bDimArray[_parentIdx] 
07   if(threadIdx < _bDim) { 
08     child body // 用 _bx 替换 blockIdx.x
09                // 用 _gDim 替换 gridDim
10   } 
11 } 
12 _global parent(...) { 
13   
14   _gDim = gDim 
15   _bDim = bDim 
16   _groupIdx = blockIdx.x/_AGG_GRANULARITY 
17   // 根据 _groupIdx 在预分配的缓冲区中找到组的内存段
18   if(_gDim > 0) { 
19     (_parentIdx, _sumPrevGDim) = 
20       atomicAdd(&(_numParents[_groupIdx], _sumGDim[_groupIdx]), (1, _gDim)) 
21     _argsArray[_parentIdx] = args 
22     _gDimScannedArray[_parentIdx] = _sumPrevGDim + _gDim 
23     _bDimArray[_parentIdx] = _bDim 
24     atomicMax(&_maxBDim[_groupIdx], _bDim) 
25   } 
26   _threadfence() 
27   _syncthreads() 
28   if(threadIdx == launcher thread in block) { 
29     _nFinishedBlocks = atomicAdd(&_numFinishedBlocks[_groupIdx], 1) + 1 
30     _isLastBlockToFinish = (_nFinishedBlocks == _AGG_GRANULARITY) 
31     if(_isLastBlockToFinish) { 
32       child <<< _sumGDim[_groupIdx] , _maxBDim[_groupIdx] >>> 
33         (_argsArray, _gDimScannedArray, _bDimArray); 
34     } 
35   } 
36   
37 }

多块粒度的额外优势。除了填补块和网格粒度聚合之间的权衡空间外，多块粒度聚合相比于网格粒度聚合还有一个优势。网格粒度聚合需要CPU参与聚合启动，而多块粒度聚合完全在GPU上运行，从而使CPU可以执行其他任务。因此，多块粒度聚合与内核调用的异步语义更兼容。

B. 聚合阈值

聚合阈值的引入。当在聚合之前应用阈值优化时，参与聚合网格的原始子网格数量可能会大幅减少。如果参与聚合的原始子网格数量不足，聚合的收益可能不足以抵消其开销。为了解决这个问题，我们通过应用聚合阈值来增强聚合。聚合逻辑之前会执行一个操作来计算参与的父线程数量。如果参与的父线程数量未达到某个阈值，子网格将由其父线程正常启动，而不是被聚合。因此，一个子网格可能以三种方式之一执行：在其父线程内串行化，由其父线程直接启动，或作为聚合网格的一部分启动。由于应用聚合阈值需要父线程同步以计算参与线程的数量，因此它只能在可以进行线程间屏障同步的warp和块粒度上应用。

VI. 编译器框架

优化的集成与应用顺序。我们将我们的三种优化——阈值、粗化和聚合——集成到一个单一的编译器框架中。为实现关注点分离，每种优化都作为独立的源到源转换过程实现，每个过程接收一个CUDA .cu文件并生成一个.cu文件。这些转换是独立的，意味着可以按任何顺序应用它们的任意组合，同时生成正确的代码。

推荐的应用顺序。尽管优化可以按任何顺序应用，我们按以下顺序应用它们：阈值、粗化，然后是聚合，如图8(a)所示。
* 阈值在粗化之前应用，因为粗化会操作网格维度，这使得提取需要与阈值比较的线程数变得更加困难。
* 阈值在聚合之前应用，因为聚合会将小网格与大网格合并为一个聚合网格，在它们被聚合成更大的网格后，再隔离小网格并进行串行化会更加困难。
* 粗化在聚合之前应用，因为反聚合逻辑应该位于粗化循环之外，以便其开销可以在多个原始子块之间摊销。

图8. 组合三种优化
图8. 组合三种优化

组合优化的效果。图8(b)展示了组合所有三种优化对图1(a)中示例的影响。在此示例中，两个父线程的子网格较小，因此子网格的工作通过阈值优化在父线程中串行化。其余两个父线程的子网格较大，因此它们协作聚合其网格并执行单次启动。聚合网格中的每个子块搜索其父线程并从内存中获取相应的参数和配置。然后，该子块执行多个原始子块的工作，因为它在聚合之前已被粗化。

实现细节。这些编译器转换是在Clang【索引23, LLVM: A compilation framework for lifelong program analysis & transformation, 2004, CGO】中作为源到源转换过程实现的。阈值和粗化转换及其支持的分析是从头开始实现的。聚合转换是通过修改先前工作之一【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】提供的实现来完成的。我们利用这项工作是因为它是开源的，但我们的技术也可以应用于任何执行聚合的先前工作【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】【索引41, Compiler-assisted workload consolidation for efficient dynamic parallelism on GPU, 2016, IPDPS】。

A4 实验环境

TABLE I 基准测试和数据集
基准测试和数据集表格1

基准测试和数据集表格2
基准测试和数据集表格2

A4 实验结果

A. 性能

图9展示了各基准测试和数据集在应用所有优化组合后的性能结果，性能以相对于CDP版本的加速比报告。

图9. 性能（越高越好）
图9. 性能（越高越好）

B. 执行时间分解

图10展示了各基准测试和数据集的执行时间花费情况，以KLAP（CDP+A）为基线。

图10. 执行时间分解（越低越好）
图10. 执行时间分解（越低越好）

C. 阈值和聚合粒度的影响

图11展示了在保持最佳粗化因子不变的情况下，性能如何随阈值和聚合粒度的变化而变化。

图11. 阈值和聚合粒度的影响
图11. 阈值和聚合粒度的影响

D. 低嵌套并行度工作负载

图12展示了在嵌套并行度较低的道路图上图基准测试的性能。

图12. 图基准测试在道路图上的性能（越高越好）
图12. 图基准测试在道路图上的性能（越高越好）

A7 补充细节

IX. 相关工作

本文的独特性。据我们所知，我们的工作是第一个提供一个编译器框架来优化动态并行代码，该框架将阈值、粗化和聚合优化结合在一起。

动态并行的低效性观察。一些基准测试工作【索引16, Characterizing performance and power towards efficient synchronization of GPU kernels, 2016, MASCOTS】【索引36, NUPAR: A benchmark suite for modern GPU architectures, 2015, ICPE】【索引40, Characterization and analysis of dynamic parallelism in unstructured GPU applications, 2014, IISWC】观察到动态并行因高启动开销和硬件利用率不足而导致的低效性，这推动了各种硬件和软件优化的发展。

硬件优化。已提出许多硬件优化来减轻动态并行的开销。例如，动态线程块启动（DTBL）【索引37, Acceleration and optimization of dynamic parallelism for irregular applications on GPUs, 2016, Ph.D. dissertation】【索引38, Dynamic thread block launch: A lightweight execution mechanism to support irregular applications on GPUs, 2015, ISCA】、LaPerm【索引39, Laperm: Locality aware scheduler for dynamic parallelism on GPUs, 2016, ISCA】、SPAWN【索引34, Controlled kernel launch for dynamic parallelism in GPUs, 2017, HPCA】和LASER【索引35, Quantifying data locality in dynamic parallelism in GPUs, 2018, POMACS】。这些硬件优化对未来几代GPU很有前景，但在当前GPU上不可用，这激发了对软件优化的需求。此外，如果实现，硬件优化可能与我们提出的软件优化具有协同作用。

编译器/软件优化。已经提出了许多编译器/软件优化来提高动态并行的性能或提供动态并行的替代方案。CUDA-NP【索引42, CUDA-NP: Realizing nested thread-level parallelism in GPGPU applications, 2014, PPoPP】和Free Launch【索引9, Free launch: optimizing GPU dynamic kernel launches through thread reuse, 2015, MICRO】等方法通过避免动态并行来减轻其开销，但需要线程处于待命状态。

聚合技术。Li等人【索引24, Exploiting dynamic parallelism to efficiently support irregular nested loops on GPUs, 2015, CoSMiC】【索引25, Nested parallelism on GPU: Exploring parallelization templates for irregular loops and recursive computations, 2015, ICPP】、Wu等人【索引41, Compiler-assisted workload consolidation for efficient dynamic parallelism on GPU, 2016, IPDPS】和KLAP【索引13, Techniques for optimizing dynamic parallelism on graphics processing units, 2018, Ph.D. dissertation】【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】是聚合技术，其中多个子网格被合并为一个聚合网格以减少启动开销。我们利用了其中之一，KLAP【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】，作为我们流程中的聚合组件。

KLAP中的其他优化。KLAP【索引14, KLAP: Kernel launch aggregation and promotion for optimizing dynamic parallelism, 2016, MICRO】还包括另一种动态并行优化，即提升（promotion），它针对单块内核递归调用自身的特定模式。我们的优化不适用于此模式。

其他不规则并行框架。已提出各种框架来优化GPU上不规则并行应用的执行，例如Wireframe【索引1, Wireframe: supporting data-dependent parallelism through dependency graph execution in GPUs, 2017, MICRO】、Juggler【索引6, Juggler: a dependence-aware task-based execution framework for GPUs, 2018, PPoPP】、ATA【索引17, Adaptive task aggregation for high-performance sparse solvers on GPUs, 2019, PACT】和BlockMaestro【索引2, BlockMaestro: Enabling programmertransparent task-based execution in GPU systems, 2021, ISCA】。我们的工作侧重于优化使用动态并行表示的GPU嵌套并行应用。

A5 结论

本文提出了一个开源的编译器框架，用于优化嵌套并行应用中动态并行的使用。该框架包含三个关键优化：阈值、粗化和聚合。我们的评估表明，与不使用动态并行或仅使用先前工作中的聚合优化相比，我们的编译器框架显著提高了使用动态并行的嵌套并行应用的性能。

A6 附录

A. 摘要

我们的工件是一个用于优化使用动态并行应用的编译器，其工作流程如图8(a)所示。我们在Clang【索引23, LLVM: A compilation framework for lifelong program analysis & transformation, 2004, CGO】中实现了该编译器，并已将编译器代码公开。由于构建编译器需要构建Clang/LLVM，这可能耗时耗资源，因此我们在一个Docker镜像中提供了编译器的预构建二进制文件，以及所需的依赖项和已评估的基准测试/数据集。审阅者可以使用编译器二进制文件来转换基准CUDA代码，然后在支持CUDA的GPU上编译并运行代码，以验证第VIII节中报告的计时/加速结果。我们提供了自动化此过程的脚本。

B. 工件清单（元信息）

C. 描述

  1. 如何访问：工件可在https://http://doi.org/10.6084/m9.figshare.17048447.v1 。
  2. 硬件依赖：需要支持CUDA的设备。
  3. 软件依赖：我们在安装了CUDA-9.1的环境中运行了测试。软件依赖已打包在Docker镜像中。
  4. 数据集：数据集列于表I中，并已包含在工件中。

D. 安装

工件包含一个带有安装说明的README文件和一个用于处理安装的便捷脚本。

E. 实验工作流程

要运行实验，请运行Docker镜像：‘./http://run.sh‘。在Docker容器内，要使用默认参数编译所有二进制文件，请切换到‘test/‘目录中相关基准测试的目录并运行‘make all‘。要运行具有最佳配置的基准测试，请运行‘http://bestcombination.sh‘。要执行详尽搜索，请运行每个基准测试目录中可用的‘http://sweep.sh‘。

F. 评估和预期结果

运行‘http://bestcombination.sh‘应提供用于报告图9和图12中加速比的执行时间。运行‘http://sweep.sh‘应提供用于报告图11中加速比的执行时间。

G. 实验定制

可以通过为每个基准测试运行所需的参数来定制实验。参数可以在每个基准测试提供的Makefile中更新。