Hardware Compute Partitioning on NVIDIA GPUs for Composable Systems

We assume an x86_64 or aarch64 platform containing at least one embedded or discrete NVIDIA GPU of the Volta (2018) generation or newer.

Tasks are assumed to be closed-source, unmodifiable, CUDA-using Linux processes. Non-CUDA GPU-using tasks may coexist, but may not use spatial partitioning.

We focus on embedded, real-time systems, where resources are limited but execution-time deadlines must be met.

Each task executing on a GPU has an associated GPU context. This context includes per-GPU-task state, such as an on-GPU virtual address space. Unless explicitly stated otherwise, we assume a one-to-one mapping of GPU-context to CPU-task.

Work is dispatched into a context via one or more compute kernels. A kernel is launched by passing a GPU-executable binary, and a number of GPU threads, to a GPU-usage library such as CUDA or OpenCL. The library will then compose and pass a Task Metadata Descriptor (TMD) to the GPU for execution. Each GPU thread concurrently executes the same instructions, but operates at a different data offset. For example, a kernel for an element-wise array addition could replace a for loop over every index value with a GPU thread per value. The threads would then execute the loop body over all index values in parallel, rather than having to iterate.

Threads are organized into groups known as blocks, and all threads within a single block execute on the same SM (Streaming Multiprocessor; a group of compute cores). A series of kernel executions may be serialized via first-in-first-out command queues known as streams in CUDA. Only kernel executions in the same stream are serialized; kernels in separate streams are permitted to execute concurrently.

In Fig. 1 we illustrate the typical architecture of an NVIDIA GPU chip, using the Ada-generation AD102 as an exemplar (used in the RTX 4090 and RTX 6000 Ada GPU models). NVIDIA GPUs are subdivided into several independent engines, with the most significant being the Compute/Graphics Engine. Smaller engines handle special functions such as bulk data movement (the Copy Engines) and video processing. Each engine may have a different context active on it at a time [6].

The Compute/Graphics Engine is subdivided into General Processing Clusters (GPCs). Each GPC is independently connected to each DRAM controller and the co-located last-level, level-two (L2) cache slices. The GPCs subdivide into Thread Processing Clusters (TPCs) of two Streaming Multiprocessors (SMs) each. Each SM contains dozens of compute cores (128 each on the AD102) and a level-one (L1) cache. The AD102 contains 18,432 compute cores total; Fig. 2 illustrates how core counts have increased in recent years.

We next discuss how GPU contexts are scheduled onto hardware engines.

We present the scheduling pipeline from CPU task to GPU compute cores in Fig. 3. CPU tasks insert kernel launch commands (represented as TMDs) into CUDA streams, and those CUDA streams are mapped onto a smaller or equivalent number of GPU channels.¹¹1Strictly, channels contain a queue called a pushbuffer. The channel count is important – using more streams than channels causes blocking [6]. Which GPU channel(s) may access the Compute/Graphics Engine at a time is dictated by the runlist, making the runlist the central arbitrator.

GPU runlists (typically one for each GPU engine [6]) are composed of time-slice groups (TSGs) of channels. TSGs are executed in a work-conserving, preemptive round-robin order by default. While a TSG is active, commands from all its contained channels are pulled and executed on the runlist-associated engine. Each TSG has an associated time-slice; upon expiration of this timeslice the runlist scheduler preempts any in-progress commands and switches to the next TSG. Such a switch is also triggered by the exhaustion of commands from all the TSG’s channels. Note that all a TSG’s channels must be from the same context.²²2See line 293 in manuals/ampere/ga100/dev_ram.ref.txt [31].

We show a runlist for three tasks in Fig. 4, following the in-memory layout. In this example, TSG 0 would be executed for 2 ms, and during that time, commands from channels 0 and 1 would be received and executed by the Compute/Graphics Engine. After 2 ms, any active commands from channels 0 and 1 would be preempted, and the GPU would move on to commands from the channels of TSG 1 for 2 ms. This process repeats for TSG 2, then loops back to TSG 0. The time-slice can be set differently for each TSG.

This GPU scheduling algorithm ensures that the Compute/Graphics Engine has only one context active at a time. This can lead to significant under-utilization, as many GPU-using tasks are unable to saturate all GPU compute cores on their own [43, 28, 47]. We demonstrate this for one inference of the YOLOv2 image-detection network in Fig. 5. At no point is the network able to utilize the entire GPU – it utilizes only 40% of the GPU’s SMs on average. Increasing GPU core counts (Fig. 2) have only worsened this problem.

