A Compiler Framework for Optimizing
Dynamic Parallelism on GPUs

Fig. 1(a) shows an example of how dynamic parallelism can be used in practice. In this example, parent threads executing on the GPU each discover some nested work that can be parallelized. Each parent thread launches a child grid to perform the nested work in parallel, and each parent thread provides its child grid with a different set of launch configurations and parameters. The amount of nested work may vary across threads. Hence, the child grids have different sizes. One source of inefficiency that may arise when using dynamic parallelism in this way is that a massive number of child grids may be launched, and many of them may be small in size [40]. In this case, the large number of child grid launches causes congestion, and the small size of the child grids causes the device to be underutilized.

Aggregation is an optimization that consolidates or aggregates child grids that are launched by multiple parent threads into a single grid. Various works [24, 25, 41, 14] apply this optimization either manually or in the compiler. Fig. 1(b) shows how aggregation can be applied to the example in Fig. 1(a). In this example, the parent threads coordinate to find the cumulative size of all their child grids in order to launch a single aggregated grid. In prior work, the scope of coordination has been across parent threads in the same warp, block, or grid. If the scope is across a warp or a block, one of the participating parent threads launches the aggregated grid on behalf of the others. If the scope is across a grid, the aggregated grid is launched from the host. The scope of coordination is referred to as the aggregation granularity.

In the original code, each parent thread may provide different parameters and launch configurations to its child grid. However, in the transformed code, only one set of parameters and launch configurations can be provided. For this reason, before launching the aggregated grid, the parent threads each store their original parameters and launch configurations in memory, and a pointer to this memory is passed to the aggregated grid. The child threads must then identify who their original parent thread is in order for them to load the right parameters and configurations from memory. To do so, each child thread block executes a search operation. The work done by the parent threads to identify the size of the aggregated grid and store their individual parameters and launch configurations in memory is referred to as the aggregation logic. The work done by the child threads to identify their original parent thread and load their original parameters and configurations from memory is referred to as the disaggregation logic.

The advantage of aggregation is that it reduces the number of child grids launched, thereby reducing congestion, and it increases the sizes of the child grids to ensure better hardware utilization. The disadvantage of aggregation is that it incurs overhead due to the aggregation and disaggregation logic, and it delays the child grid launches until all parent threads are ready to launch. The choice of aggregation granularity (warp, block, or grid) involves making a trade-off between these advantages and disadvantages. Using a larger granularity reduces congestion and increases utilization, but incurs higher overhead from the aggregation and disaggregation logic and delays the child grids longer before launching them.

Thresholding is an optimization where a child grid is only launched dynamically if the number of child threads exceeds a certain threshold. Otherwise, the child threads are executed sequentially by the parent thread. Fig. 2 illustrates how thresholding can be applied to the example in Fig. 1(a). In this example, two of the parent threads have a small number of child threads. The benefit gained from parallelizing these child threads is unlikely to be worth the launch overhead. For this reason, the parent threads instead execute the work of the child threads sequentially.

Thresholding is commonly applied manually by programmers [24, 25, 41, 34]. However, manual application of thresholding complicates the launch code, requires code duplication, and hurts code readability. For this reason, we propose to automate thresholding via a compiler transformation (Section III-B) and discuss the challenges associated with doing so (Sections III-C and III-D).

Fig. 3 shows an example of how our compiler applies the thresholding transformation. The original code in Fig. 3(a) consists of a parent kernel (lines 04-08) and a child kernel (lines 01-03). The parent kernel calls the child kernel using dynamic parallelism (line 06) and configures it with a grid dimension gDim and a block dimension bDim.

Fig. 3(b) shows the code after the thresholding transformation is applied. The transformation consists of two key parts: constructing a serial version of the child to be executed by the parent thread (lines 09-15) and applying a threshold to either perform the launch or call the serial version (lines 21-26).

To construct the serial version, the child kernel is replicated and its attribute is changed to __device__ so that it becomes a device function (line 09). Two parameters are appended to the parameter list: _gDim, which represents the original grid dimension of the parallel version, and _bDim, which represents the original block dimension. Loops are inserted around the child body (lines 10-11) to serialize the child threads. The first loop (line 10) iterates over the child thread blocks while the second loop (line 11) iterates over the child threads in each child block, using the bounds passed as parameters. Finally, all uses of the reserved index and dimension variables are replaced with the corresponding loop indices and bounds. The example shows a 1-dimensional child kernel for simplicity, however, if the child kernel is multi-dimensional, loops would be inserted for each dimension.

