SubGPU Pipelining: Intra-GPU Fine-Grained Spatial Parallelism for LLM Inference

SubGPU Pipelining: Intra-GPU Fine-Grained Spatial Parallelism for LLM Inference Muralidhar Andoorveedu

Marius Juston Tingkai Liu Sanjay Patel NVIDIA {mandoorveedu, mjuston, tingkail, sanjpatel}@nvidia.com

arXiv:2611.14904v1 [cs.DC] 11 Nov 2026

November 2026

Abstract

step can execute dozens of GPU kernels per transformer layer—normalization, attention, projection, routing, expert computation, residual addition, and more—each with distinct compute, memory, and communication requirements. Modern GPUs such as the NVIDIA B200 [18] are powerful enough that individual decode kernels at small batch sizes complete in microseconds, making launch overhead, idle compute cycles, and wasted HBM bandwidth important effects that collectively leave significant performance on the table. Standard approaches—tensor parallelism (TP), data parallelism (DP), expert parallelism (EP), context parallelism (CP), and pipeline parallelism (PP)—scale the inference problem across GPUs, and do not directly address the intra-GPU inefficiencies just described. Specifically, three pathologies remain regardless of how many GPUs are added:

Large language model (LLM) decode is dominated by short, predominantly memory-bound kernels that make bursty use of on-chip execution resources and memory bandwidth. This paper introduces SubGPU Pipelining, a fine-grained, intra-GPU spatial execution model that cuts an operator graph into ordered contiguous segments, binds those segments to spatial stages backed by disjoint streaming-multiprocessor (SM) pools, and streams microbatches through the resulting schedule. Selecting the graph cuts, segment-to-stage binding, per-stage SM allocation, and microbatch size is a coupled optimization problem: each choice changes the overlap pattern, shared-resource contention, cache state, and execution times of the others. We evaluate SubGPU Pipelining through both simulation and silicon experiment using the GPT-OSS 120B workload with 32-GPU B200 deployment settings. The analytical shared-resource simulator searches across SubGPU pipeline structures and batch sizes to map the per-user/perGPU throughput Pareto frontier. At matched per-GPU throughput, the modeled SubGPU frontier provides up to 2.2× higher per-user throughput than temporal execution, and at matched per-user throughput it provides up to 1.9× higher per-GPU throughput. To test physical feasibility and further validate the simulation results, silicon experiments are performed, using one B200 GPU to emulate a 32-GPU B200 deployment. The experiment evaluates one selected two-stage pipeline configuration at global batch 2,048 using CUDA Green Contexts and TensorRT-LLM kernels. Results show that SubGPU Pipelining improves mean DRAM read-plus-write throughput from 71.8% to 83.6% of peak, and reduces the 36-layer decoder iteration time by 10.33%.

1. Kernel launch overhead. At small decode batch sizes, the dispatch latency for each kernel represents a nontrivial fraction of total step time. 2. Under-utilized compute. Memory-bound operators (e.g., expert projections at low-batch sizes) can saturate the HBM bandwidth roofline before consuming all SMs. Using all the SMs in such a case can mean the GPU is underutilized. 3. Bursty bandwidth demand. Alternating kernels with different bottlenecks create periods of under-utilization. For example even though the execution as a whole may be memory bound, compute-bound and latency-bound kernels during execution can cause memory bandwidth to be underutilized.

This paper introduces a concept called SubGPU Pipelining. SubGPU Pipelining addresses inefficiency with LLM decode by cutting the temporal operator sequence within a layer (or, more generally, a repeating compute graph) into contiguous graph segments and executing each segment 1 Introduction concurrently on disjoint SM pools. The central problem is The decode phase of LLM inference is the major source to jointly choose the graph cuts, segment-to-pool binding, of cost in latency-sensitive serving. A single decode and microbatch size so as to minimize target-workload 1

makespan under shared-resource contention. We study this problem with simulation and provide proof-of-feasibility on a Blackwell GPU using the GPT-OSS 120B [20] decode forward pass. The key ideas are:

Megatron-LM [15] combines pipeline, tensor, and data parallelism and uses interleaved pipeline schedules at large scale. CTA-Pipelining [12] moves the dependency granularity below whole kernels to cooperative thread arrays (CTAs), also called CUDA thread blocks, on tightly connected, shared-memory multi-GPU systems. It allows producer and consumer CTAs from dependent kernels to make progress concurrently on different GPUs and can compose with tensor parallelism. These approaches distribute pipeline work across GPUs and do not directly address intra-GPU resource under-utilization.

1. Cut a layer’s operator graph into K ordered contiguous segments and bind them to D stages, each backed by a controlled, disjoint SM pool. K = D or K = D + 1; the latter represents the case where first and last segments may share a stage. 2. Divide the per-GPU batch into m microbatches and stream them through the segments while preserving graph dependencies and serializing work bound to the same stage. 2.2 Intra-Device Parallelism 3. Jointly optimize graph cuts, stage count, segment-tostage binding, per-stage SM allocation, and microbatch Earlier work proposed hardware, compiler, and runtime size because every segment’s duration depends on the support for pipelines within a GPU. Kim et al. [11] use complete overlapping schedule. reference-count scoreboarding and a pipeline-aware CTA scheduler to overlap independent portions of consecutive SubGPU Pipelining is conceptually composable with existkernels. VersaPipe [27] includes persistent-thread execuing parallelism strategies: it operates within the per-GPU tion models, explicit stage-to-SM mapping, and an autowork allocation produced by an outer TP, DP, EP, CP, or PP tuned hybrid of coarse- and fine-grained GPU pipelines. configuration, provided that communication and sharedcuSync retains separate CUDA kernels but replaces kernelresource interactions are included in the schedule. The wide barriers with tile-grained producer–consumer synlogical execution model does not prescribe a particular chronization, allowing independent tiles of dependent CUDA mechanism. kernels to overlap [9]. Rammer lowers operators into fine-grained tasks and constructs static spatio-temporal Contributions schedules that expose both inter- and intra-operator paral• A formal definition of SubGPU as a coupled graph seg- lelism [13]. Several recent LLM systems concurrently use a mentation, resource allocation, and microbatching probGPU’s compute, memory, and communication resources. lem with state-dependent segment times (§3). NanoFlow [28] splits serving batches into nano-batches • A search and performance-modeling methodology for and jointly searches their sizes, execution order, and fracevaluating candidate SubGPU configurations (§4). tional GPU resource allocations while accounting for con• Analytical simulation of GPT-OSS 120B on modeled current interference. Flux [3] fuses collective communiB200 deployments that estimates the opportunity from cation with computation, Comet [25] applies fine-grained spatial co-scheduling under explicitly stated modeling fused overlap to MoE workloads, and TokenWeave [7] assumptions (§5.2). uses wave-aware token splitting to overlap computation • A Blackwell prototype in TensorRT-LLM, requiring zero changes to existing CUDA kernels, that demonstrates and communication. DynaFlow [21] decouples the logical coarse spatial co-scheduling, raises measured average model from its physical schedule through a programmable frontend and an asynchronous runtime. DRAM throughput, and reduces iteration time (§5.3). Megakernels A complementary approach replaces a sequence of operator kernels with one persistent megakernel. ThunderKittens provides CUDA-embedded ab2.1 Pipeline Parallelism Across GPUs stractions for constructing high-performance kernels [23]; Pipeline Parallelism (PP) is a widely used technique for a ThunderKittens-based Llama implementation fuses a scaling deep learning training and inference across mul- complete forward pass, assigns precomputed instruction tiple GPUs. GPipe [8] introduced microbatch pipeline streams to SMs, and tracks dependencies within the kerparallelism, mapping disjoint sets of model layers onto nel [24]. MPK compiles tensor programs into SM-level different devices and flowing microbatches through the task graphs and executes them with decentralized schedulresulting pipeline. PipeDream [14] uses asynchronous ing inside a persistent megakernel, enabling cross-operator pipeline schedules to reduce fill and drain bubbles, while software pipelining and fine-grained overlap [4].