To efficiently use the GPU, idle compute capacity must be reclaimed for other tasks. Unfortunately, concurrently running multiple tasks on the GPU (e.g., by sharing a single context) leads to a new set of problems.

When multiple tasks execute concurrently on a GPU, shared-resource interference can occur. This is where contention for shared hardware resources such as caches or compute cores creates slowdowns between tasks. Such interference creates unpredictability, as knowing when and how tasks will interfere on what shared resource is an unresolved area of study. Thus, to ensure timing predictability, few hardware resources should be shared.

Prior work has avoided such interference on the GPU by only allowing one task to access the GPU at a time [13, 12, 3]. This can lead to underutilization – concurrently running multiple tasks would allow for more efficient use of GPU hardware.

Spatial partitioning allows for concurrency without interference. It prevents hardware resources from being shared between tasks by partitioning them into mutually exclusive subsets, and assigning each subset to a concurrently running task. Without shared resources, interference is prevented, and execution times remain predictable.

In designing a spatial partitioning mechanism, it should satisfy the following properties:

Portable:

The mechanism should work on a wide set of GPU models.

Logically Isolated:

The mechanism should preserve logical isolation between tasks, e.g., virtual address space isolation and independent exception handling.

Transparent:

The mechanism should not require changes to tasks, e.g., recompilation.

Low-overhead:

The mechanism should add negligible overhead to critical-path operations, e.g., a kernel launch.

Hardware-enforced:

Partitions should be enforced by hardware to protect from malicious or misbehaving tasks.

Dynamic:

Partitions should be reconfigurable without restarting a task.

Granular:

Partitions should be defined in granular units, e.g., a TPC for compute partitioning.

The first four properties are desirable for any software system, but the last three are partitioning-specific. We now discuss prior work in light of these requirements.

The history of DRAM partitioning is extensive ([20, 21, 45] are exemplars), and it has been extended onto the GPU by Fractional GPUs [16] and SGDRC [46]. These works use a memory-organization approach known as page-coloring [20] to force each task onto prescribed GPU DRAM and L2 units. Such memory reorganization requires difficult-to-obtain model-specific hashing functions, but the principle is generally applicable to any GPU. NVIDIA has since developed a proprietary alternative, but this is not available on most GPUs and cannot be enabled independently of MiG [30, 10].

The DRAM partitioning technique of Fractional GPUs and SGDRC satisfies all desired spatial isolation properties stated in Sec. 2.5, so we assume the application of such an approach and focus on the remaining problem: compute partitioning.

Prior work on spatial partitioning for GPU compute cores can be divided into academic-provided and NVIDIA-provided solutions. We summarize these works in Table 1.

Unfortunately, devising better solutions for spatial partitioning of GPU compute has been stymied by the scarce details released about NVIDIA GPU architecture and capabilities – even instruction encodings are secret [15, 18, 17]. Many years of investigation have begun to contravene this limitation.

Otterness et al. [35] and Amert et al. [2] elucidated the scheduling of CUDA kernels via black-box experiments and introduced the cuda_scheduling_examiner tool, which we use. Olmedo et al. [33] and our prior work [5] used these results and other sources to construct a model of how the underlying GPU kernel dispatch hardware works. Parallel work by Capodieci et al. [7], Spliet and Mullins [39], and us [6] clarified the preemption and high-level scheduling capabilities of NVIDIA GPUs. This last work [6] also introduced two tools we use: the nvdebug tool for directly extracting GPU state, and the gpu-microbench suite for examining scheduling behavior. Outside of the academic community, the Nouveau [25] and Mesa [24] reverse-engineered NVIDIA GPU driver projects have documented GPU hardware capabilities and CPU-to-GPU interfaces. We lean heavily on all these prior works throughout our paper.

To determine how MPS interacts with the GPU scheduling pipeline, we applied the nvdebug tool [6], gpu-microbench suite [6], and cuda_scheduling_examiner toolkit [35] to test and observe GPU state and behavior with and without MPS. Adding MPS changed more aspects of GPU state than nvdebug was able to display, so we extended it on Ada (2022) and older GPUs, drawing layout information from open-source NVIDIA code [31, 27, 29] and the nouveau driver [25]. Our improved version of nvdebug is available.⁴⁴4Available at http://rtsrv.cs.unc.edu/cgit/cgit.cgi/nvdebug.git and within our artifact.