To apply the threshold, the number of child threads is first identified and stored in the _threads variable (line 21). Section III-D discusses how the number of child threads is identified. Next, an if-statement is inserted around the child kernel call (line 22) to ensure that the dynamic launch is performed only if the number of child threads is greater than or equal to _THRESHOLD. Here, _THRESHOLD is a macro variable that can be overridden at compile time for tuning purposes. If the threshold is not met, then the device function implementing the serial version is called instead (line 25), thereby serializing the child work in the parent thread.

Not all child kernels are amenable to the kind of transformation described in Section III-B. In particular, there are two kinds of child kernels that we do not transform: (1) child kernels that perform barrier synchronization across threads via __syncthreads() or warp-level primitives, and (2) child kernels that use shared memory.

For child kernels that perform barrier synchronization, serializing GPU threads while supporting such synchronization has been done in the literature [33, 20, 18, 21, 12]. The key strategy is to divide the code into regions separated by the barriers and insert loops around each region. However, these techniques target serializing multiple GPU threads in one CPU thread. Extending them to serialize multiple GPU threads in one GPU thread is not practical for two reasons. The first reason is that these techniques perform scalar expansion of all local variables to preserve the state of all threads across barriers. Such a scalar expansion on the GPU would convert all register accesses to memory accesses which would be prohibitively expensive. The second reason for not serializing child threads that perform barrier synchronization is that code that includes barrier synchronizations often implements a parallel algorithm that is not efficient when serialized. For example, a parallel reduction operation uses a reduction tree and leverages barriers to synchronize between levels of the tree. However, reduction trees are not an efficient way to perform sequential reductions. It is more efficient to use a simple reduction loop. In this case, it is better to let the programmer apply the thresholding optimization manually because the best sequential and parallel algorithms are different.

For child kernels that use shared memory, we do not construct a serial version of the kernel because every parent thread would require as much shared memory as an entire child block which would make the shared memory requirements of a parent block too high. Besides, kernels that use shared memory most often use __syncthreads() to coordinate access to shared memory across threads. Hence, these kernels will most likely not be transformable anyway for the reasons related to barrier synchronization previously mentioned.

The transformation described in Section III-B needs to identify the desired number of child threads in order to compare it with the threshold. Identifying the desired number of child threads is challenging because this information is not what the programmer provides in the kernel call. Instead, the programmer provides the grid dimension (number of blocks) and the block dimension (number of threads per block). The programmer uses the desired number of child threads in calculating the grid dimension.

One way to identify the desired number of child threads is to multiply the grid dimension with the block dimension. However, this approach gives the total number of threads in the child grid including threads that will not be used. This value is not a representative value to compare with the threshold. For example, consider a kernel with a nested kernel call where the child block dimension is configured to 1024 threads. One of the parent threads identifies 2 units of nested parallel work to be processed by the child. This parent thread will configure its child kernel with 1 block. In this case, multiplying the grid dimension with the block dimension gives 1024 total threads which is much larger than the actual number of threads desired (2 threads). Ideally, the value that should be compared to the threshold is 2, not 1024. Hence, multiplying the grid dimension with the block dimension is not a good approach.

The approach we use to identify the desired number of child threads is based on the observation that programmers usually calculate the grid dimension as a ceiling-division of the desired number of threads by the block dimension. There are arbitrary ways in which a ceiling-division can be expressed so it is not possible to have a static analysis that always determines the desired number of threads with certainty. Instead, we identify the most common patterns used by programmers to calculate ceiling-division and employ a simple static analysis based on these patterns.

Fig. 4 shows common expressions that programmers write for calculating the grid dimension using ceiling-division. Options (a)-(c) use integer arithmetic, while options (d)-(e) convert to floating point and use the ceil function. Option (f) is used in multi-dimensional blocks, where the operands to the dim3 constructor could each be an expression that resembles options (a)-(e). For all options, the expression may be expressed as a whole, or it may be expressed in parts where subexpressions are stored in intermediate variables. Note that N and b can be arbitrary expressions.