2

Background and Related Work

2

SubGPU focuses on a latency-oriented formulation over contiguous graph cuts, segment-to-stage binding, disjoint SM allocations, and microbatch size. Its logical model is independent of a particular dispatch mechanism; Section 6 discusses persistent kernels and other CUDA mechanisms as possible realizations and extensions.

2.3

erator DAG for a repeating region, such as one transformer layer. Here, E contains only dependencies internal to the region; cross-region dependencies are handled when regions are composed into the full forward-pass DAG. Each contiguous region of operators between two graph cuts is a segment, which will be provisioned a disjoint SM pool on which to execute. From a scheduling perspective, the resource-mapped segment forms a stage in the SubGPU pipeline. Let CK (G) contain the legal ways to cut G into K nonempty ordered, dependency-contiguous segments c = (S1 , . . . , SK ). Each prefix S1 ∪ · · · ∪ Sk must be predecessor-closed, so every boundary is a dependencyrespecting graph cut. A configuration uses D physical stages and either K = D or K = D + 1 segments. The latter case captures the periodic boundary: the last segment of one region and the first segment of the next may use the same stage. Numbering stages by first appearance, let w ∈ {0, 1} indicate this boundary reuse, so K = D + w/. For segment indices k ∈ {1, . . . , K}, define the segment-to-stage binding ( 1, w = 1 and k = K, ρw (k) = (1) k, otherwise.

Spatial Dataflow Execution

Kitsune augments a GPU with architectural primitives and compiler support for co-resident spatial dataflow pipelines [5]. Purpose-built dataflow accelerators pursue a related spatial execution model at the architecture level. SambaNova’s RDU maps computation, local storage, and data movement onto a reconfigurable fabric of compute units, memory units, and programmable interconnect [22]. Recent work on the SN40 applies this substrate to LLM decoding by fusing large decoder subgraphs, streaming intermediate values on chip, and overlapping samples across layers through BatchStreaming [6].

2.4

GPU Resource Partitioning

SM-level spatial sharing predates recent CUDA mechanisms. Adriaens et al. [1] partition SMs among concurrent GPU applications; Janzén et al. [10] enforce disjoint SM subsets in software and use partition-specific models to schedule a precedence-constrained task graph; and Zhao et al. [26] search SM partitions while accounting for bandwidth and cache interference. CUDA Streams allow concurrent kernel submission but provide no guarantee of co-resident execution or an explicit SM partition. Cooperative Groups and Multi-Process Service provide other forms of coordination and co-scheduling. CUDA Green Contexts [16] provide first-class SM partitioning within a single CUDA context: streams within a Green Context are bound to an SM subset, and kernels launched on those streams execute on that subset. Persistent kernels, CUDA events, and CUDA Graphs [17] provide different choices for dispatch and dependency signaling. These are possible realization mechanisms; the SubGPU execution model itself requires controlled resource partitioning and dependency-preserving progress, while useful performance also requires sufficiently low handoff overhead.

Thus the binding is one-to-one when w = 0; when w = 1, only the first and last segments share a stage. We represent a configuration as θ = (D, w, c, p, b),

(2)

where p = (p1 , . . . , pD ) gives the SM allocation of the D physical stages and b is the microbatch size. Let H denote the target GPU hardware, including its SM count, NSM and minimum per-stage SM allocation, pmin . The D stage allocations correspond to pairwise-disjoint SM subsets, so no SM is assigned to more than one physical stage. (D) Let PH contain the legal ordered P D-stage SM-allocation vectors for H, with pd ≥ pmin and d pd = NSM , and let B(B) contain the supported microbatch sizes that divide the per-GPU batch B. Feasibility is summarized by w ∈ {0, 1},

2 ≤ D ≤ |V | − w,

c ∈ CD+w (G), (D) p ∈ PH ,

(3)

b ∈ B(B).

Let F denote the configurations satisfying Equation 3. The choice of b gives m = B/b ∈ Z+ microbatches. Smaller 3 SubGPU Execution and Optimiza- microbatches expose more independent work and can fill more stages, but they may also reduce the efficiency of tion Problem each kernel invocation; this is why b is optimized jointly with the graph cuts and SM allocation. 3.1 Execution Model and Design Space Figure 2 gives scheduling intuition for the chain-like, Figure 1 contrasts conventional temporal execution with an one-to-one case. Each segment is gated by both its data illustrative SubGPU schedule. Let G = (V, E) be the op- dependencies and the preceding segment assigned to the 3

Figure 1: Temporal execution (left) versus an illustrative one-to-one SubGPU binding (right). Temporal execution executes one operator at a time, with the full GPU. SubGPU binds contiguous graph segments to disjoint SM pools and streams microbatches through them, allowing work with complementary resource demands to overlap. In this example, the temporal case is executing one batch across multiple waves. The spatial case is executing the same batch across 8 microbatches. same physical stage; the next subsection formalizes these constraints. Figure 3 shows the one-to-one case: five contiguous segments bind to five physical stages with different SM allocations.

and inter-segment handoff latency are negligible, so data and stage edges impose precedence but no separate delay. For any realized effective durations, let Ai and Ci denote start and completion times. Topological evaluation then gives (

3.2

Augmented Task DAG and Objective

Ai = max

For a configuration θ, unrolling one forward pass over its region instances and m microbatches yields segment tasks Jθ . A task i = (ℓ, µ, k) executes segment Sk in region instance ℓ for microbatch µ, on stage ρw (k). A region instance is one occurrence of the repeating graph G in the unrolled forward pass; for a transformer model, it typically corresponds to one layer. For task i, Pidata contains the tasks whose outputs it consumes, including cross-region dependencies. A fixed global dispatch order is chosen as a topological ordering of these data dependencies. The set Pistage is empty or contains the preceding task on the same physical stage after that global order is restricted to the stage. These stage edges serialize all work on a partition, including the boundary segments that share stage 1 when w = 1. Both edge types therefore point forward in the same order and together define an acyclic augmented task DAG. Intuitively, a task starts only after all required outputs have arrived and its stage is free. It then occupies that stage for an effective duration τi , supplied by Section 4; the forward pass ends when the last task finishes. In the formal model, we assume dependency signaling

) max Cj , max Cj stage

j∈Pidata

j∈Pi

,

(4)

Ci = Ai + τi , We take an empty maximum to be zero. The forward-pass makespan is T (θ) = max Ci . (5) i∈Jθ

For fixed τi , Equations 4 and 5 are a max-plus timing recurrence on a weighted event DAG and reduce to a longest-path calculation [2]. Section 4 separately describes how the simulator estimates τi under overlap and shared-resource contention. The optimization objective is θ⋆ ∈ arg min T (θ). θ∈F

(6)

The timing calculation is simple, but the optimization remains coupled: graph cuts change the tasks and data edges, boundary reuse changes the stage edges, and SM allocation and microbatching change the durations and critical path. The reported results use static configurations; Section 6 discusses dynamic reconfiguration. 4

practical as the outer grid grows. More importantly, changing a cut, allocation, or microbatch changes overlap, shared-resource contention, cache behavior, and comOutput/Residual Router/Dispatch Expert/Combine Norm + QKV Attention pletion times, so candidates cannot be scored by adding (a) Temporal baseline: segments execute serially on the full GPU. isolated kernel costs. Section 4 therefore uses inexpensive proposals and refinement to explore a small part of the space, and the event-driven simulator to evaluate the Norm + QKV coupled schedules that remain. Attention

Output/Residual

4

Router/Dispatch

Search and Performance Model

We now describe the event-driven model used to evaluate candidates and the staged procedure used to search the overlap drain fill space in Section 3.3. Lower-level implementation con(b) SubGPU: segments on distinct stages overlap subject to data stants and policies are outside the scope of this design and same-stage dependencies. report. Expert/Combine

Figure 2: Illustrative schedules for six microbatches under a chain-like, five-segment, one-to-one binding. Rectangle widths use hand-chosen relative durations to expose unequal task lengths and scheduling bubbles; they are not simulator outputs and do not imply that a segment has the same duration across microbatches or overlap states. Actual times follow Equation 4 using contention-dependent per-task durations supplied by the simulator. The panels are independently scaled and do not compare absolute makespan.

3.3

4.1

Event-Driven Simulator

The augmented task DAG fixes data precedence and the FIFO order on each physical stage. Internally, the simulator expands every segment task into its constituent kernel tasks. A stage executes only its head task and cannot bypass blocked work. Whenever a dependency releases a task or an active work counter reaches zero, the simulator recomputes completion times for the active set and advances to the next event. This captures fill, steady state, drain, recurrence across regions, and head-of-line blocking without assuming constant or balanced stage times. The resulting start and completion times determine the effective duration τi = Ci − Ai used by Section 3. Each kernel task carries analytical compute and memory work together with a modeled full-GPU compute rate. The model assumes that compute throughput scales linearly with the number of SMs assigned to the task’s stage: pd(u) rucomp = µu , (8) NSM

Scale of the Search Space

An outer parallelism configuration q specifies the GPU count and the applicable TP, DP, EP, CP, and PP degrees used to distribute the model. It induces the per-GPU operator graph Gq , local tensor shapes, and collective communication tasks. Let X be the finite set of feasible search points x = (q, B, b, D, π) for a given model, hardware, sequence length, and repeating-region definition, where B and b are the per-GPU batch and microbatch, D is the physical-stage count, and π is the fixed simulator-policy setting. Let Px contain the legal SM allocations at x. An exhaustive sweep would contain X Nraw = |Px | (|CD (Gq )| + |CD+1 (Gq )|) , (7)

where u is a kernel task, d(u) is its physical stage, and µu is its modeled full-GPU rate. Exposed memory concurrency grows with the assigned SM count but is capped by a calibrated per-SM estimate, a hardware limit, and the task’s remaining demand. L2-hit and HBM-bound traffic jointly consume L2 bandwidth; aggregate miss-side concurrency selects a nonlinear HBM bandwidth response; and each task is limited by the issue bandwidth of its assigned SMs. Concurrent tasks also share network bandwidth. Compute and data movement may progress concurrently, and a task completes only after all of its modeled work is finished. The same fixed sharing and cache policies are used for Temporal and SubGPU candidates within an evaluation sweep. The cache model tracks capacity-bounded residency of logical weight and activation objects across stages, layers,

x∈X

where the two cut terms count the one-to-one and boundary-reuse cases; an invalid cut set is empty. At one illustrative grid point, a 20-operator chain with five stages has 15,504 such cut-and-binding choices. Under the reported full-budget, two-SM-quantum policy, dividing 148 SMs among five nonempty ordered stages gives 1,088,430 allocations—about 1.7 × 1010 combinations before counting any other point in X . Brute-force evaluation therefore quickly becomes im5

Total = 148 SMs on one GPU (B200, NSM =148) 22 SMs

28 SMs

20 SMs

16 SMs

62 SMs

Norm + QKV

Attention

Output Proj / Residual

Router / Dispatch

Expert MLP / Combine

RMSNorm Q proj KV proj

RoPE Q/K FMHA

Output proj Add+Norm

Router gating TopK dispatch

Expert up (BMM) SwiGLU Expert down Combine

microbatch flow → Figure 3: Illustrative five-stage SubGPU configuration for one MoE transformer layer on a single B200 GPU. Five contiguous graph segments bind one-to-one to five SM pools; box width is proportional to SM allocation. This is one feasible configuration, not a prescribed segmentation. FMHA denotes fused multi-head attention, and BMM denotes batched matrix multiplication. and microbatches, charging misses to HBM. The resident Algorithm 1 Simulation-guided SubGPU search table begins empty for each simulation and evolves ac- Require: base DAG G, hardware H, outer grid X cording to deterministic admission and eviction policies. Ensure: reported nondominated configurations Grid point: x = (q, B, b, D, π) specifies outer parallelism configuActivation writes may seed later reads and evict resident ration, per-GPU batch, microbatch, physical-stage count, and fixed objects. KV-cache accesses are modeled as streaming trafsharing/cache policies. fic and are not admitted to the resident table. Weight reuse Derived objects: Gq is the instantiated local DAG; ϕx contains kernel resource features; Px is the set of legal SM allocations. is therefore an outcome of the simulated schedule rather Partition: c gives the contiguous cuts and w indicates boundary than an assumed SubGPU benefit. Collective operations reuse. are explicit tasks with setup latency and payload-dependent Objectives: for fixed (x, c, w), fbx (p) estimates makespan and fabric work, and compute kernels are configured with base fx (p) obtains it from the event simulator. Controls: g = 2 SMs, Lest ∈ N ∪ {∞}, and Lsim ∈ N; the launch overhead.

4.2

default uses Lest = ∞. 1: Q ← ∅ 2: for all x = (q, B, b, D, π) ∈ X do 3: Gq ← I NSTANTIATE(G, q) 4: if ¬F EASIBLE(x, Gq , H) then 5: continue 6: end if 7: ϕx ← K ERNEL F EATURES(Gq , x, H) 8: (c, w) ← P ROPOSE PARTITION(Gq , D, ϕx ) 9: if no legal partition is returned then 10: continue 11: end if 12: S ← I NITIAL A LLOCATIONS(c, w, ϕx , Px ) 13: if S = ∅ then 14: continue 15: end if 16: Cx ← ∅ 17: for all p0 ∈ S do 18: pb ← PAIRWISE L INE(fbx , Px , p0 , g, Lest ) 19: for all s ∈ {p0 , pb} do 20: ps ← PAIRWISE L INE(fx , Px , s, g, Lsim ) 21: Cx ← Cx ∪ {ps } 22: end for 23: end for 24: p⋆ ← arg minp∈Cx fx (p) 25: p⋆ ← PAIRWISE L INE(fx , Px , p⋆ , g, Lsim ) 26: Q ← Q ∪ {R ECORD(x, c, w, p⋆ , fx (p⋆ ))} 27: end for 28: return PARETO(Q)

Search Procedure

For each workload and deployment choice, the search instantiates the local operator graph and rejects infeasible configurations. It then proposes a graph segmentation, refines the SM allocation, and evaluates the resulting candidate. Prior work likewise uses performance models or profiling to account for co-scheduled interference while searching SM partitions and intra-device schedules [26,28]. The search is expensive enough that it is best to run offline, with pre-computed configurations selected by the inference engine. The partition proposal summarizes each kernel’s expected compute, memory, cache, and communication behavior using the simulator inputs described in algorithm 1. Dynamic programming tries to find an anchored, dependency-respecting contiguous partition that minimizes differences between kernels in each segment. This is to promote the sharing of different resources as segments overlap. A separate step evaluates legal circular rotations; an improving rotation may produce the boundary-reuse layout in Equation 1. The feature objective proposes plausible cuts rather than predicting execution time. Uniform and service-proportional SM allocations seed 6

Algorithm 2 Pairwise-line SM-allocation refinement

mine how quickly each kernel exposes memory-level parallelism as SMs are added. Collective schedules are selected from modeled latency/bandwidth costs before global fabric contention is applied. The outer sweep and feasible allocation sets are finite. Each allocation search terminates when no feasible pairwise move improves the objective or when an applicable pass limit is reached. Small graphs and SM budgets can be exhaustively enumerated to diagnose the segmentation and allocation heuristics. The reported points are therefore the best configurations found over the finite search grid under these modeling assumptions, not proofs of global optimality.

Require: objective f , feasible set P, initial allocation p ∈ P, SM quantum g, pass limit L ∈ N ∪ {∞} Ensure: feasible refined allocation p Move: pd→r [z] moves zg SMs from donor stage d to receiver stage r, leaving every other allocation unchanged. Local variables: M lists legal one-quantum transfers; t is the current objective; z counts transferred quanta; h is the line-search step; and tline is the best value found on that line. If L = ∞, refinement continues until an early return. 1: t ← f (p) 2: for pass = 1, . . . , L do 3: M ← {(d, r) : d ̸= r, pd→r [1] ∈ P} 4: if M = ∅ then 5: return p 6: end if ′ ′ 7: (d, r) ← arg min(d′ ,r′ )∈M f (pd →r [1]) 8: z ← 1; h ← 1; tline ← f (pd→r [1]) 9: if tline ≥ t then 10: return p 11: end if 12: while h > 0 do 13: z′ ← z + h 14: if pd→r [z ′ ] ∈ P ∧ f (pd→r [z ′ ]) < tline then 15: z ← z ′ ; tline ← f (pd→r [z ′ ]); h ← 2h 16: else 17: h ← ⌊h/2⌋ 18: end if 19: end while 20: p ← pd→r [z]; t ← tline 21: end for 22: return p

5

Evaluation

5.1

Evaluation Methodology

We evaluate SubGPU Pipelining in two complementary ways on the same workload: GPT-OSS 120B [20] decode at input sequence length 32,768 and KV history 32,768, with 32-GPU B200 deployment settings: TP8/DP4 for attention and EP32 for MoE. Section 5.2 uses the analytical shared-resource simulator to search the full design space and map the per-user/per-GPU throughput Pareto frontier; Section 5.3 implements one pipeline design on a B200 GPU with native TensorRT-LLM [19] kernels and the allocation search. For each seed, a cheap dependency- measures its latency and resource utilization. In both evaluations, we compare: aware makespan estimate first guides pairwise SM transfers. The event simulator then polishes both the original al• Temporal: conventional kernel-by-kernel execution, location and its estimator-refined counterpart, followed by with each kernel eligible to use the full SM pool. a final polish of the best candidate. Simulator evaluations • SubGPU: the forward-pass graph is cut into contiguous are memoized. By default, estimator-guided refinement segments, where each segment is bound to a stage with continues until no improving transfer remains, while each a static SM allocation, and microbatches are streamed simulator-guided polishing call has a bounded pass count. through the stages. The reported search assigns each physical stage a positive, even number of SMs and uses the full SM budget. The two evaluations serve different purposes. The simulaIt tests microbatch counts below, near, and above the tion explores the design space: each point on its frontier physical-stage count; Temporal uses one full-batch micro- may use a different sharding, batch size, microbatch size, batch. These are search choices rather than requirements graph segmentation, stage count, and SM allocation, so it of the SubGPU model. After the sweep, the reported con- shows where SubGPU helps and which pipeline structures figurations are Pareto-filtered. Algorithms 1 and 2 give the to build. The hardware experiment runs the pipeline decorresponding procedures. signs on real hardware at global batch 2,048 and reports the best-performing one in detail. It shows that the design can be realized with existing CUDA mechanisms and kernels, 4.3 Calibration and Scope and measures the latency and utilization changes bring by Kernel work estimates come from analytical operator SubGPU Pipelining. shapes and deployment sharding. Hardware inputs proBoth evaluations use pipeline designs that we develvide datatype-specific compute throughput, HBM and L2 oped for GPT-OSS, whose layers alternate between slidingcapacity and bandwidth, per-SM issue bandwidth, commu- window and full attention. Sliding-window attention is nication latency and bandwidth, and an empirical HBM substantially shorter than full attention, so placing both atbandwidth-versus-concurrency curve. A calibrated per- tention operators in one stage leaves that stage imbalanced kernel bytes-in-flight estimate and its global scale deter- across successive layers. Therefore, we explore designs 7

Figure 4: Modeled Pareto frontiers for GPT-OSS 120B decode on 32 B200 GPUs at input sequence length 32,768. The SubGPU curve pools the best retained configuration across stage counts.

Figure 5: The same Pareto candidates separated by SubGPU stage count. Each stage is one physical SM partition; the temporal baseline uses the full SM pool.

over a two-layer macro-period rather than a single layer. We consider several such designs, including ones with additional stages spanning two layers, and ones with a stage dedicated to sliding-window attention. The simplest is the “Macro” two-stage design: one stage runs only fullattention FMHA, and the other runs all remaining operators of the two-layer macro-period, including sliding-window attention. We evaluate these designs in both simulation and hardware.

GPU provides up to 2.2× higher per-user throughput than temporal execution. The gain is largest at the throughputoriented end of the frontier: it is 1.8–2.2× above 5,000 tokens/s/GPU and 1.3–1.8× between 3,000 and 5,000 tokens/s/GPU. At matched per-user throughput, SubGPU provides 1.1–1.4× higher per-GPU throughput between 100 and 600 tokens/s/user and up to 1.9× toward the latency-oriented end.

The two frontiers approach parity at their extremes. At very small batches, there are too few independent microbatches to expose substantial overlap. At the opposite ex5.2 Analytical Simulation treme, operators become more compute-bound and better 5.2.1 Simulation Setup utilize the temporal SM pool, reducing the capacity available for productive spatial co-scheduling. The modeled For GPT-OSS 120B decode on 32 B200 GPUs, the simgains therefore arise from the joint configuration selected ulator of Section 4 is used to estimate the makespan of for the intermediate regime rather than from a universally one complete forward pass for each temporal and SubGPU faster segmentation or binding. candidate, and we retain the best configurations found by the search. The simulator unrolls all layers into the forward-pass DAG, represents collectives as explicit tasks, and carries shared-resource and L2-residency state across layer boundaries. We report the Pareto frontier of per- 5.2.3 Effect of Stage Count GPU throughput versus per-user throughput, where peruser throughput serves as an inverse measure of latency. Figure 5 groups the SubGPU candidates by stage count. The curves examine the benefit of spatial execution un- More stages offer finer resource right-sizing and additional der our stated modeling assumptions and are not meant to concurrency, but also divide the SM pool more finely and predict measured execution. require enough microbatch parallelism to keep every stage productive. Since the first and last graph segments may share a stage, a D-stage configuration contains either D or 5.2.2 Simulated Pareto Opportunity D + 1 contiguous segments. Stage count is therefore couFigure 4 compares the best retained frontiers. We com- pled to graph cuts, boundary reuse, microbatch size, and pare them in two ways, interpolating linearly between per-stage allocation rather than being a fixed architectural frontier points. At matched per-GPU throughput, Sub- constant. 8

Figure 6: Kernel-level timeline for a representative 1.8-ms steady-state segment from the middle of the complete GPT-OSS replay at global batch 2048. The temporal baseline executes sliding-window attention, full attention, and the remaining operators on all 148 SMs. The SubGPU pipeline assigns full-attention FMHA to 68 SMs and everything else—including sliding-window attention and kernels from later layers—to 80 SMs. Each Green Context contains two CUDA stream rows. This is an instrumented representative replay; authoritative end-to-end timing comes from uninstrumented CUDA-event measurements.

5.3

Hardware Prototype on Blackwell

5.3.1

Experimental Setup

Context. As actual kernel performance differs from the SOL estimates used by the simulator, which shifts the relative durations of the stages, for each of the pipeline designs, we tune the per-stage SM allocation to balances the execution of pipeline stages and minimize the measured iteration latency, instead of directly reusing the SM allocations from the simulator.

The silicon prototype of SubGPU Pipelining for GPTOSS 120B is implemented with TensorRT-LLM [19] on NVIDIA B200 GPU. We used one B200 to emulate the per-rank workload of the 32-GPU deployment described in Section 5.1. For this experiment, the global decode batch is chosen to be 2048 tokens. All computation kernels leverage the TensorRT-LLM kernels with version 1.3.0rc22.post1, while communication are estimated using virtual kernels, performing the corresponding attention allreduce, MoE dispatch, and MoE combine endpoint work. In our virtual communication model, each communication is floored at 3 µs latency plus payload bytes divided by 900GB/s, representing the unidirectional B200 NVLink bandwidth.

Across all designs and tuned allocations, the “Macro” two-stage pipeline design with full-attention FMHA in a 68-SM Green Context and all other operators across two layers in an 80-SM Green Context achieves the lowest measured latency. We report this configuration in the rest of this section. Each Green Context owns two streams, and two equal microbatches flow through them. The temporal baseline executes the same layer sequence as one batch on all 148 SMs. CUDA events and CUDA Graphs enforce the 5.3.2 Pipeline Design and Implementation producer/consumer and same-stream dependencies. FigWe implement all purposed pipeline designs on hardware, ure 6 shows a representative steady-state segment from the realizing every stage and SM partitions using CUDA Green middle of the complete 36-layer replay. 9

(a) DRAM read-plus-write throughput. The SubGPU pipeline sustains 83.6% of peak on average, compared with 71.8% for the temporal baseline.

(b) Fraction of the 148 SMs active. The SubGPU pipeline averages 71.7% and the temporal baseline averages 85.9%

Operator category

Temporal baseline (ms)

SubGPU pipeline (ms)

Norm + QKV Full attention Sliding attention Router/Dispatch Expert MLP Output/Combine

0.556 ± 0.001 10.957 ± 0.008 0.190 ± 0.001 1.406 ± 0.000 0.981 ± 0.003 0.930 ± 0.005

1.358 ± 0.002 15.345 ± 0.020 0.426 ± 0.002 3.473 ± 0.010 3.059 ± 0.012 2.789 ± 0.017

Figure 7: Measured GPU utilization for one 36-layer GPT-OSS 120B decode iteration on an NVIDIA B200. Nsight Systems counters are sampled at 200 kHz and aggregated into 100 µs bins for the full-iteration panels and 40 µs bins for the zoomed panels. Means are calculated from the unbinned 200-kHz samples. Left panels show the full iteration. Right panels zoom in on repeated sequences of two pipeline cycles, where each cycle spans one sliding-attention layer and one full-attention layer. Dashed and dotted vertical lines indicate the mean durations of one and two pipeline cycles, respectively. The table gives the total kernel residence time for each category, reported as mean ± standard deviation across the four replays. 5.3.3

Latency and Resource Utilization Analysis

execution incurs lower kernel efficiency and additional endpoint overhead. These small costs accumulate along After 20 warmups, three rounds of 100 uninstrumented the serialized temporal baseline. In the SubGPU pipeline, CUDA Graph replays give a mean 36-layer makespan of costs off the critical path can overlap with work in the other 14.970 ms for the temporal baseline and 13.424 ms for the context and be partially hidden. Therefore, the temporal SubGPU pipeline, showing a 10.33% latency reduction. In baseline can degrade more than the pipeline relative to the corresponding 200-kHz Nsight captures, mean DRAM their modeled performance, producing a larger measured read-plus-write throughput over four steady replays in- speedup. This tolerance of small per-kernel inefficiencies creases from 71.8% to 83.6% of peak. Mean active-SM is a second benefit of SubGPU Pipelining, in addition to throughput decreases from 85.9% to 71.7%: partitioning improving shared-resource utilization. uses fewer SMs for each individual kernel, but overlap raises memory-system utilization and reduces end-to-end time. Figure 7 shows these measured profiles.

5.4 Comparing Simulation and Hardware Results The two evaluations agree in direction: at batch 2048, both find that SubGPU Pipelining reduces decode latency relative to temporal execution, despite differ in magnitude. At batch 2,048, the best simulated two-stage configuration reduces makespan by 6.6% (Section 5.2), whereas the hardware measures a 10.33% reduction. We attribute part of this difference to optimistic per-kernel assumptions in the simulator. For example, the model allows kernels to operate near 99% of peak DRAM bandwidth, whereas real

Taken together, the two results support each other. The simulation shows where in the design space SubGPU Pipelining helps and which pipeline structures are worth building. The hardware experiment shows that such a structure can be realized with Green Contexts and unmodified kernels, and that it delivers a measured 10.33% latency reduction with higher DRAM utilization on a real GPU. We view these results as a proof of feasibility, with additional optimization still possible through different cuts, allocations, and kernels.

10

6 Realization, Extensions, and Com- 6.4 Composition with Outer Parallelism position SubGPU operates within the per-GPU work allocation pro6.1

Implementation Considerations

The logical model requires controlled partitioning of execution capacity, asynchronous progress of work on different stages, serialization of segments bound to the same stage, and dependency signaling. The timing model assumes that signaling and handoff overhead are negligible. Several CUDA mechanisms can provide parts of this functionality. Green Contexts. CUDA Green Contexts [16] bind streams to SM subsets and provide the resource isolation used by the current prototype.

duced by outer tensor, data, expert, context, or pipeline parallelism, although the ideal parallelism strategy may change when SubGPU is used. The outer strategy determines each GPU’s operator graph, tensor shapes, communication, and batch; the SubGPU optimizer then cuts that graph and binds its segments to local stages. When pipeline parallelism is present, SubGPU forms an inner, intra-GPU schedule within each pipeline rank. Outer communication must appear as tasks or data edges in the augmented task DAG, and its shared-resource effects must be reflected in the Section 4 duration model.

7

Discussion and Limitations

CUDA events and graphs. The prototype uses events for dependency ordering and CUDA Graphs for repeatable Applicability conditions. SubGPU relies on enough inlow-overhead submission while retaining existing kernels. dependent work to overlap work across stages. It is most promising when the batch exposes multiple microbatches, operators have complementary resource profiles, and some Persistent kernels. A more specialized implementation temporal operators reach their bandwidth or parallelism could keep per-stage workers resident and feed them micro- limit before using the full SM pool. The opportunity diminbatches through work queues, potentially reducing launch ishes when the batch is too small to populate the pipeline or overhead and enabling additional on-chip state reuse. This when large, compute-bound kernels already use the GPU is an implementation option rather than a defining prop- efficiently. erty of SubGPU and requires separate engineering and Microbatching also changes weight-reuse behavior. If evaluation. concurrently active segment working sets do not remain in L2, repeated HBM traffic can offset the benefit of spatial overlap. This limitation becomes more important for 6.2 Dynamic Reconfiguration larger models with larger weights and activations. TechSection 3 notes that graph cuts, stage count, boundary niques such as prefetching data into L2 may mitigate this reuse, per-stage allocation, and microbatch size may vary overhead, but require further evaluation. between runtime regimes. A runtime could select different configurations as KV-cache length, layer type, routed expert load, or available batch parallelism changes. The Modeling limitations. The simulation results are modappropriate transition points and the cost of reconstruct- eled opportunities rather than formal upper bounds. They ing or relaunching the underlying execution resources are use static within-configuration allocations and configured mechanism-specific. The simulator results and prototype per-kernel base overheads rather than modeling launch and synchronization dynamics in detail. L2 residency is use static choices within each evaluated configuration. enabled for modeled weights and activations, while KVcache accesses remain streaming traffic that is not admitted to the resident-state table. Their credibility depends on the 6.3 Weight Prefetch resource-sharing, cache-initialization, admission, eviction, Prefetching the weights needed by segment k for layer communication, and operator-scaling policies described in L + 1 while other work from layer L is active may reduce Section 4. Sensitivity to these assumptions and calibration the first-use latency of the next layer. It is not free: the against measured kernels are necessary before extrapolatprefetch consumes HBM and L2 bandwidth, cache capac- ing the Pareto gains to an end-to-end deployment. ity, and possibly execution resources, and can therefore change the durations and overlap of every concurrently active segment. Temporal execution may also expose Implementation complexity. Any implementation must prefetch opportunities. A thorough evaluation of prefetch- manage graph-cut and SM-partition granularity, depening is beyond the scope of this report, and is included in dency signaling, buffering, failures, and transitions beour on-going work on SubGPU Pipelining. tween configurations. Persistent-kernel implementations 11

add work-queue, SMEM-lifetime, and recovery concerns, while framework-level implementations must control launch and synchronization overhead. These costs are mechanism-specific and are not part of the logical definition.

Future work. The next steps are to validate the simulator and optimizer more extensively, compare selected configurations with exhaustive search on small instances, and evaluate end-to-end dense and MoE deployments across GPU generations. Dynamic allocation, persistent-kernel realizations, weight prefetch, and longer-lived on-chip weight persistency remain extensions to be evaluated with their contention and transition costs included.

8

Declaration of Generative AI and AI-assisted technologies in the writing process During the preparation of this work, the authors used generative AI to improve the clarity and fluency of the English text as well as some scripts. After using these tools, the authors reviewed and edited the content as needed and take full responsibility for the publication’s content.

References [1] Jacob T. Adriaens, Katherine Compton, Nam Sung Kim, and Michael J. Schulte. The case for GPGPU spatial multitasking. In Proceedings of the 18th IEEE International Symposium on High Performance Computer Architecture (HPCA), pages 79–90. IEEE, 2012. [2] François Baccelli, Guy Cohen, Geert Jan Olsder, and JeanPierre Quadrat. Synchronization and Linearity: An Algebra for Discrete Event Systems. John Wiley & Sons, Chichester, UK, 1992.

Conclusion

We presented SubGPU Pipelining, a technique for improving GPU compute and memory utilization. SubGPU Pipelining partitions the forward-pass graph into contiguous segments, binds each segment to a dedicated pool of SMs, and streams microbatches through these pools so that operators with different resource demands run concurrently. The approach is defined by this schedule rather than by any particular CUDA primitive. We evaluated SubGPU Pipelining on GPT-OSS 120B in a 32-GPU B200 deployment setting. In simulation, SubGPU configurations outperform temporal execution across most of the per-user/per-GPU throughput Pareto frontier, providing up to 2.2× higher per-user throughput at matched per-GPU throughput and up to 1.9× higher per-GPU throughput at matched per-user throughput. On real hardware, a twostage implementation built from CUDA Green Contexts and existing TensorRT-LLM kernels raises mean DRAM throughput from 71.8% to 83.6% of peak and reduces the iteration time by 10.33%. This result shows that SubGPU Pipelining is physically feasible on current GPUs and delivers a measurable benefit without custom kernels. Validating additional configurations and exploring more specialized implementations remain future work.

Acknowledgments We thank various teams within NVIDIA for their support. This includes Jeremy Wilke, Wonchan Lee, Anurag Gupta, Haoming Lu, Yuanqing Zhao, Jatin Mitra, Ankit More, Brian Pharris, Angshuman Parashar, Vikram Sharma Mailthody, Neal Crago, Guilleme Sole, Andrew Foote. 12

[3] Li-Wen Chang, Wenlei Bao, Qi Hou, Chengquan Jiang, Ningxin Zheng, Yinmin Zhong, Xuanrun Zhang, Zuquan Song, Chengji Yao, Ziheng Jiang, Haibin Lin, Xin Jin, and Xin Liu. Flux: Fast software-based communication overlap on GPUs through kernel fusion. arXiv preprint arXiv:2406.06858, 2024. [4] Xinhao Cheng, Zhihao Zhang, Yu Zhou, Jianan Ji, Jinchen Jiang, Zepeng Zhao, Ziruo Xiao, Zihao Ye, Yingyi Huang, Ruihang Lai, Hongyi Jin, Bohan Hou, Mengdi Wu, Yixin Dong, Anthony Yip, Zihao Ye, Songting Wang, Wenqin Yang, Xupeng Miao, Tianqi Chen, and Zhihao Jia. MPK: A compiler and runtime for Mega-Kernelizing tensor programs. In 20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 26), pages 1909–1926, Seattle, WA, July 2026. USENIX Association. [5] Michael Davies, Neal Crago, Karthikeyan Sankaralingam, and Stephen W. Keckler. Kitsune: Enabling dataflow execution on GPUs with spatial pipelines. ACM Transactions on Architecture and Code Optimization, 22(4), December 2025. [6] Darshan Gandhi, Pushkar Nandkar, David Koeplinger, Nasim Farahini, Romy Tsoupidi, Samuel Rydh, Matheen Musaddiq, Tuowen Zhao, Reid Goodbar, Nathan Sheeley, Leon Zhang, Matthew Shaffer, John Long, Han Wang, Angela Wang, Arjun Sabnis, Joshua Brot, Yun Du, Hakan Zeffer, Mingran Wang, and Raghu Prabhakar. Dataflow is all you need. In A. Chowdhery and Z. Jia, editors, Proceedings of Machine Learning and Systems, volume 8, pages 422–436. MLSys, 2026. [7] Raja Gond, Nipun Kwatra, and Ramachandran Ramjee. TokenWeave: Efficient compute-communication overlap for distributed LLM inference. In Proceedings of Machine Learning and Systems, volume 8, 2026. [8] Yanping Huang, Youlong Cheng, Ankur Bapna, Orhan Firat, Mia Xu Chen, Dehao Chen, HyoukJoong Lee, Jiquan Ngiam, Quoc V. Le, Yonghui Wu, and Zhifeng Chen. GPipe:

Efficient training of giant neural networks using pipeline parallelism. In Advances in Neural Information Processing Systems (NeurIPS), 2019.

[19] NVIDIA Corporation. NVIDIA TensorRT-LLM. https://docs.nvidia.com/tensorrt-llm/ index.html, 2026.

[9] Abhinav Jangda, Saeed Maleki, Maryam Mehri Dehnavi, [20] OpenAI. gpt-oss-120b & gpt-oss-20b model card. arXiv preprint arXiv:2508.10925, 2025. https://arxiv. Madan Musuvathi, and Olli Saarikivi. A framework for fineorg/abs/2508.10925. grained synchronization of dependent GPU kernels. In 2024 IEEE/ACM International Symposium on Code Generation [21] Yi Pan, Yile Gu, Jinbin Luo, Yibo Wu, Ziren Wang, Hongand Optimization (CGO), pages 93–105, 2024. tao Zhang, Ziyi Xu, Shengkai Lin, Baris Kasikci, and Stephanie Wang. DynaFlow: Transparent and flexible [10] Johan Janzén, David Black-Schaffer, and Andra Hugo. Parintra-device parallelism via programmable operator schedultitioning GPUs for improved scalability. In 2016 IEEE 28th ing. In Proceedings of Machine Learning and Systems, International Symposium on Computer Architecture and volume 8, 2026. High Performance Computing (SBAC-PAD), pages 42–49. IEEE, 2016.