To understand the semantic meaning of this newly accessible GPU state, we leveraged context and definitions from NVIDIA patents touching on the runlist scheduler [11], kernel scheduling pipeline [37, 1], MPS [9] and MiG [10]. Unfortunately, patents may describe infeasible or impossible devices, and so we only tenuously relied on them, verifying described behavior with experiments.

MPS uses a client-server paradigm, where each CUDA-using task is a client of at most one MPS server task. The MPS server acts as an intermediary to the GPU for its clients, and allows clients to run concurrently with one another. Relative to the GPU, a system of two MPS clients and one MPS server would appear as a single task, since clients only access the GPU through the MPS server. Multiple MPS servers may exist, each with different clients.⁵⁵5Support for multiple MPS servers is only implicitly documented. The environment variable CUDA_MPS_PIPE_DIRECTORY can be set to control where an MPS server advertises itself, and where CUDA searches for the MPS server. Two MPS servers cannot advertise to the same path, and so this variable must be set uniquely for each to allow multiple servers to exist. In this case, each MPS server would appear to be a separate task to the GPU.

We now investigate how adding MPS modifies arbitration between GPU-using tasks, i.e., how it modifies the runlist and associated data structures. To enable our subsequent analysis of MPS’s pitfalls, this section is particularly detailed.

We begin by reillustrating Fig. 4, with detail from our extensions to nvdebug, in Fig. 6. This figure retains the same runlist as Fig. 4, but expands on the configuration of each channel. The channel-configuration data structure is known as the channel’s instance block, and – besides describing the command queue (not shown for space) – specifies the virtual address space to be used for commands from the channel. Virtual address spaces are configured in a peculiar way; each channel includes an indexed list of page tables, and the page table to use is selected by specifying an index into that list. The list of page tables is called the “Subcontext Table,” and the channel’s index into it is called the channel’s “Subcontext ID.”⁶⁶6Subcontext ID is also “Virtual Engine ID” (VEID) in some sources. Adding to the confusion, this mechanism only selects the page table for the Compute/Graphics Engine; the instance block includes a separate page table configuration field for other engines, such as Copy.⁷⁷7See line 297 in manuals/ampere/ga100/dev_ram.ref.txt [31]. In our experiments, we always observed the same page table configuration for all engines; we mirror that in Fig. 6. Note that all channels in a TSG have an identical subcontext table – this is a requirement for channels sharing a context.⁸⁸8See line 293 in manuals/ampere/ga100/dev_ram.ref.txt [31]. All of these details are important in order to discuss how MPS effects address-space isolation.

In Fig. 7, we illustrate how runlist and channel configurations change when MPS is enabled. This is still a system of the same three GPU-using tasks, but an MPS server has been started with Tasks 2 and 3 as clients (Task 1 remains independent of MPS). Both the runlist location and virtual address space configuration have changed for the tasks now running as MPS clients. We discuss each change in turn.

With MPS enabled, the two tasks now running as MPS clients no longer have independent TSGs in the runlist. The runlist only contains two TSGs: one for Task 1 (TSG 0), and one for the MPS server (TSG 1). While Tasks 2 and 3 do not retain independent TSGs, they do retain independent channels within the MPS server’s TSG (Channels 3–4 for Task 2 and 5–6 for Task 3). (Note that the MPS server has taken Channel 2 for its own use, and so Tasks 2 and 3 are forced to use channels with higher IDs.) When this runlist is scheduled, Task 1 will, as before, execute for 2 ms before its budget expires and it is preempted. The GPU will then switch to the next TSG in the runlist, the one for the MPS server. This will likewise be run for 2 ms before the GPU loops back to Task 1. While the MPS server’s TSG is active, commands from all its channels are sent to the Compute/Graphics Engine concurrently. The effect is such that all the MPS client tasks’s CUDA streams concurrently execute as though they were in the same program. Note that the two MPS clients share a single 2 ms time-slice with a 4 ms period, whereas before they each got a single 2 ms time-slice at a 6 ms period. This is a reduction in GPU time available jointly to the two tasks, from two-thirds to one-half. This is a side-effect of enabling MPS to be aware of: if any other tasks in the system continue to run without MPS, the total GPU time available to all MPS clients will be less than the total time those tasks would have collectively received if run apart.⁹⁹9The time-slice length of the MPS server could be extended to ensure tasks retain access to an equivalent fraction of GPU time, but this is unsupported by NVIDIA’s software.