We observe from the examples in Fig. 4 that N is usually in the subexpression on the left hand side of the division. Moreover, the subexpression containing N may also contain constants such as 1 or b (which is usually a constant). Based on this observation, our analysis pass looks for a division operation, takes the subexpression on the left hand side, and removes additions and subtractions of constants, considering the remaining subexpression as the desired number of threads. This analysis is heuristic by nature and is not guaranteed to find the true desired number of threads. However, using a heuristic is acceptable in this context because the result will only be used to choose whether to serialize or parallelize the work. This choice does not impact program correctness in any way.

The subexpression that is found is assigned to _threads in Fig. 3 on line 21. The occurrence of the subexpression in gDim is then replaced with _threads to ensure that the expression is not duplicated in the code just in case the expression has side effects.

Coarsening is an optimization where the work of multiple thread blocks in the original code is assigned to a single thread block. The number of original thread blocks assigned to each coarsened thread block is referred to as the coarsening factor. When the GPU hardware is oversubscribed with more thread blocks than it can accommodate simultaneously, the hardware serializes these thread blocks, scheduling a new one whenever an old one has completed. There are multiple advantages of serializing these thread blocks in the code rather than letting the hardware do it. First, it reduces the number of thread blocks that need to be scheduled, and allows some warps in the coarsened block to proceed to executing the work of the next original block before other warps have completed their part in the previous original block. Second, if there is common work across the original thread blocks, that work can be factored out and executed once by the coarsened thread block, allowing its cost to get amortized. The disadvantage of coarsening is that it reduces parallelism, thereby underutilizing the device if the coarsening factor is too high.

Coarsening as an optimization is often applied by programmers in many different contexts [22]. Prior works have also applied coarsening in the compiler [26, 32], but not in the context of dynamic parallelism. We propose to apply coarsening in the context of dynamic parallelism, and automate its application via a compiler transformation.

Fig. 5 illustrates how coarsening can be applied to the child thread blocks in the example in Fig. 1(a). In this example, each coarsened child thread block in the transformed code executes the work of two child thread blocks in the original code. The advantage of applying coarsening in the context of dynamic parallelism is that it reduces the number of child thread blocks that need to be scheduled. More importantly, when applied before aggregation, coarsening also has the advantage of providing child thread blocks that do more work per block, which makes them more capable of amortizing the disaggregation logic overhead.

Fig. 6 shows an example of how our compiler applies the coarsening transformation to the original code in Fig. 3(a). The transformation consists of two key parts: coarsening the child kernel (lines 01-05) and modifying the launch configurations to launch the coarsened child (lines 08-10).

To coarsen the child kernel, a parameter _gDim is appended to the parameter list (line 01) which represents the original grid dimension without coarsening. A coarsening loop is inserted (line 02) that iterates over the work of the original child thread blocks assigned to the coarsened block. Uses of the reserved index and dimension variables are replaced with the corresponding loop indices and bounds. The example shows coarsening in one dimension only for simplicity, however, if the child grid is multi-dimensional, loops would be inserted for each dimension.

To modify the launch configurations to launch the coarsened child, the original grid dimension gDim is stored in a variable _gDim (line 08). The value is also copied to _cgDim, which represents the coarsened grid dimension. The x-dimension of the coarsened grid dimension _cgDim is then ceiling-divided by the coarsening factor _CFACTOR (line 09). Here, _CFACTOR is a macro variable that can be overridden at compile time for tuning purposes. Again, the example shows coarsening in one dimension for simplicity. Finally, the child kernel is configured with the coarsened grid dimension, and the original grid dimension is passed as a parameter (line 10).

Recall from Section II-B that prior work has performed aggregation at warp, block, and grid granularity. Using a larger granularity reduces congestion and increases hardware utilization, but incurs higher overhead from the aggregation and disaggregation logic and delays the child grids longer before launching them. The choice of aggregation granularity presents a trade-off between these performance factors. However, there is a wide gap between the block and grid granularity aggregation schemes that leaves a large part of the trade-off space unexplored. A parent thread block typically has hundreds to up to 1,024 parent threads, whereas a parent grid can have many thousands to millions of parent threads. To better explore the trade-off space, we propose an intermediate granularity between block and grid granularity, which is multi-block granularity aggregation.