[22] Raghu Prabhakar, Sumti Jairath, and Jinuk Luke Shin. SambaNova SN10 RDU: A 7nm dataflow architecture to accel[11] Gwangsun Kim, Jiyun Jeong, John Kim, and Mark Stephenerate software 2.0. In 2022 IEEE International Solid-State son. Automatically exploiting implicit pipeline parallelism Circuits Conference (ISSCC), pages 350–352. IEEE, Februfrom multiple dependent kernels for GPUs. In Proceedings ary 2022. of the 25th International Conference on Parallel Architectures and Compilation Techniques, pages 341–352, 2016. [23] Benjamin Spector, Simran Arora, Aaryan Singhal, Arjun Parthasarathy, Dan Fu, and Christopher Ré. ThunderKit[12] Tingkai Liu, Muralidhar Andoorveedu, Sanjoy Das, Santens: Simple, fast, and Adorable kernels. In International jay Patel, and Volodymyr Kindratenko. CTA-Pipelining: Conference on Learning Representations (ICLR), 2025. A latency-oriented spatial scaling method for multi-GPU [24] Benjamin Spector, Jordan Juravsky, Stuart Sul, Owen systems. arXiv preprint arXiv:2607.07862, 2026. Dugan, Dylan Lim, Dan Fu, Simran Arora, and Christopher [13] Lingxiao Ma, Zhiqiang Xie, Zhi Yang, Jilong Xue, Youshan Ré. Look ma, no bubbles! designing a low-latency Miao, Wei Cui, Wenxiang Hu, Fan Yang, Lintao Zhang, and megakernel for Llama-1B. Hazy Research technical Lidong Zhou. Rammer: Enabling holistic deep learning blog, https://hazyresearch.stanford.edu/ compiler optimizations with rTasks. In 14th USENIX Symblog/2025-05-27-no-bubbles, May 2025. posium on Operating Systems Design and Implementation Code: https://github.com/HazyResearch/ (OSDI 20), pages 881–897. USENIX Association, NovemMegakernels. ber 2020. [25] Shulai Zhang, Ningxin Zheng, Haibin Lin, Ziheng Jiang, [14] Deepak Narayanan, Aaron Harlap, Amar Phanishayee, Wenlei Bao, Chengquan Jiang, Qi Hou, Weihao Cui, Size Vivek Seshadri, Nikhil R. Devanur, Gregory R. Ganger, Zheng, Li-Wen Chang, Quan Chen, and Xin Liu. Comet: Phillip B. Gibbons, and Matei Zaharia. PipeDream: GenerFine-grained computation-communication overlapping for alized pipeline parallelism for DNN training. In Proceedmixture-of-experts. In Proceedings of Machine Learning ings of the 27th ACM Symposium on Operating Systems and Systems, volume 7, 2025. Principles (SOSP), 2019. [26] Xia Zhao, Zhiying Wang, and Lieven Eeckhout. Classification-driven search for effective SM partitioning [15] Deepak Narayanan, Mohammad Shoeybi, Jared Casper, in multitasking GPUs. In Proceedings of the 32nd InternaPatrick LeGresley, Mostofa Patwary, Vijay Korthikanti, tional Conference on Supercomputing (ICS), pages 65–75. Dmitri Vainbrand, Prethvi Kashinkunti, Julie Bernauer, ACM, 2018. Bryan Catanzaro, Amar Phanishayee, and Matei Zaharia. Efficient large-scale language model training on GPU clus- [27] Zhen Zheng, Chanyoung Oh, Jidong Zhai, Xipeng Shen, ters using Megatron-LM. In Proceedings of the InternaYoungmin Yi, and Wenguang Chen. VersaPipe: A versatile tional Conference for High Performance Computing, Netprogramming framework for pipelined computing on GPU. working, Storage and Analysis (SC’21), 2021. In Proceedings of the 50th Annual IEEE/ACM International Symposium on Microarchitecture, pages 587–599, 2017. CUDA toolkit 12.4 re[16] NVIDIA Corporation. lease notes: Green Contexts. https:// [28] Kan Zhu, Yufei Gao, Yilong Zhao, Liangyu Zhao, Gefei docs.nvidia.com/cuda/archive/12.4.0/ Zuo, Yile Gu, Dedong Xie, Zihao Ye, Keisuke Kamahori, cuda-toolkit-release-notes/index.html, Chien-Yu Lin, Ziren Wang, Stephanie Wang, Arvind Krish2024. namurthy, and Baris Kasikci. NanoFlow: Towards optimal large language model serving throughput. In 19th USENIX [17] NVIDIA Corporation. CUDA Programming Symposium on Operating Systems Design and ImplemenGuide. https://docs.nvidia.com/cuda/ tation (OSDI 25), pages 749–765, Boston, MA, July 2025. cuda-programming-guide/, 2026. USENIX Association. [18] NVIDIA Corporation. NVIDIA Blackwell Tuning Guide. https://docs.nvidia.com/cuda/ blackwell-tuning-guide/, 2026.

13

Related documents

• SubGPU Pipelining: Intra-GPU Fine-Grained Spatial Parallelism for LLM Inference#1648934
• Training an MoE LLM in MXFP4: How We Closed the Gap to BF16#1447596
• PipeLive: Efficient Live In-place Pipeline Parallelism Reconfiguration for Dynamic LLM Serving#1063894
• Simulate Before You Scale: High-Fidelity LLM Inference Planning Without GPUs#1324856
• Elastic Precision: Software-Hardware Co-Design of a Mixed-Precision Architecture for Long-Context LLM Inference#1536730