We now consider how virtual address space configurations change – or stay the same – with the addition of MPS. Visible in Fig. 7, at the bottom of the Channel 2–6 Instance Blocks, the Subcontext Table for every channel of the MPS server TSG is triple the size of the tables in Fig. 6. This allows the table to include separate page tables for MPS, Task 2, and Task 3; the appropriate one is selected by the Subcontext ID on each channel. Each channel also has a non-compute page table identical to the one selected by its Subcontext ID. In this manner, Tasks 2 and 3 retain distinct virtual address spaces, just as without MPS.¹⁰¹⁰10Why the convoluted TSG-wide Subcontext Table, rather than a per-channel page table? This may speed up context switches by allowing a TSG’s page tables to be read all at once, rather than requiring a scan of all a TSG’s channels.

What does this mean for MPS’s suitability to real-time embedded systems? We answer by considering MPS under each of our desired properties (Sec. 2.5). Without loss of generality, we assume that all MPS clients are associated with a single server in this subsection.

nvtaskset works like Linux’s taskset tool. While taskset sets the CPU cores that any task may use, nvtaskset sets the GPU cores that a task may use. To co-run two unmodified DNN tasks on the GPU with nvtaskset, one would run the shell commands:

The partitions could later be resized:

And tasks started without nvtaskset could later be moved to a partition:

Partitions can also be specified in terms of GPCs instead of TPCs (if nvdebug is loaded):

nvtaskset is distributed as an extension to libsmctrl, has no dependencies beyond CUDA, and can be built and installed via a simple make install invocation.

nvtaskset uses MPS to allow tasks to co-run, uses parts of libsmctrl to enforce partitioning, uses shared-library interception to apply to every CUDA-using task, uses shared memory to communicate dynamic partition updates, and uses GPU topology registers to associate GPCs with TPCs.

nvtaskset has several limitations inherited from MPS and libsmctrl. nvtaskset:

• $\mathrm{■}$

  Is only compatible with CUDA-using tasks.

• $\mathrm{■}$

  Does not support tasks that use multiple GPUs.

• $\mathrm{■}$

  Cannot affect partition changes on already-launched kernels.

• $\mathrm{■}$

  Is still subject to Pitfalls 1–3 of MPS.

• $\mathrm{■}$

  Only supports up to 15 co-running tasks.²²²²22TSGs are limited to a maximum of 128 channels [31], so an MPS server can only service 15 clients of 8 channels each (after subtracting 8 MPS-server-internal channels). We verified this 15-client limit on both Volta- and Ada-generation GPUs.

nvtaskset is also subject to Pitfall 5 by default. While the $\%smid$ register can be restored to its consistent, non-MPS behavior by toggling off the WDU’s dynamic partitioning register, we chose not to do this. To better support tasks which ignore the CUDA programming model and attempt to execute work on self-selected SMs (such as persistent threads [44, 14]), we can “hide” the existence of unallocated SMs with the WDU’s dynamic partitioning capability. To do this, leave on fake SM IDs, and set MPS’s partition limit to match the number of SMs allocated by nvtaskset. This can ensure that tasks only see allocated SMs, with seemingly-contiguous SM IDs.

However, nvtaskset is portable to any recent discrete NVIDIA GPU, does not require task modifications, preserves address space isolation, and is available as open-source software.²³²³23Available at http://rtsrv.cs.unc.edu/cgit/cgit.cgi/libsmctrl.git and within our artifact.

We measure startup overhead and launch overheads for all mechanisms via benchmarks run in the 57% partition. Startup overhead is the time from exec() to first kernel running on the GPU of a minimal program, and launch overhead is the time from cuLaunchKernel() to the kernel running on the GPU. We added these benchmarks to the gpu-microbench suite [6],²⁵²⁵25Available at http://rtsrv.cs.unc.edu/cgit/cgit.cgi/gpu-microbench.git and within our artifact. and ran each experiment once to prime caches before gathering the data displayed in Fig. 11.

As shown in Fig. 11(a) and Fig. 11(b), no partitioning mechanism worsens observed worst- or average-case startup or launch times.²⁶²⁶26Fig. 11(b) omits 100th percentile (max) times, as use of Linux’s isolcpus, sched_yield() and nohz=full options were insufficient to eliminate all noise, potentially due to SMIs [22] or interrupts.

Uniquely, MiG statically binds a hardware scheduling pipeline to each partition’s TPCs [10] such that only a subset of TPCs have to be set up, and considered as part of a kernel launch. We suspect this is what lowers launch and startup overheads with MiG.

With MPS, newly launched tasks (MPS clients) only initialize channels and subcontexts, with the MPS server providing the parent context and TSG. This appears to significantly reduce startup overheads for MPS and the MPS-based nvtaskset, more than compensating for the added client-server communication cost of MPS.