In multi-block granularity aggregation, we divide a parent grid into groups of blocks, each group having a fixed number of blocks. The threads within the same group of blocks collaborate to aggregate their child grids into a single aggregated grid. In prior work [14], the aggregation logic involves a scan operation on the original grid dimensions and a max operation on the block dimension to identify the aggregated grid configuration. It also involves a barrier synchronization to wait for all participating threads to store their individual configurations and arguments to memory before the aggregated launch is performed. Since we cannot synchronize across multiple thread blocks, we perform the scan and max operations for multi-block granularity using atomic operations similar to what is done at grid granularity in prior work. As for the barrier synchronization, we replace it with a group-wide counter that is atomically incremented by each thread block when it finishes. The last thread block in the group to increment the counter performs the aggregated launch.

Fig. 7 shows an example of how our compiler applies the multi-block granularity aggregation transformation to the original code in Fig. 3(a). The transformation consists of two parts: the aggregation logic in the parent kernel (lines 14-35) and the disaggregation logic in the child kernel (lines 01-11).

For the aggregation logic in the parent kernel, we save the gDim and bDim expressions in temporary variables to avoid recomputing them every time we use them in case they have side effects (lines 14-15). We then identify the group that the thread block belongs to (line 16) and find the group’s memory segment in a pre-allocated memory buffer (line 17). This memory segment is used to store the configurations and arguments that threads in the group pass to their children. Next, each thread that launches a child grid (has a non-zero grid dimension) atomically increments two global counters simultaneously (lines 19-20): (1) _numParents to assign an index to the parent thread so that the thread knows where to store its arguments and configuration, and (2) _sumGDim to find the total number of child blocks of prior parent threads which we use to initialize the scanned array of grid dimensions. The two global counters are incremented simultaneously by treating them as a single 64-bit integer. Each thread then stores its arguments, scanned grid dimension, and block dimension to memory so they can be passed to the child grid (lines 21-23). Each thread also performs an atomic operation to find the maximum block dimension (line 24).

After each thread writes its configuration and arguments to global memory, it performs a fence operation to ensure that the configuration and arguments are visible to its child blocks (line 26). This fence is not necessary in prior work because the visibility of the data is ensured either by the semantics of the dynamic kernel call (for warp and block granularity) or by grid termination (for grid granularity). However, it is necessary here because the launch may be performed by a different thread block than the one that writes the data. Hence, the writing thread block must ensure that its data is visible in global memory before notifying the other blocks that it is ready. A local barrier (line 27) is also needed to ensure that all threads in the block finish storing their data before notifying.

Finally, it is time to perform the aggregation launch. One thread in the block (line 28) atomically increments the group-wide counter (line 29) and checks if its block is the last block in the group to finish (line 30-31). If so, the thread launches the aggregated grid, configuring it with the sum of all grid dimensions and the maximum of all block dimensions, and passing pointers to the memory arrays containing the original arguments and configurations (lines 32-33).

For the disaggregation logic in the child kernel (lines 01-11), the code remains largely the same as prior work [14]. Each child block performs a binary search through the scanned grid dimension array to identify its original parent thread (line 02). It then loads its parameters and configuration (lines 04-06) and performs the work of the child kernel based on them (lines 07-10).

Note that besides filling the gap in the trade-off space between block and grid granularity aggregation, multi-block granularity aggregation has another advantage over grid granularity aggregation in particular. Grid granularity aggregation requires the CPU to be involved in the aggregated launch whereas the multi-block granularity aggregation runs entirely on the GPU, which frees the CPU to perform other tasks. Hence, multi-block granularity aggregation is more compatible with the asynchronous semantics of kernel calls. If the CPU is needed for performing other tasks, multi-block granularity aggregation may be better for overall execution time even if grid granularity aggregation has better kernel time.