To evaluate how strongly partitions are enforced, we measured the execution time of a $6144\times 6144$ matrix multiply (yielding $36\times 2^{10}$ blocks of 1024 threads) (“MM6144”) in a 57% partition while interfering tasks executed on the remainder of the GPU. We used the mandelbrot and random_walk tasks from the cuda_scheduling_examiner toolkit as compute- and memory-heavy interfering tasks respectively. Interfering tasks executed continuously, out of sync with each other and the matrix multiply. We carefully configured this experiment to sidestep the pitfalls of MPS discussed in Sec. 4.4, focusing exclusively on how well a partition is enforced in an otherwise-ideal scenario. Good partition enforcement means that the execution time distribution of our MM6144 task does not shift in the presence of interfering tasks.

The results against memory- and compute-heavy interference are shown in Fig. 12(a) and Fig. 12(b) respectively. Each plot includes a baseline “Scaled” time for comparison – this is not a measured value, but is a scaled-up value derived from the execution time of the benchmark without any partitioning or interference. For example, if MM6144 took 244 ms when run alone on the GPU, the “Scaled” value for a 57% partition would be $144/0.57=253$ ms. This value is the minimum execution time possible for the MM6144 task in a 57% partition before accounting for sympathetic caching effects.

We specifically choose this matrix multiply task as it was most sensitive to interference among several other synthetic tasks we tested, and it could be more-precisely timed than a full neural network. Since we use the same partitioning mechanism as libsmctrl, the enforcement results previously shown for real tasks [5] continue to apply.

When competing work is memory-bound, as in Fig. 12(a), MPS does nothing to prevent memory contention, only limiting the compute available for MM6144. This results in average (line in box) and worst-case (top of whisker) execution times for our MM6144 task only slightly better than with no partitioning at all.

For memory-bound competing work (Fig. 12(a)), MiG beats nvtaskset, likely because MiG also partitions DRAM. Interestingly, even though nvtaskset includes no explicit memory partitioning, its implicit partitioning of the L0, L1, and TLB caches by aligning partitions to GPCs appears to be enough to beat MPS and approach MiG’s performance.

Surprisingly, for compute-bound competing work (Fig. 12(b)), nvtaskset beats MiG (and every other approach) on average- and worst-case execution times – without the hardware modifications of MiG. Upon further investigation, we found that MiG compromises compute speed to serve other design goals. This unexpected behavior is not documented, and causes MiG to fail in other ways, as we will explore under Obs. 7.

In all, nvtaskset enforces partitioning much better than MPS, and even better than MiG in some cases – all without requiring task, driver, or hardware modification.

To evaluate partitioning granularity, we recorded the total execution time of a $8192\times 8192$ matrix multiply (yielding $64\times 2^{10}$ blocks of 1024 threads) (“MM8192”) at every possible partition size for each partitioning method and plot the results as points in Fig. 13. The lines in Fig. 13 represent the closest available configuration which allocates at most the specified number of TPCs. For example, there is no MiG configuration for 10 TPCs, so we plot time for the closest available allocation of no more than 10 TPCs – 7 TPCs.

Both MPS and nvtaskset can specify partition sizes down to the per-TPC level, visible as the different MM8196 execution times for each setting in Fig. 13. However, nvtaskset can assign specific TPCs to a partition, in contrast to MPS’s generic percentage value. For this 54-TPC GPU, that means nvtaskset supports a total of $2^{54}$ different partition settings per-task, MPS supports 54 per-task, and MiG only supports 5 per-task.

Visible in Fig. 13, the largest-available MiG partition contains only 49 TPCs, whereas MPS or nvtaskset are able to access up to 54 TPCs. Upon further investigation, we find that no configuration of MiG partitions on the A100 can access more than 49 TPCs total. Every partition size is an even multiple of 7 TPCs, and 7 does not divide 54 evenly, wasting the remainder – 5 TPCs. This surprising issue is not documented, and means that enabling MiG immediately and inherently disables 9% of the A100 GPU. Concerningly, we found that this issue is even worse on NVIDIA’s newer GPUs. On the H100 GPU (SXM-80GB version tested), we found a loss of 6–15% (depending on the MiG partition size chosen).²⁷²⁷27We suspect the capacity loss stems from a decision to make unit-size MiG slices appear identical, despite differences in the underlying GPCs. Due to floorsweeping, some GPCs will have more working TPCs than others – those TPCs must be disabled to emulate identical GPCs when using MiG.