When thresholding is applied before aggregation, the number of original child grids that participate in the aggregated grid may be substantially reduced. If there is an insufficient number of original child grids participating in the aggregated grid, the benefit of aggregation may not be worth its overhead. To address this issue, we enhance aggregation with another optimization that applies an aggregation threshold. The aggregation logic is preceded by an operation to count the number of participating parent threads. If the number of participating parent threads does not meet a certain threshold, the child grids are launched normally by their parent threads instead of being aggregated. Therefore, a child grid may be executed in one of three ways: it may be serialized within its parent, launched directly by its parent, or launched as part of an aggregated grid. Since applying an aggregation threshold requires parent threads to synchronize to count the number of participating threads, it can only be applied at warp and block granularity where barrier synchronization across threads is possible.

Fig. 9 shows the performance results for each benchmark and dataset with all combinations of optimizations applied. Performance is reported as speedup over the CDP version.

Applying CDP alone leads to a performance degradation in almost all cases compared to not applying CDP. Aggregation alone recovers from this degradation and substantially improves performance, where CDP+A is 12.1$\times$ faster than CDP and 2.4$\times$ faster than No CDP. These observations are consistent with those made in prior works [24, 25, 14, 41], which note that the large number of launches causes congestion and aggregating the grids mitigates this overhead.

Thresholding alone gives substantial speedup, where CDP+T is 13.4$\times$ (geomean) faster than CDP. Thresholding also gives speedup in the presence of aggregation, where CDP+T+A is 2.9$\times$ (geomean) faster than CDP+A, and CDP+T+C+A is also 3.1$\times$ (geomean) faster than CDP+C+A. The incremental benefit of thresholding with aggregation is not as pronounced as without aggregation because the benefit of thresholding is reducing the number of launches which aggregation also does. Nevertheless, the speedup is still significant.

Coarsening without aggregation gives modest speedup, where CDP+C is 1.01$\times$ (geomean) faster than CDP and CDP+T+C is 1.09$\times$ (geomean) faster than CDP+T. On the other hand, coarsening with aggregation gives more speedup, where CDP+C+A is 1.16$\times$ (geomean) faster than CDP+A, and CDP+T+C+A is 1.22$\times$ (geomean) faster than CDP+T+A. Notice how coarsening is synergistic with aggregation (its speedup in the presence of aggregation is greater than its speedup in the absence of aggregation). The reason is that in the presence of aggregation, coarsening helps amortize the disaggregation logic across more work in the coarsened child block in contrast with much less work in each original child block. Although the benefit of coarsening is much smaller than that of thresholding and aggregation, it is still substantial. A speedup of 1.22$\times$ is significant considering that it comes from a compiler optimization that requires no programmer intervention or additional architecture support.

Ultimately, our compiler framework as a whole substantially improves performance with the three optimizations combined, where CDP+T+C+A is 43.0$\times$ (geomean) faster than CDP. CDP+T+C+A is also 8.7$\times$ (geomean) faster than No CDP, showing that dynamic parallelism is a powerful programming feature when combined with the right optimizations. Moreover, CDP+T+C+A is 3.6$\times$ (geomean) faster than KLAP (CDP+A), showing that our framework provides substantial speedup over prior works that perform aggregation alone [14].

Fig. 10 shows how the execution time is spent for each benchmark and dataset. We use KLAP (CDP+A) as baseline instead of CDP because prior work [14] has already shown how aggregation affects the execution time breakdown, and including CDP makes the figure illegible. We compare with CDP+T+A and CDP+T+C+A to show the incremental impact of thresholding and coarsening, respectively.

The first observation is that thresholding increases parent work and decreases child work. Thresholding serializes child work in parent threads so it is natural for parent threads to have more work to do and child threads to have less.

The second observation is that thresholding decreases the overhead from aggregation, launching, and disaggregation. Thresholding results in fewer parent threads launching and fewer child threads being launched. The aggregation overhead decreases because fewer parent threads participate, the launch overhead decreases because there are fewer launches, and the disaggregation overhead decreases because there are fewer child threads searching for their parents.

The third observation is that coarsening decreases the launch overhead. Since coarsening reduces the number of child blocks that need to be scheduled, the launch overhead decreases.

The fourth observation is that coarsening decreases the disaggregation overhead. Coarsening amortizes the disaggregation logic across multiple child thread blocks rather than having each child thread block perform its own disaggregation.

The final observation is that coarsening decreases parent work and increases child work for some benchmarks, namely BFS and SSSP. This result is unintuitive because coarsening does not change the work that the parent does. The reason is that for these two benchmarks, after coarsening is applied, the reduction in launch and disaggregation overhead results in a lower optimal threshold to be found. The lower threshold results in more work being offloaded from parent to child. This observation demonstrates an important interaction between thresholding and coarsening.

Fig. 11 shows how performance varies for different threshold values and aggregation granularity while keeping the coarsening factor constant at the best coarsening factor found. For space constraints, we only show the results for one dataset.

The first observation is that for most benchmarks (all except SP), as the threshold increases initially, performance also improves. Increasing the threshold initially results in more of the small child grids getting serialized in their parent threads, which reduces the number of launches, hence the congestion, and avoids launching small grids that underutilize the device.

The second observation is that for some benchmarks (e.g., BFS, BT, MSTF, SSSP), increasing the threshold too much causes performance to degrade again. Increasing the threshold too much results in large child grids getting serialized within their parent threads, which reduces parallelism and causes high control divergence.

The third observation is that different benchmarks perform best with different levels of aggregation granularity. For example, SP and TC perform best with grid granularity, BFS and SSSP perform best with multi-block granularity, BT and MSTF perform best with block granularity, and MSTV performs best without aggregation. The fact that multi-block granularity aggregation performs best in some of the cases reflects the importance of providing an intermediate granularity between grid and block granularity.

As mentioned in Section VII, we use an exhaustive search to perform tuning to show the maximum potential of the optimizations and to present a complete view of the design space. However, from our experience, such a broad search is unnecessary. First, the best threshold is typically the one that allows approximately 6,000-8,000 child grid launches. Second, performance is not very sensitive to the coarsening factor provided that it is sufficiently large ($>$8) so the coarsening factor does not need to be searched with a broad range or with high resolution. Third, aggregation at warp granularity is never favorable. With these observations in mind, users can typically find a combination of parameters that is very close to the best with less than ten runs. Moreover, the compiler framework exposes these parameters in a configurable manner to make it easy for users to leverage off-the-shelf autotuners [4]. Such tuning is worthwhile for kernels that run repeatedly on similar datasets. On the other hand, if a user cannot afford to tune the kernel, it is not necessary for the user to find the best parameters to benefit from the optimizations. Selecting a sub-optimal set of parameters would still yield a speedup, just not the maximum possible speedup. For example, if we fix the threshold to 128 for all benchmarks and datasets, then CDP+T+C+A will have a geomean speedup of 1.9$\times$ over CDP+C+A, as opposed to 3.1$\times$ when the best threshold is used. Hence, our optimizations are still useful even if the best parameters are not searched for or found.

Dynamic parallelism is useful when the amount of nested parallelism is high such that the launch overhead is worth the parallel work extracted. If the amount of nested parallelism is low for all parent threads, the benefit of dynamic parallelism is limited. To demonstrate the impact of dynamic parallelism on applications with a small amount of nested parallelism, we evaluate the graph benchmarks on a road graph (USA-road-d.NY [11]). The graph has 264,346 vertices, 730,100 edges, an average degree of 3, and a maximum degree of 8. Hence, each vertex has a small number of outgoing edges so processing the graph has a small amount of nested parallelism.

Fig. 12 shows the performance results for each graph benchmark on the road graph with all combinations of optimizations applied. It is clear that CDP versions perform substantially more poorly on this graph relative to the No CDP versions. Our proposed optimizations are able to recover much of the performance degradation, but not all of it. In fact, in this experiment, we tune the threshold beyond the largest launch size such that CDP+T degenerates to serializing all child threads like No CDP. However, CDP+T still cannot recover fully. The reason is that the mere existence of a dynamic launch in the code, even if it is never executed, results in a performance degradation. To verify this observation, we compare two kernels where the only difference is a dynamic launch that is guarded by a condition that is always false such that the launch is never performed. Upon inspecting the assembly code of the two kernels, we observe that a large number of additional instructions are generated besides the instructions for performing the launch. Upon profiling the execution of the two kernels, we observe that a large number of additional instructions are executed even though the launch is never performed.

Note that the SP benchmark on the RAND-3 dataset performs poorly in Fig. 9 also because the amount of nested parallelism is low (all child grids have fewer than 32 threads). These results show the importance of being aware of the application and dataset when choosing whether or not to apply dynamic parallelism and its associated optimizations.