arXiv:2605.03190v1 [cs.DC] 4 May 2026
VDCores: Resource Decoupled Programming and Execution for Asynchronous GPUs Zijian He
Adrian Sampson
[email protected] University of California, San Diego San Diego, California, USA
[email protected] Cornell University Ithaca, New York, USA
Yiying Zhang
Zhiyuan Guo
[email protected] University of California, San Diego and GenseeAI Inc. San Diego, California, USA
[email protected] Cornell University Ithaca, New York, USA
Abstract
However, GPU programming and execution models have not evolved accordingly. Most existing systems still rely on a resource-coupled monolithic kernel model, originally designed for largely synchronous, data-parallel execution. They continue to use this model even as GPU hardware becomes increasingly asynchronous and task-parallel. This model fundamentally misaligns with modern GPU hardware and creates two drawbacks. First, the monolithic kernel abstraction substantially increases the complexity of asynchronous GPU programming. GPU programmers must carefully orchestrate memory movement, tensor-core execution, synchronization, and pipelining to avoid bubbles [33, 41, 49]. Second, monolithic kernels pack dependency management and overlap decisions into a single opaque schedule and execution unit [7, 8]. This creates utilization bubbles and prevents opportunistic execution across kernel boundaries based on runtime resource availability. Dynamic workloads further exacerbate this limitation, where changing inputs quickly invalidates compile-time overlap and resource-allocation decisions [22, 40]. Recent frameworks embrace new programming paradigms, libraries, or compilers to ease the programming and performance tuning for asynchronous GPUs. For example, CUTLASS uses warp specialization to overlap memory and compute within a kernel, while Mirage Persistent Kernel (MPK) [16] and ThunderKittens (TK) [41] extend this idea with cross-task pipelining. However, these approaches only mitigate, rather than remove, the mismatch between monolithic kernels and asynchronous hardware. They improve overlap inside kernels or across statically planned tasks, but still preserve the kernel/task as the main unit of composition. To better match asynchronous GPUs, we propose a direct but fundamental shift in GPU programming and execution model: Expose each asynchronous hardware unit on GPUs as an isolated programming and execution unit. Developers therefore program asynchronous units directly and independently, rather than wiring them inside a single monolithic abstraction. At runtime, each hardware unit is driven
Modern GPUs increasingly rely on specialized and asynchronous hardware units to deliver high performance. Yet these units are often underutilized because today’s GPU software stacks still organize programming and execution around a monolithic kernel model that mismatches asynchronous hardware. To address this issue, Virtual Decoupled Engines (VDCores) presents a new decoupled programming and execution model for asynchronous GPUs. VDCores abstracts asynchronous hardware execution units as resource isolated virtual cores and represents workloads as dependency-connected micro-operations (𝜇ops). this abstraction removes static orchestration from the programmer, enables automatic overlap of memory and compute based on dependency and resource readiness, and thereby improves utilization of asynchronous hardware resources. Realizing such a decoupled abstraction efficiently on today’s GPUs is itself challenging, VDCores addresses this through a GPU-specialized programming model and GPU runtime design that preserves the flexibility while minimizing implementation overhead. Across four LLM inference workloads on GH200, H100, and RTX 6000 Pro GPUs, VDCores significantly improves decoding throughput by 24% on average and by up to 77% under dynamic inputs, while reducing kernel programming and specialization effort by 90%. We have open sourced VDCores at https://github.com/vdcores/ vdcores.
1
Introduction
Modern GPUs increasingly expose specialized and asynchronous internal resources, such as tensor cores and hardwareassisted asynchronous memory movement [30, 32]. Architectures such as NVIDIA Hopper further strengthen this trend with mechanisms like the Tensor Memory Accelerator (TMA), allowing memory transfer and computation to proceed concurrently [27]. In principle, such hardware should improve utilization by reducing idle time and enabling overlap across different execution units. 1
Zijian He, Adrian Sampson, Yiying Zhang, and Zhiyuan Guo Monolith Kernel Model H100 GPU kernel kernels
GPU Compute Unit
VDCores Decoupled Model (§ 4.1)
VDCores Executor
uops ops
Virtual Virtual Comptue ComputeCore Core
Generator (§ 4.2)
uops
(§ 4.3)
Virtual Memory Core
SM
SM
x 132
Abstraction
Streaming Multiprocessor (SM) CUDA CUDA Core Core
Async Compute Units
Tensor Core
Tensor Memory Accelerator (TMA)
Async Memory Units
Figure 1. Comparing VDCores and monolith kernel programming and execution models. On an NVIDIA H100 GPU with asynchronous hardware units.
by resource-isolated and opportunistic scheduling rather than static orchestration. Based on this idea, we build virtual decoupled cores (VDCores), the first system providing a decoupled programming interface and execution stack for asynchronous GPUs. VDCores realizes the decoupled model by virtualizing memory and compute resources into software-managed execution units: virtual memory cores (VMCs) and virtual compute cores (VCCs), each of which could execute independently. Given operators from high-level ML frameworks, e.g. PyTorch, VDCores lowers workloads into dependencyconnected 𝜇op streams, and maps them onto corresponding virtual units for execution. VDCores provides a unified solution to asynchronous hardware. First, it lets programmers express computation and data movement as fine-grained and resource-isolated 𝜇ops, without manually and statically specifying overlap, orchestration, or synchronization strategies. Second, at runtime, the system can automatically overlap memory and compute operations based on actual dependency satisfaction and resource availability. Finally, once a library of fine-grained 𝜇ops is built for this virtual architecture, future tasks can reuse and re-orchestrate them to realize different execution plans, without rebuilding specialized kernels for every workload variant. Despite its promise, major challenge lies in efficiently realizing the decoupled model on existing GPUs. Modern GPUs still use a SIMT programming interface that favors regular, uniform control flow to expose asynchronous units. Materializing a decoupled runtime on top of this substrate can resemble building a software microarchitectural emulator: it must explicitly track resource availability, manage finegrained dependencies, and schedule work across asynchronous units in real time. Such logic is inherently branch-heavy, stateful, and synchronization-intensive, making it a poor fit for SIMT execution. Without careful design, these control and bookkeeping overheads can easily erase the flexibility
VDCores
Figure 2. VDCores Inference Performance and Programming Effort. Bottom left is better.
and utilization gains of decoupling. VDCores addresses this performance challenge with the following key ideas: First, VDCores fully exploits the different forms of parallelism available on GPUs to sustain high 𝜇op throughput. Rather than interpreting 𝜇ops through a purely scalar control path, VDCores organizes each virtual core as a software pipeline with separate control and execution stages, and uses SIMT threads cooperatively within each path for instruction handling, address generation, allocation, and data movement. This design lets VDCores retain the flexibility of decoupled execution while still issuing and executing 𝜇ops fast enough to fullfil the capacity of GPU computation and memory. Second, VDCores reduces the cost of dependency tracking through co-design of the programming model, 𝜇op generation, and execution. The key insight is to restrict and structure inter-𝜇op dependencies so they can be encoded compactly and partially resolved without bookkeeping. This allows the 𝜇op generator to simplify the dependency handling offline, avoiding much of the expensive cross-virtual-core synchronization at runtime. We implement VDCores on asynchronous GPUs with different architectures including NVIDIA GH200, H100, and RTX 6000 Pro GPUs. We evaluate it with four representative LLM inference workloads. As highlighted in Figure 2, VDCores achieves 24% higher decoding throughput on average than state-of-the-art hand-optimized kernel and megakernel baselines, and improves performance further by up to 77% under dynamic inputs. At the same time, VDCores reduces kernel programming and specialization effort by 90%. We have open sourced VDCores at https://github.com/vdcores/vdcores.
2
Motivation
2.1
Asynchronous GPU Programming and Performance
Existing ML frameworks [35] lower high-level operators into a sequence of GPU kernels and launch them on the device. 2
Mirage Persistent Kernel kernel launch
hidden data dep
VDCores
Kernel
Kernel A Memory Unit (TMA)
M1
Compute Unit (CUDA/Tensor Core)
Bubbles M2
Compute Unit (CUDA/Tensor Core)
M1
Kernel B M1
Cross-task Bubble M2
M3
C1
C2
C1
C2
M2
M2
M1
M3
...
C1
C2
C1
C2
C3
C3
Runtime Reorder
uop Memory Unit (TMA)
uop dependency
VDCores Dynamic Fusion Execution Time
Auto-overlapping
Figure 3. Asynchronous unit utilization under the kernel execution model. VDCores overcomes this limitation with asynchronous and independent execution.
Figure 4. Memory-bandwidth utilization comparison on a single Llama3-8B layer. MPK underutilize memory bandwidth due to inefficient inter-task overlapping.
A kernel is the standard GPU programming and execution unit: a host-launched device function that executes the same program across many threads over different data elements. This abstraction is becoming increasingly strained as GPU hardware grows more heterogeneous and asynchronous. As shown in Figure 1, achieving high performance increasingly requires fully utilizing asynchronous units (e.g., Tensor Cores and TMAs) at the same time. This trend reflects a broader architectural shift in GPUs (e.g., Blackwell [29], AMD [2]) and accelerators (e.g.TPU [24], Groq [19]) toward asynchronous execution. However, on asynchronous GPUs, the kernel abstraction creates substantial programming complexity. The difficulty appears along two fronts. First, programmers must manually coordinate heterogeneous warp roles inside SIMT code, e.g., using warp specialization to embed producer and consumer behaviors into conditional control flow and explicitly synchronizing their overlap. Second, performance depends on manually orchestrating deep software pipelines, e.g., multibuffering. Existing solutions reduce parts of this burden but do not eliminate the require of manual, static and monolith orchestration. Further, the monolith kernel abstraction packs asynchronous resources into a single opaque execution unit, itself limits the utilization of asynchronous units. A kernel bundles multiple resources and execution phases under a single start and finish, so the hardware cannot reuse partially idle resources across kernel boundaries for sub-operations, causing bubbles in execution. In Figure 3, between two adjacent kernels mapped to the same hardware unit, one kernel may still be performance writing back while some of its resources, such as the compute engine, are already idle, and cannot be used by the next kernel. Existing solutions like
programmable launch control and prologue/body/epilogue staging only mitigates this overhead. As in Figure 4, even MPK applies inter-task prologue-body overlapping, significant task-boundary-aligned drops in memory utilization are still visible during execution, indicating that much of the available hardware parallelism remains difficult to express and easy to leave underutilized. 2.2
Kernel Specialization Efforts
Modern ML execution is no longer a fixed sequence of uniform operators over stable input shapes. e.g., S-LoRA [36] shows that serving many adapters introduces dynamic weights, varying ranks, and heterogeneous batches, requiring both flexible memory management and specialized execution strategies. Monolithic kernels struggle to adapt to dynamic cases: a kernel optimized for one regime (e.g., low-latency decoding or single-adapter execution) often performs poorly in another (e.g., long-context decoding, large heterogeneous batches, or mixed-adapter serving). As a result, efficient execution increasingly relies heavily on kernel specialization, including kernel fusion, and megakernel designs. However, much of this specialization on asynchronous GPUs changes only how resources are orchestrated, including fixed operator bundles, buffer pipeline adjusting, and nearby producer–consumer fusion. e.g., the fusion support interfaces in Megatron-Core [37] (11 fusion modules) and vLLM [25] (9 fusion passes) are dominated by epilogue/prologue-style fusion. This specialization is simple in logic, but creates a huge programming overhead with a monolithic kernel model, which creates a cross-product problem: even small orchestration changes can cause large performance shifts and therefore require new kernel variants. 3
Zijian He, Adrian Sampson, Yiying Zhang, and Zhiyuan Guo
load.dep V
VCC msg queues (m2c, c2m)
load.dep V
VMC
SM
matvec.tile K
load Mt
rope.tile M/4
store Ot load T
matvec.tile store.smem
GMEM
load.dep Ot store Rt
load.smem
1
Tiled schedule (on 4 SMs) Auto Overalpping
DRAM VMC
Ot 2
1
VCC
4
3
5
1
6
2
rope.tile
VMC
VCC
VMC
VCC
SM
store O
load.dep(V)
SM
matvec.tile K
load.mat Mt
rope.tile M/4
store.smem S0 load T
Figure 5. VDCores Abstract machine model.
load.smem S0 store Rt
Beyond the monolithic kernel abstraction. Taken together, these limitations call for a direct, resource-facing model for asynchronous GPUs. Rather than packaging resources, the model should expose them as independent work units that can be composed explicitly and scheduled opportunistically at runtime.
3
2
Tiled schedule bypass global memory Dynamic Fusion
SRAM VMC
S0 2
1
VCC
4 1
3 5
6
2
Figure 6. VDCores Execution Example. First, VDCores virtual cores reorders the execution to fill execution idle gaps, i.e., auto-overlapping between 𝜇ops. In 1 , the load Mt 𝜇op can be reordered and executed before load.dep V because the former have no dependency, while load.dep V is could still be waiting on its dependent instructions to finish. On VDCores, any 𝜇op with ready dependency be executed first, reducing the runtime stalls casued by static scheduling and coarse-grained dependencies. Further, VDCores 𝜇op builder (§ 4.2) exploits this dynamic execution model to generate 𝜇op sequences with fewer dependency-caused stalls, enabling optimization across operator boundaries. One example is dynamic fusion, which promotes global-memory communication to shared memory for one-to-one dependency edges between memory 𝜇ops on the same virtual memory core. As shown in schedule 2 , the VDCores 𝜇op optimizer detects this pattern and rewrites store and load into store.local and load.local. At runtime, the dependency is resolved through local queues with low overhead. In effect, this achieves operator fusion without requiring a manually written fused variant.
VDCores Overview
VDCores serves a role similar to CUDA in GPU programming: it provides a programming abstraction and execution model for asynchronous GPUs (§ 4.1). It is built around a new hardware execution-unit abstraction, decoupled cores, including virtual memory cores (VMCs) and virtual compute cores (VCCs), and a new programming and execution unit, micro-operators (𝜇ops). VDCores serves as a drop-in replacement for the CUDA backend. For GPU applications, e.g., performing matrix multiplication in PyTorch, CUDA converts them into a list of kernel launches, while VDCores converts them into a 𝜇op graph (§ 4.2), consisting of 𝜇op streams for virtual cores, defining the detailed operations to execute, and connected by dependency edges, defining the execution order of the 𝜇ops, as illustrated in Figure 5. VDCores then submits the constructed 𝜇ops to its on-GPU executors, virtual decoupled cores (§ 4.3), for execution. Within each virtual core, 𝜇ops execute in a dependency-driven manner, following dataflow order [6] rather than instruction order. After completing one 𝜇op, the executor selects another ready 𝜇op for execution and continues until the entire 𝜇op graph finishes. Execution Example. We illustrate VDCores execution flow using an instruction flow of two ML operators: a matrix-vector multiplication (𝑂 = 𝑀@𝑁 ), connected by a RoPE rotation (𝑅 = 𝑟𝑜𝑝𝑒 (𝑂,𝑇 )), as shown in Figure 6. The VDCores scheduler decomposes operators into a a list of dependency-connecting compute and memory 𝜇ops. With 4 pair of VCCs and VMCs available, VDCores decomposition is to tile the matrix along the 𝑀 dimension and assign the work to each virtual core pair. Each virtual core’s 𝜇op sequence is shown as 1 . On each virtual core, VDCores execute 𝜇op based on dependency and resource readiness. We illustrate the process and it’s benefits with two execution examples.
4
VDCores Design
In this section, we introduce the three major components materialize the decoupled programming and execution model on GPUs: VDCores decoupled model, VDCores 𝜇op generator and VDCores virtual cores. 4.1
The VDCores Decoupled Model
This section introduces the decoupled programming model VDCores adopt and how it shifts the programming and execution paradigm on asynchronous GPUs. 4.1.1 VDCores Abstraction of GPU Hardware. VDCores decomposes GPU hardware execution units (streaming processors, SMs, or compute units, CUs) into resourcespecialized cores: Compute Cores and Memory Cores. Each core resembles a lightweight register machine; VDCores further specifies its architectural state (e.g., register file, localmemory management) and programming interface to access 4
frees the programmer from manually orchestrating asynchronous coordination across hardware resources. VDCores pre-defines a list of compute, memory, and shared control 𝜇ops that support common types of ML operations.
1 VDC_DEFINE_UOP(COP_MATVEC, h_matvec) 2 __vdc_cop__ h_matvec(VdcInst &inst, VdcCtx &ctx) { 3 auto c = ctx.alloc_registers(NCREGS); 4 // the number of tiles are encoded in size field 5 for (int i = 0; i < inst.size; i++) { 6 float *a = ctx.m2c.pop_wait(); 7 float *b = ctx.m2c.pop_wait(); 8 matvec_tile_accumulate(a, b, c); 9 ctx.c2m.push(a, b); 10 } 11 float *c_shared = ctx.m2c.pop_wait(); 12 copy(c, c_shared); 13 ctx.m2c.push(c_shared); 14 }
4.1.3 VDCores Execution Model. VDCores implies dependency-driven execution model. Each 𝜇op may depend on resources produced or owned by other virtual cores in order to complete its task. For example (Figure 7), a matvec 𝜇op requires shared-memory tiles containing the input matrix and vector tiles during execution. In VDCores, such relationships are represented explicitly as dependencies edges between 𝜇ops, marking a execute-after relationship. This execution model improves asynchronous hardware utilization in two ways. First, by removing ML-operator execution boundaries and merging work into a single 𝜇op graph, it prevents resources from being blocked behind unrelated work. Second, by resolving execution at runtime from actual resource readiness, it naturally adapts to dynamic execution conditions instead of relying on a fixed static schedule. However, VDCores’s dependency-driven execution model can be expensive to realize on GPUs, due to the major overhead of dependency tracking. Conventional dataflow and outof-order architectures [4] rely on dependency scoreboards that track every issued operation, incurring substantial cycle and memory overhead on GPUs. VDCores addresses this challenge in two ways. First, it reduces the complexity of the dependency graph by restricting and simplifying the dependency model. Second, instead of relying on runtime, cross-core dependency resolution, VDCores uses co-design to encode dependency structure into 𝜇ops ahead of time and to localize dependency resolution as much as possible, while preserving the same execution flexibility. Encoding dependency in 𝜇op. VDCores restricts each memory 𝜇op to have at most one inter-memory-𝜇op dependency, plus an optional dependency to computation 𝜇op executing on the same SM. This simplification eliminates the need to represent and resolve an arbitrary dependency graph in hardware. Instead, VDCores encodes dependency information directly into 𝜇ops in a GPU-friendly form. Dependencies between 𝜇ops in local memory core and compute cores are represented using flags on the memory 𝜇op side. For instance, when a memory 𝜇op carries a send flag, it forwards its allocated memory region to a local compute core. Conversely, a recv flag indicates that the 𝜇op consumes a region sent from a local compute core. Similarly, dependency between two 𝜇ops are also encoded into a depId field, and its direction is determined by the send/recv flags. Thus, when generating 𝜇ops, VDCores does not store explicit pointers to dependent 𝜇ops, and instead tracks dependencies entirely through these encoded fields.
Figure 7. Simplified new compute 𝜇op matvec. Kernel programmers use dependency queues m2c and c2m, to acquire and release memory resources with push/pop.
them. Each type of core executes a flow of specific microoperations (𝜇ops): Compute Cores execute compute 𝜇ops, while Memory Cores execute memory 𝜇ops. A 𝜇op is the smallest unit of execution and programming unit in VDCores. Each 𝜇op describes a unit of task within the capability of the execution unit when the resource is ready. Memory 𝜇ops in VDCores include memory movement and management operations, e.g., load and store. For example, the memory 𝜇op, load.dep in Figure 5 loads a memory block from addr in global memory to the local memory. Compute 𝜇ops are computation atoms that can be completed within a single resource domain (e.g., a local tensor core) and encode tile-granular semantics. For example, the compute 𝜇op attn performs the fused attention score computation. Control 𝜇ops implement control flow like conditions and branching. VDCores further provides control 𝜇ops, including loop, continue_if, which can execute on all virtual cores. Together, these 𝜇ops give each virtual core a small programmable execution context, enabling it to express complex GPU computation patterns beyond a fixed instruction stream. 4.1.2 VDCores Programming Model. Similar to programming kernels in CUDA, on VDCores, GPU programmer programs 𝜇ops to extend VDCores’s capabilities or support new machine learning operators, e.g., a new attention operator such as LinearAttention or new memory operations for data movement over GPU-GPU interconnects. As shown in Figure 7, adding a new 𝜇op involves two steps. First, the developer defines a new opcode and specifies how the fields in the 𝜇op instruction word should be interpreted. Second, the developer implements and registers a GPU device-side 𝜇op handler that executes the 𝜇op. It takes the instruction word and the virtual core context as input. Programmers uses the context to request the service of runtime, including resolving dependency, allocating resource and changing register states. As shown in Figure 7, it can use the m2c and c2m dependency queues to acquire memory resources from memory cores and release them after use. Compared with writing monolithic kernels, adding a new 𝜇op 5
Zijian He, Adrian Sampson, Yiying Zhang, and Zhiyuan Guo
Tiling = m64n16k128 load W gemm.m64n16 4
load X store O
DAG. Each operator is associated with a set of valid decompositions into 𝜇ops, provided by the 𝜇op builder, and each 𝜇op and dependency edge carries a cost model that estimates its execution cost on virtual cores. A critical path is considered dominant if its cost exceeds the average per-core cost by more than a threshold. The scheduler then refines operators on these paths by changing their tilings until the expected benefit becomes marginal or no dominant critical path remains. This process drives globally optimized execution while also reducing direct dependency-resolution overhead.
Tiling = m128n16k64 + reduce
gemm.m64n16 2 gemm.m64n16 2
load W load X reduce O
Figure 8. VDCores Composes 𝜇ops to Dynamically Tiles Shapes.
Virtual-flow-based dependency-driven 𝜇op execution. To enables the runtime to quickly identify 𝜇ops that are independent and can therefore execute in parallel, without performing general dependency checks at runtime. To support this, VDCores 𝜇op generate (§ 4.2) assign each VDCores 𝜇op a virtualFlowId that encodes coarsegrained dependency structure among 𝜇ops. 𝜇ops connected by direct dependency edges are assigned to the same virtual flow, while independent 𝜇ops are placed in different virtual flows. At runtime, instructions within the same virtual flow execute in order, whereas instructions from different virtual flows may be reordered and overlapped when resources permit. 4.2
4.2.2 Dependency Graph Lowering. VDCores next embeds inter-𝜇op dependencies directly into the per-virtual-core 𝜇op bytecode. It first performs virtual flow assignment, decomposing the dependency graph on each virtual core into independent chains and assigning one virtual flow to each chain. This decomposition remains possible even with control flow loops and branches, because computation 𝜇ops introduce no direct dependency between each other, and each memory 𝜇op has at most one incoming inter-memory dependence. As a result, independent 𝜇ops can bypass stalled ones instead of being serialized behind them in a single producer pipeline. VDCores then checks for and eliminates deadlocks by elaborating resource allocation along 𝜇op flow on each core, especially slot-based local memory allocation, and detecting potential over-allocation. When necessary, it reorders instructions without violating dependence constraints. For example, it may move an independent store ahead of a later load to free resources earlier. This step is particularly important when one VMC serves multiple VCCs. Finally, once dependencies are embedded, VDCores removes redundant dependency checking. When two 𝜇ops are placed on the same virtual flow, in-order execution already enforces their readiness order, so explicit barriers implied by transitive dependences can be eliminated. This reduces polling and control overhead while preserving correctness.
VDCores 𝜇op Generator
VDCores 𝜇op generator takes highlevel machine learning operator graph (e.g., PyTorch computation DAG) and convert them into 𝜇op, assign them to virtual cores and build dependency between them. VDCores resolves this progressively across two tiers to balance scheduling effectiveness and scheduling overhead. First, it chooses how to decompose ML operators to jobs on multiple cores. With a global view of the cross-operator DAG, this level focuses on exposing task-level parallelism and balancing load across virtual cores. Then, it optimizes the generated per-core 𝜇op flows. This level focuses on memory-𝜇op placement, fine-grained dependence resolution, and instruction execution efficiency within each virtual core.
4.2.3 Dynamic Data Placement and Fusion. Several optimizations of VDCores 𝜇op generation requires multiple steps to work together, including dynamic data placement. dynamic data placement change the placement of memory 𝜇ops between neighboring 𝜇ops with one-to-one dependency. For example, intermediate data can be promoted to shared memory to enable fusion-like locality, or kept separate to expose more overlap opportunities. The scheduler is aware of this optimization and accounts for the reduced cost of one-to-one dependencies. This design differs from conventional kernel fusion. Because VDCores already decouples memory movement and computation into separate 𝜇ops, composing them requires neither code generation nor recompilation. The same computation can therefore be realized either as a locality-oriented fused pipeline or as a more weakly coupled flow that exposes additional overlap opportunities.
4.2.1 Adaptive Scheduling and Tiling. VDCores first schedules the DAG of logical operators to virtual cores, similar to existing GPU backends. The key difference is that VDCores schedules over the full DAG (instead of single DAG node) and adapts execution at runtime by composing 𝜇ops directly, rather than by selecting from precompiled specialized kernels. We refer to this flexibility as adaptive scheduling. Examples include partitioning a GEMM along the 𝑀, 𝑁 , or 𝐾 dimensions, or partitioning attention across token and head blocks. Different tilings results on different workload to be assigned on each computation unit, where VDCores could dynamically build through composing 𝜇ops instead of requiring a new kernel. Given this flexibility, VDCores chooses tilings by iteratively eliminating dominant critical paths in the execution 6
Virtual Compute Core thread×128
CFU
CFU
SMEM
thread×32
Alloc
Gen. Address dep. resolve dep.resolve
EXEC
M2C
placing all work in a single GPU loop, this design improves effective 𝜇op throughput by 4.2×. Within each pipeline stage, VMC further parallelize control and data movement by leveraging parallel GPU threads. In the CFU, thread-level parallelism and software pipelining allow 𝜇op decode and address generation to proceed concurrently with instruction fetch and register access. Managing local memory is likewise parallelized through a SIMT allocator. Each 𝜇op may request an arbitrary number of contiguous slots, and allocated slots may be freed in an order different from allocation (e.g., when implementing FlashAttention 𝜇op). VMC tracks memory availability as a 32-bit bitmask in shared memory and uses multiple threads to search candidate allocation positions in parallel. Each allocator thread probes for a slot starting from its thread ID, and the warp then uses a warp-level voting primitive to identify the first available location. Parallel virtual flow execution: by dynamically mapping virtual flows onto physical execution units. At runtime, the CFU maintains a mapping from each virtual flow to a physical execution unit. For a given 𝜇op, the CFU dispatches it according to this mapping. 𝜇ops within the same flow are always assigned to the same execution unit and execute in order, preserving intra-flow ordering. When the number of active virtual flows exceeds the number of physical units, the CFU multiplexes multiple flows onto the same unit. This design preserves dependency semantics independent of the underlying runtime implementation. Non-dependent operations can execute in parallel across multiple LDUs and STUs. However, within each individual LDU or STU, 𝜇ops execute strictly in order, following the sequence of dependency resolution, operation execution, and dependency message delivery. As a result, a 𝜇op may spinwait on an unresolved dependency and temporarily block later 𝜇ops assigned to the same unit. For example, within an LDU, one 𝜇op may wait on a pending dependency and thereby stall subsequent 𝜇ops in that unit. This design preserves ordering and dependencies without creating a single hot spot of dependency tracking in VMCs.
Virtual Memory Core
reclaim STU
LDU1
LDU2
thread×32
thread×32
thread×32
writeback
readiness
C2M
Figure 9. VDCores virtual executor. The example shows executors on a single H100 SM, with two VCC and single VMC.
4.3
VDCores GPU Executor
Virtual cores are software 𝜇op interpreter built on top of GPU hardware execution units. Each virtual core is launched once as a persistent kernel at the beginning of execution. VDCores runtime streams new 𝜇ops to virtual cores for execution when new requests arrive. A core design challenge of virtual cores is how to preserve the flexibility of VDCores decoupled model without introducing excessive runtime overhead. VDCores addresses this challenge by drawing inspiration from classical microarchitecture while restructuring the design to match GPU hardware. At a high level, it forms a software analogous to a pipelined superscalar microarchitecture: multiple virtual execution units are organized as a wide pipeline, with each unit mapped onto GPU warps and register state, while explicitly leveraging SIMT parallelism within them. 4.3.1 Virtual Memory Core. Virtual memory cores (VMCs) execute memory and control 𝜇ops and manage their internal state (e.g., loop counters) and memory resources (e.g., shared-memory capacity within an SM). VMC 𝜇op execution throughput is critical to VDCores performance. For example, an NVIDIA H200 GPU has up to 4TB/s DRAM bandwidth. The VDCores runtime must sustain over 10K tile-level load/store 𝜇op/s to avoid becoming the performance bottleneck. VMC employs two levels of parallelism. The first comes from building an 2-stage execution pipeline for each 𝜇op. Each 𝜇op is first processed by a front-end control-flow unit (CFU), which manages register state, handles control 𝜇ops, and performs address generation. The decoded 𝜇op and generated address is then issued to multiple parallel load units (LDUs) and store units (STUs), where dependency resolution and memory movement are handled independently. This pipelined design reduces the initiation interval for processing a 𝜇op: throughput is determined by the slowest pipeline stage rather than the sum of all stages, and instruction sequences no longer suffer from head-of-line blocking. Compared with
4.3.2 Virtual Compute Core. Virtual compute cores handle compute resources, including register files, SIMT cores and asynchronous matrix-computation units (e.g., Tensor Cores). Each VCC contains a control-flow unit (CFU) and multiple execution units (EU) for executing computation 𝜇ops. As an example, in Figure 9, VDCores by default create two EU per SM, each handling 128 SIMT threads and attached CUDA core resources, when applicable. Unlike VMC, the CFU and EU in VCC are not mapped to separate parallel threads. Instead, they are organized as a software pipeline. The CFU state is stored in EU-shared memory, and for each instruction, the CFU runs first on EU’s thread set, maintaining and preparing state (e.g., updating the GPR value), in a manner similar to an operating-system 7
Zijian He, Adrian Sampson, Yiying Zhang, and Zhiyuan Guo
GPU #SMs ShMem KB/SM DRAM BW H100 132 228 KB 3.35 TB/s GH200 132 228 KB 4.00 TB/s RTX6000 Pro 188 100 KB 960 GB/s
kernel, before yielding to the EU to execute the target 𝜇op. This design is motivated by the low density of control 𝜇ops in VCC execution: more than 98.6% of execution time is spent in the EU performing actual 𝜇op execution.
Table 1. Hardware configuration used in end-to-end evaluation. On each system VDCores deploys 1 VMC and 2 VCCs per SM, with 8KB memory slot size.
4.3.3 Dependency Resolving. On VDCores, dependency are resolved at runtime by passing dependency messages through first-in first-out (FIFO) messages queues connecting virtual cores. Following VDCores execution model, VCCs only communicate with VMCs on the same execution unit, while all VMCs can talk to each other. VMC-to-VCC messages include shared memory regions that have been loaded from global memory and thus can be used by the compute units. VCC-to-VMC message queues return used or generated memory regions to be recycled or written back. Similarly, VMCs communicate through a global queue allocated based on its depid. Between VMCs, messages passes the ownership of global memory resources (e.g., after writing to a global memory block).
Portable Runtime. We port the runtime on three different NVIDIA GPUs with different microarchecture specifications, showing the portability of VDCores. As summarized in Table 1, each of them has different number of SMs, shared memory size, and global memory bandwidth. They also feature different asynchronous execution unit configurations. e.g., GH200 and H100 additionally support asynchronous tensor core execution with WGMMA instructions. These difference is hidden by the VDCores virtual core and 𝜇op abstraction. same 𝜇op graph executes on both architectures, while matvec 𝜇op are backed by different implementations. Message Passing on Virtual Cores. Both dependency queues and inter-execution unit communication in VDCores are implemented through message passing. We realize this mechanism using 𝑚𝑏𝑎𝑟𝑟𝑖𝑒𝑟 instructions [31], which provide efficient synchronization while enabling space-sharing of hardware resources. Inside each SM, all execution units, including VMC and VCCs, share the same four physical execution slots. If an execution unit blocks on a message (e.g., at pop_wait()), it is de-scheduled and resumed only when the awaited message arrives. As a result, VDCores avoids spin waiting and reduces wasted cycles from idle execution units. Dynamic address generation VDCores supports dynamic address generation for every 𝜇op to produce runtime-dependent addresses. In VDCores, an instruction may be marked with a dynamic flag, indicating that its address field (which may encode an address, an 𝑛-dimensional tensor coordinate, or the index of a TMA descriptor) is generated dynamically and needs to be computed by adding an offset from an accumulator register. Further, the accumulators are designed to be updated by control 𝜇ops, e.g., by repeat based on loop counter. This mechanism substantially reduces the number of instructions needed for recurring patterns, which are common in LLM inference workloads, e.g., full inference of llama-8b is encoded in only 224 𝜇ops per virtual core.
4.3.4 Avoiding Deadlock. Dependency-driven execution introduces a risk of deadlock when outstanding 𝜇ops across virtual cores form a cyclic wait relationship on shared resources. A typical case arises between a VMC and a VCC: the compute side waits for inputs to be loaded into shared memory before producing outputs, while the memory side waits for memory slots to be released before issuing additional load 𝜇ops. Since VDCores forbids cyclic 𝜇op dependencies, deadlocks can only arise from resource-allocation cycles rather than dependency cycles. VDCores prevents such deadlocks by linearizing resource acquisition through a joint scheduling and runtime protocol, summarized as in-order allocation and out-of-order execution. The VDCores 𝜇op generator (§ 4.2) first constructs a baseline 𝜇op order that is deadlock-free under in-order execution. The runtime then preserves this order during local resource allocation (e.g., shared-memory slots), while still allowing later execution steps to proceed out of order once their dependencies are satisfied (e.g., loading a memory slot and sending notification). This avoids deadlock because resource allocation remains acyclic in the preserved baseline order, and once allocated, blocked 𝜇ops wait only on acyclic dependencies, so some ready 𝜇op can always execute and eventually release resources.
5
Implementation
We implement VDCores’s executors and 𝜇op generator for NVIDIA Hopper and Blackwell GPUs in 5K lines of C++/CUDA. On top of this runtime, VDCores provides both low-level 𝜇op execution APIs and high-level ML-operator execution interfaces in 4K Python LoC, enabling integration with PyTorch as a backend. We highlight the following implementation details further help VDCores to improve execution efficiency:
6
Evaluation
We answer three key questions in this section: 1. Does VDCores improve end-to-end model execution over state-of-the-art kernel and megakernel systems? (§ 6.1) 2. How much does each key optimization in VDCores contribute to the performance gain? (§ 6.2) 8
VDCores Qwen3-1.7B
Torch + TK Qwen3-8B
TK-llama1B vLLM Llama3.2-1B
SGLang Llama3.1-8B
1.33x
1.37x
1.41x
1.53x
1.28x
1.31x
1.34x
1.47x
1.04x
1.46x
1.54x
1.68x
1.27x
1.32x
1.43x
1.47x
1.38x
1.38x
1.35x
1.55x
1.11x
1.13x
1.11x
1.30x
1.01x
1.46x
1.44x
1.61x
1.16x
1.18x
1.24x
1.27x
1.37x
1.48x
1.49x
1.59x
1.18x
1.11x
1.11x
1.16x
1.34x
1.38x
1.39x
1.39x
1.18x
1.10x
1.10x
1.16x
1
2
4
8
1
2
4
8
1
2
4
8
1
2
4
8
0.5
GH200
0.0 1.0 0.5 0.0
RTX6000 Pro
Normalized Throughput
H100
1.0
Mirage
1.0 0.5 0.0
Batch Size
Figure 10. End-to-end decoding performance over LLM inference. Throughput is normalized to VDCores. The number on each group shows VDCores’s performance gain over the best baseline system. TK: ThunderKittens.
3. Is software-managed decoupled execution practical given the runtime overhead introduced by VDCores? (§ 6.3)
6.1
exposes only limited overlap opportunities. ThunderKittensLlama1B removes most stalls through manually designed fusion and prefetching specialized for this model architecture and batch-size-one regime, reaching roughly 96% of VDCores’s performance. The remaining gap is consistent with our argument that compile-time task-coupled scheduling, even when highly optimized, still leaves some overlap opportunities uncovered. Programming effort. Beyond performance, VDCores substantially reduces the need for manual specialization. As shown in Table 2, VDCores implements the evaluated Llama31B pipeline with only 6 reusable 𝜇ops and 741 lines of GPU code, while requiring no task-specific fused implementation. In contrast, Mirage and TK-Megakernel rely on 8 and 7 independent operators, respectively, together with 3 and 5 fused tasks, and require roughly 3x more GPU code. This gap reflects a fundamental difference in where optimization complexity lives. Kernel-centric systems realize performance through task-specific fusion and specialized implementations, so both code size and maintenance cost grow as operators, surrounding stages, or target hardware change. By contrast, VDCores reuses the same small set of 𝜇op building blocks and changes only their composition and scheduling, shifting optimization effort from per-task kernel variants to a reusable runtime substrate.
End-to-End Performance and Coding Effort
We first evaluate VDCores on end-to-end LLM inference under offline decoding. We study four representative models—Qwen3-1.7B, Qwen3-8B, Llama3.2-1B, and Llama3.1-8B. Our setup includes KV cache and paged attention. Without continuous batching, we experiment with offline decoding of a fixed batch for 64 steps from a 128-token context. We report token-generation throughput for batch sizes 1 to 8. We compare against strong baselines spanning the main current design points: vLLM [25] and SGLang [50] represent optimized kernel-per-operator execution with JIT generation and CUDA Graphs, while Mirage [48] and ThunderKittensllama1B [42] implementations represent expert-tuned megakernels. We also include Torch+ThunderKittens as a lowerengineering baseline using PyTorch’s execution stack plugged with an optimized attention kernel. Figure 10 shows the end-to-end throughput results. Across all 48 evaluated combinations of model, hardware, and batch size, VDCores consistently achieves the lowest per-token decoding latency among both kernel-per-operator and megakernel baselines. Relative to the best baseline in each setting, VDCores delivers a 1.31x geometric-mean speedup, corresponding to a 23% average reduction in per-token latency, and up to a 1.68x performance improvement. The gains are especially pronounced on smaller models, where shorter task durations amplify the cost of underutilization across kernel boundaries. MPK reduces kernel launch overhead through persistent-style execution, but its pipeline structure is still packed into each individual task and thus
6.2
Performance Deep Dive
To understand where VDCores’s benefits come from, we evaluate each key mechanism with targeted case studies to examine when each is effective. 6.2.1 Auto Overlapping. VDCores derives intra- and intertask overlap directly from dependency-driven execution of 𝜇ops. To isolate the benefit of this mechanism, we start from 9
Zijian He, Adrian Sampson, Yiying Zhang, and Zhiyuan Guo
System
# Independent LoC # Fused Fused LoC Avg. H100 Perf. Reusability kernels/tasks/𝜇ops tasks (norm.) (for new model and hardware) VDCores 6 741 0 – 1.00 High Mirage 8 2,339 3 1,073 0.83 Low (Monolith Tasks) TK-llama1b 7 2,065 5 596 0.95 Low (BS=1, Special Fusion) vLLM 14 6,424 8 3,788 0.63 High
Table 2. GPU Program Line of Code Comparison across Systems for implementing llama3-1b, along with qualitative reusability under new inputs and new architectures.
1.0 0.8 0.6 0.4 0.2 0.0
VDCs (no-overlap) VDCs (no-virtual-flow)
1.00x 1.00x 0.90x 0.91x 0.86x 0.92x 0.79x 0.77x 0.71x 0.72x
VDCs (full-overlap)
Normalized Performance
Normalized Performance
vLLM Mirage
1.00x 0.71x 0.73x 0.62x 0.32x
Llama3-8B
Qwen3-8B
Llama3-1B
VDCores
1.00 0.75 0.50 0.25 0.00
0.5K
0.5K-B
TK
64K
Mirage
RAND
Context-Length / Skew Setting
Figure 11. Auto overlapping deep dive. VDCs: Virtual Decoupled Cores.
Figure 13. Dynamic uneven-context evaluation (H100). x-axis
Execution Time (us)
TK (no-fuse) TK 50 40 30 20 10 0
44.23 36.94
Mirage (no-fuse) Mirage
VDCs (no-fuse) VDCs
labels: 0.5K = 1×512, 0.5K-B = 64×512, 64K = 1×65536, and RAND = 64×random[256,1024].
43.83 33.53
10.56 8.73
MLP Layer
6.2.2 Dynamic Fusion. Dynamic fusion in VDCores is realized by rewriting the memory path of intermediate data between neighboring tasks, rather than by introducing a new fused task implementation. We evaluate this mechanism on two representative manually fused patterns, QKV-Projection + RoPE and the MLP block. As Figure 12 shows, VDCores closely matches the performance benefit of manual fusion in both cases, demonstrating that composing 𝜇ops differently can exploit similar locality benefits without the need for specialized kernels. We further study Embedding + RMS, a fusion opportunity that emerges automatically from VDCores’s runtime placement decisions but is typically left unfused in expert-tuned systems. In this case, fusion reduces execution time from 4.30 to 3.10µs, a 28% improvement over its own non-fused path
17.47 6.80 5.30
QKV-Proj+RoPE
4.30 3.10
Embedding+RMS
Figure 12. VDCores Dynamic Fusion Compared to Manual Fused Kernels. a restricted version of VDCores that converts 𝜇op data dependencies into task issue barriers, forcing all instructions from a later task to wait even when their dependencies are already satisfied. This resembles a kernel-per-operator execution style. We then enable cross-task overlap by restoring fine-grained dependency-driven instruction issue. Finally, we enable virtual-flow assignment, which moderates head-of-line blocking by allowing ready 𝜇ops to bypass stalled ones. As Figure 11 shows, enabling cross-task overlap accounts for the dominant share of VDCores’s benefit across all workloads. Without it, VDCores becomes comparable to, and in some cases slower than, baseline systems, showing that VDCores’s advantage does not come from faster individual operators but from the overlap exposed by dependency-driven scheduling. Enabling virtual-flow assignment yields a further 5% improvement on top.
6.2.3 Adaptive Scheduling. VDCores adapts to changing workloads by recomposing 𝜇op flows and remapping work at runtime, rather than by building specialized task implementations. We study the effectiveness of this mechanism with two case studies: dynamic attention length and dynamic LoRA serving. Dynamic Attention Length. We evaluate dynamic attention under four H100 workload regimes that vary both sequence length and batch structure, spanning short-context, long-context, and mixed-context cases (Figure 13). Mirage 10
sLoRA 2.44x
2000 1000 0
3.00x 3.47x
3.03x
3.32x
System
Sched. Time (ms)
Mirage TK (attn) vLLM SGLang
27000 2.28 352 261
VDCores
1.4
1 2 4 8 16 Distinct LoRA requests per batch
Throughput (GB/s)
Execution time (ms)
VDCores
16K
CUDA-core overhead
3.1%
1000 1K
2K
4K
8K
Block size
32K
Figure 16. VDCores (VDCs) memory operation efficiency at different block sizes.
PyTorch
1.5 1.0 0.5 L S Attn- Attn-
4 3 2 1
Figure 17. Runtime latency metrics and aggregate runtime overhead. Captured by NCU Compute Profiler.
3.28x 1.00x Single cycle
3.89x
4.19x
1.52x + Alloc + Co-routine+ Pipelined Current
Figure 18. Memory-core Optimization Breakdown.
Figure 15. Single-operator performance of VDCores compared with VDCores (manual), a manually tuned warpspecialization implementation, and with PyTorch.
Compute-bounded and memory-bounded kernels. We measure the runtime overhead of VDCores on representative compute- and memory-intensive operators. We compare against VDE-manual, a hand-crafted warp-specialized implementation with similar pipelining but minimal overhead, and against PyTorch as a reference single-kernel baseline. As shown in Figure 15, VDCores stays within 8% of peak performance on average, while sustaining over 82% of peak FLOPS and 93% of peak memory bandwidth on H100. Memory virtual core overhead. We sweep I/O block size and compare VDE’s effective bandwidth against the raw TMA issue rate to measure when virtual-core overhead is amortized. As Figure 16 shows, bandwidth approaches raw-instruction performance at 4KB for writes and 16KB for reads. Because real LLM operators typically move data at or above these granularities, the overhead of the virtual memory core is well amortized in practice. Figure 17 further breaks down the runtime cost over the lifetime of a virtual core and estimates that the aggregate overhead consumes only 3.1% of total core time. We ablate the key optimizations used to achieve this low overhead. We focus on instruction issue, because VMC throughput is largely determined by the control path that prepares and dispatches memory 𝜇ops. As Figure 18 shows, adding allocation support reduces latency from 366.9 to 242.2 cycles. Introducing coroutine execution further reduces latency to 112.0 cycles by improving overlap within the memory core, and pipelining lowers it again to 94.2 cycles. The final design reaches 87.6 cycles overall, yielding a 4.2× speedup over the naive implementation.
employs a fixed kernel template across these cases, and VDCores outperforms it by a large margin in all regimes, with up to 6.18× lower latency. ThunderKittens adopts a host-side schedule and processor assignment from the input workload, but this adaptation is still built around a fixed specialized kernel. VDCores consistently outperforms it by 15% on average. The gap is especially pronounced in the short-context regime, where the overheads imposed by that fixed kernel structure are harder to amortize. Dynamic LoRA serving. LoRA [23] specializes a shared base model with lightweight per-domain adapter weights. In serving, a single batch may mix requests that share the same adapter while others access different adapters. This creates highly variable effective matrix shapes, making fixed kernel schedules suboptimal. We compare against sLoRA [36], which uses stage-specific scheduling for the expand and shrink phases of LoRA execution. However, within each stage, the fixed kernel is largely agnostic to how requests are distributed across adapters. By contrast, VDCores adapts the 𝜇op flows of each adapter to the realized batch composition, improving SM occupancy and reducing makespan by up to 3.47x, as shown in Figure 14. 6.3
VDCs-read VDCs-write TMA-read TMA-write
422 ns 22ns 34ns 43ns 22ns
2000
Normalized performance
Performance normalized
VDCores (manual)
Copy GEMV GEMM-S GEMM-L RMS
Cycles
Start up latency 𝜇 op Initialization interval Ave memory load Ave memory st Ave memory ctrl
(b) Scheduling time
Figure 14. Adaptive Scheduling with dynamic LoRA serving. In-batch LoRA-adapter distribution follows.
0.0
Metric
3000
0
(a) Uneven-context performance
VDCores
4000
Runtime Efficiency Study
In this section, we focus on evaluating the design decisions in the virtual machine and executor by studying their performance implications. 11
Zijian He, Adrian Sampson, Yiying Zhang, and Zhiyuan Guo
7
Related Work
𝜇op layer, Compilers could now generate efficient handler for single leveraging their static packing and local schedule optimization. At system level, similar to dataflow compilers, compilers could be integrated into 𝜇op generator and help with translating ML operators into 𝜇op streams, leveraging the VDCores runtime harness of performance and correctness. This split preserves compiler efficiency where it works best, while avoiding ad-hoc runtime tuning and leave them for VDCores’s efficient runtime, also further support future learning-based or AI-assisted optimization. Evolving Asynchronous GPU. New GPU generations continue to introduce new asynchronous resource domains beyond traditional CUDA-core execution. For example, recent architectures such as Blackwell expose new tensor-memory inside each SM, and modern GPUs increasingly provide distributed shared-memory capabilities across SM groups. VDCores can absorb these hardware changes without changing the decoupled model: we can either extend existing virtual-core implementations (e.g., VMC/VCC behaviors) or add new virtual-core types for the new resource domain. For instance, a additional specialized computation virtual cores can be introduced with keeping the same dependency-driven coordination. On new virtual cores, VDCores runtime design could be reused, and it could integrate and overlap seamlessly into VDCores system. Supporting Heterogeneous Accelerators and Tiered Memory System. VDCores shows good fit for emerging heterogeneous accelerators which explicitly decouple the hardware units for computation, memory movement and communication. For example, in AWS Trainium [9], those jobs are handled by Neuron Core, DMA and CC-Core respectively. VDCores decoupled model extends naturally beyond GPUs, each resource domain is represented as a virtual core, and cross-domain coordination is expressed through explicit dependency exchange rather than implicit launch order. When all backends implement the same decoupled interface and memory model, VDCores can also serve as a bridge abstraction across different heterogeneous hardware. the compiler/runtime can compose end-to-end pipelines across dissimilar devices without rewriting operator semantics for each target. Compared with ad-hoc cross-device orchestration, this design preserves overlap opportunities while maintaining a uniform integration boundary, similar to inside the single GPU device.
Framework, Abstaction and Compiler for Asynchronous GPUs. CUTLASS, FlashInfer, and ThunderKittens provide optimized GPU libraries and templates that package expert implementation techniques such as tiling, pipelining, prefetching, and warp specialization [28, 41, 44]. Flux and Mirage Persistent Kernel further optimize communication and persistentkernel execution by structuring kernels into stages that overlap data movement and computation [10, 16]. These frameworks make it easier to build high-performance GPU kernels, but they still express execution as specialized kernels with preplanned overlap and synchronization. TVM, Relax, Fireiron, Triton, Cypress, and TileLang raise the abstraction level of GPU programming by generating optimized kernels from higher-level tensor, scheduling, or tile-based descriptions [13, 20, 26, 43, 45, 46]. TAWA [11] further targets asynchronous GPU kernels by automatically applying warp specialization and pipeline scheduling. These systems reduce programmer effort and improve intra-kernel optimization, but their generated programs remain primarily kernel-centric: tiling, fusion, layout, and synchronization decisions are made before execution. VDCores differs by fundamental programming model shift: exposing memory and compute work as separate dependency-connected 𝜇ops, and results in assigning the coordination into a runtime layer. Dataflow and Access-Execution Decoupled Architectures. Classic dataflow machines execute operations when their operands become ready, exposing parallelism through explicit dependencies rather than fixed instruction order [4, 5, 18, 34]. Decoupled Access/Execute architectures separate memoryaccess and compute streams to hide latency through producer– consumer queues [3, 12, 17, 21, 39, 47]. Many ML accelerators similarly use dataflow-style execution to make tensor movement, reuse, and pipeline structure explicit in hardware [1, 14, 15, 24]. These systems rely on specialized hardware, fixed dataflows, or general token-matching mechanisms. VDCores keeps the principles, but restricts them to GPU-friendly ML 𝜇ops and virtual flows and further use software cores to realize them on existing GPUs. Compilers for dataflow, DAE, and spatial accelerators map programs into explicit graphs, streams, or hardware dataflows. They typically perform dependence analysis, tiling, placement, buffering, and scheduling to expose locality and parallelism to the target architecture. Ember applies compiler optimization for decoupled accelerator to irregular embedding operations in recommender models, improving performance and performance per watt [38]. VDCores converts the GPU to virutal decoupled hardware, further leveraging the existing compiling methods for global optimization.
8
9
Conclusion
We present VDCores, a virtual decoupled-core programming framework and runtime for modern asynchronous GPUs. VDCores decouples asynchronous resource units, enables autooverlapping and dynamic-fusion optimizations for LLM serving, and enables composable optimization. We believe this design offers a practical path toward a unified, high-performance systems stack for next-generation LLM workloads.
Discussion
Compiler support for VDCores. We view VDCores as a suitable abstraction layer for high-performance compilers. At 12
References
Programming Languages and Operating Systems, ASPLOS ’14, pages 269–284. ACM, 2014. [15] Yunji Chen, Tao Luo, Shaoli Liu, Shijin Zhang, Liqiang He, Jia Wang, Ling Li, Tianshi Chen, Zhiwei Xu, Ninghui Sun, and Olivier Temam. DaDianNao: A machine-learning supercomputer. In Proceedings of the 47th Annual IEEE/ACM International Symposium on Microarchitecture, MICRO-47, pages 609–622. IEEE Computer Society, 2014. [16] 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. Mirage persistent kernel: A compiler and runtime for mega-kernelizing tensor programs, 2025. [17] Neal Clayton Crago and Sanjay J. Patel. OUTRIDER: Efficient memory latency tolerance with decoupled strands. In Proceedings of the 38th Annual International Symposium on Computer Architecture, ISCA ’11, pages 117–128. ACM, 2011. [18] Jack B. Dennis and David P. Misunas. A preliminary architecture for a basic data-flow processor. In Proceedings of the 2nd Annual Symposium on Computer Architecture, pages 126–132. ACM, 1975. [19] Groq Inc. Groq lpu architecture. https://groq.com/architecture/, 2023. Dataflow-style execution with explicit decoupling. [20] Bastian Hagedorn, Archibald Samuel Elliott, Henrik Barthels, Rastislav Bodik, and Vinod Grover. Fireiron: A scheduling language for highperformance linear algebra on gpus. In PLDI, London, UK, 2020. [21] Tae Jun Ham, Juan L. Aragón, and Margaret Martonosi. DeSC: Decoupled supply-compute communication management for heterogeneous architectures. In Proceedings of the 48th Annual IEEE/ACM International Symposium on Microarchitecture, MICRO-48, pages 191–203. ACM, 2015. [22] Ke Hong, Guohao Dai, Jiaming Xu, et al. Flashdecoding++: Faster large language model inference on gpus, 2023. [23] Edward J Hu, yelong shen, Phillip Wallis, Zeyuan Allen-Zhu, Yuanzhi Li, Shean Wang, Lu Wang, and Weizhu Chen. LoRA: Low-rank adaptation of large language models. In International Conference on Learning Representations, 2022. [24] Norman P. Jouppi et al. In-datacenter performance analysis of a tensor processing unit. In ISCA, Toronto, Canada, 2017. [25] Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph E. Gonzalez, Hao Zhang, and Ion Stoica. Efficient memory management for large language model serving with pagedattention. In Proceedings of the 29th Symposium on Operating Systems Principles (SOSP), Koblenz, Germany, 2023. [26] Ruihang Lai, Junru Shao, Siyuan Feng, Steven Lyubomirsky, Bohan Hou, Wuwei Lin, Zihao Ye, Hongyi Jin, Yuchen Jin, Jiawei Liu, Lesheng Jin, Yaxing Cai, Ziheng Jiang, Yong Wu, Sunghyun Park, Prakalp Srivastava, Jared Roesch, Todd C. Mowry, and Tianqi Chen. Relax: Composable abstractions for end-to-end dynamic machine learning. In ASPLOS, Rotterdam, Netherlands, 2025. [27] NVIDIA Corporation. Nvidia hopper architecture in-depth. https:// developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/, 2022. Technical blog. [28] NVIDIA Corporation. Cutlass: Cuda templates for linear algebra subroutines. https://github.com/NVIDIA/cutlass, 2023. [29] NVIDIA Corporation. Nvidia blackwell architecture. https://resources.nvidia.com/en-us-blackwell-architecture?ncid=pasrch-goog-587708, 2024. White paper. [30] NVIDIA Corporation. Nvidia tensor cores. https://www.nvidia.com/enus/data-center/tensor-cores/, 2024. Accessed: 2026. [31] NVIDIA Corporation. Cuda c++ programming guide: Asynchronous barriers. https://docs.nvidia.com/cuda/cuda-programming-guide/04special-topics/async-barriers.html, 2025. [32] NVIDIA Corporation. Cuda c++ programming guide: Asynchronous data copies. https://docs.nvidia.com/cuda/cuda-programming-guide/
[1] Dennis Abts, Jonathan Ross, Jonathan Sparling, Mark Wong-VanHaren, Max Baker, Tom Hawkins, Andrew Bell, John Thompson, Temesghen Kahsai, Garrin Kimmell, Jennifer Hwang, Rebekah Leslie-Hurd, Michael Bye, E. R. Creswick, Matthew Boyd, Mahitha Venigalla, Evan Laforge, Jon Purdy, Purushotham Kamath, Dinesh Maheshwari, Michael Beidler, Geert Rosseel, Omar Ahmad, Gleb Gagarin, Richard Czekalski, Ashay Rane, Sahil Parmar, Jeff Werner, Jim Sproch, Adrián Macías, and Brian Kurtz. Think fast: A tensor streaming processor (TSP) for accelerating deep learning workloads. In Proceedings of the 47th Annual IEEE/ACM International Symposium on Computer Architecture, ISCA ’20, pages 145–158. IEEE, 2020. [2] AMD. Amd cdna architecture: Enabling high-performance compute. https://www.amd.com/en/technologies/cdna, 2023. Highlights asynchronous compute and memory pipelines. [3] José-María Arnau, Joan-Manuel Parcerisa, and Polychronis Xekalakis. Boosting mobile GPU performance with a decoupled access/execute fragment processor. In Proceedings of the 39th Annual International Symposium on Computer Architecture, ISCA ’12, pages 84–93. IEEE Computer Society, 2012. [4] Arvind and David E. Culler. Dataflow architectures. Annual Review of Computer Science, 1:225–253, 1986. [5] Arvind and Rishiyur S. Nikhil. Executing a program on the MIT tagged-token dataflow architecture. IEEE Transactions on Computers, 39(3):300–318, 1990. [6] Arvind Arvind and D Culler. Dataflow architectures. Annual Review of Computer Science, 1:225–253, 11 2003. [7] Michael Bauer, Henry Cook, and Brucek Khailany. Cudadma: optimizing gpu memory bandwidth via warp specialization. In Proceedings of 2011 International Conference for High Performance Computing, Networking, Storage and Analysis, SC ’11, Seattle, Washington, 2011. Association for Computing Machinery. [8] Michael Bauer, Sean Treichler, and Alex Aiken. Singe: leveraging warp specialization for high performance on gpus. In Proceedings of the 19th ACM SIGPLAN Symposium on Principles and Practice of Parallel Programming, PPoPP ’14, page 119–130, Orlando, Florida, USA, 2014. Association for Computing Machinery. [9] Nafea Bshara. Aws trainium: the journey for designing and optimization full stack ml hardware. In Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 3, pages 4–4, 2024. [10] 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, 2024. [11] Hongzheng Chen, Bin Fan, Alexander Collins, Bastian Hagedorn, Evghenii Gaburov, Masahiro Masuda, Matthew Brookhart, Chris Sullivan, Jason Knight, Zhiru Zhang, et al. Tawa: Automatic warp specialization for modern gpus with asynchronous references. In 2026 IEEE/ACM International Symposium on Code Generation and Optimization (CGO), pages 255–267. IEEE, 2026. [12] Tao Chen and G. Edward Suh. Efficient data supply for hardware accelerators with prefetching and access/execute decoupling. In Proceedings of the 49th Annual IEEE/ACM International Symposium on Microarchitecture, MICRO-49, pages 46:1–46:12. IEEE Computer Society, 2016. [13] Tianqi Chen, Thierry Moreau, Ziheng Jiang, Lianmin Zheng, Eddie Yan, Haichen Shen, Yuwei Wang, Yida Hu, Luis Ceze, Carlos Guestrin, and Arvind Krishnamurthy. Tvm: An automated end-to-end optimizing compiler for deep learning. In OSDI, Carlsbad, CA, USA, 2018. [14] Tianshi Chen, Zidong Du, Ninghui Sun, Jia Wang, Chengyong Wu, Yunji Chen, and Olivier Temam. DianNao: A small-footprint highthroughput accelerator for ubiquitous machine-learning. In Proceedings of the 19th International Conference on Architectural Support for 13
Zijian He, Adrian Sampson, Yiying Zhang, and Zhiyuan Guo
04-special-topics/async-copies.html, 2025. Accessed: 2026. [33] NVIDIA Corporation. Cutlass blackwell forward attention main loop. https://github.com/NVIDIA/cutlass/ blob/a2439551c765c5393aebe557ee75d3a0412d2211/ examples/77_blackwell_fmha/collective/ sm100_fmha_fwd_mainloop_tma_warpspecialized.hpp, 2025. Accessed: 2025-11-20. [34] Gregory M. Papadopoulos and David E. Culler. Monsoon: An explicit token-store architecture. In Proceedings of the 17th Annual International Symposium on Computer Architecture, ISCA ’90, pages 82–91. ACM, 1990. [35] Adam Paszke, Sam Gross, Francisco Massa, Adam Lerer, James Bradbury, Gregory Chanan, Trevor Killeen, Zeming Lin, Natalia Gimelshein, Luca Antiga, et al. Pytorch: An imperative style, high-performance deep learning library. In Advances in Neural Information Processing Systems (NeurIPS), Vancouver, Canada, 2019. [36] Ying Sheng, Shiyi Cao, Dacheng Li, Coleman Hooper, Nicholas Lee, Shuo Yang, Christopher Chou, Banghua Zhu, Lianmin Zheng, Kurt Keutzer, Joseph E. Gonzalez, and Ion Stoica. S-lora: Serving thousands of concurrent lora adapters, 2024. [37] Mohammad Shoeybi, Mostofa Patwary, Raul Puri, Patrick LeGresley, Jared Casper, and Bryan Catanzaro. Megatron-lm: Training multibillion parameter language models using model parallelism. arXiv preprint arXiv:1909.08053, 2019. [38] Marco Siracusa, Olivia Hsu, Victor Soria-Pardos, Joshua Randall, Arnaud Grasset, Eric Biscondi, Doug Joseph, Randy Allen, Fredrik Kjolstad, Miquel Moretó Planas, et al. Ember: A compiler for embedding operations on decoupled access-execute architectures. In 2026 IEEE/ACM International Symposium on Code Generation and Optimization (CGO), pages 150–163. IEEE, 2026. [39] James E. Smith. Decoupled access/execute computer architectures. In Proceedings of the 9th Annual Symposium on Computer Architecture, ISCA ’82, pages 112–119. IEEE Computer Society, 1982. [40] Rupanshu Soi, Rohan Yadav, Fredrik Kjolstad, Alex Aiken, Maryam Mehri Dehnavi, Michael Garland, and Michael Bauer. Optimal software pipelining and warp specialization for tensor core gpus, 2025. [41] Benjamin Spector, Jordan Juravsky, Stuart Sul, Owen Dugan, Dylan Lim, Dan Fu, Simran Arora, and Christopher Ré. Thunderkittens: Simple, fast, and adorable ai kernels. In International Conference on Learning Representations (ICLR), Vienna, Austria, 2024. [42] Benjamin Spector, Jordan Juravsky, Stuart Sul, Owen Dugan, Dylan Lim, Dan Fu, Simran Arora, and Christopher Ré. Look ma, no bubbles! designing a low-latency megakernel for llama-1b. https: //hazyresearch.stanford.edu/blog/2025-05-27-no-bubbles, 2025. Technical blog. [43] Cypress Team. Cypress: A tile-based dsl for gpu programming. https: //github.com/cypress-dsl/cypress, 2024. [44] FlashInfer Team. Flashinfer: Efficient and flexible inference kernels for large language models. https://github.com/flashinfer-ai/flashinfer, 2024. [45] TileLang Team. Tilelang: A tile-level programming model for deep learning. https://github.com/tile-ai/tilelang, 2024. [46] Philippe Tillet, H. T. Kung, and David Cox. Triton: An intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages (MAPL), Phoenix, AZ, USA, 2019. [47] Kai Wang and Calvin Lin. Decoupled affine computation for SIMT GPUs. In Proceedings of the 44th Annual International Symposium on Computer Architecture, ISCA ’17, pages 295–306. ACM, 2017. [48] Mengdi Wu, Xinhao Cheng, Shengyu Liu, Chunan Shi, Jianan Ji, Man Kit Ao, Praveen Velliengiri, Xupeng Miao, Oded Padon, and Zhihao Jia. Mirage: A multi-level superoptimizer for tensor programs.
In 19th USENIX Symposium on Operating Systems Design and Implementation (OSDI), Boston, MA, USA, 2025. [49] Ted Zadouri, Markus Hoehnerbach, Jay Shah, Timmy Liu, Vijay Thakkar, and Tri Dao. Flashattention-4: Algorithm and kernel pipelining co-design for asymmetric hardware scaling, 2026. [50] Lianmin Zheng, Liangsheng Yin, Zhiqiang Xie, Chuyue Sun, Jeff Huang, Cody Hao Yu, Shiyi Cao, Christos Kozyrakis, Joseph E. Gonzalez, Ion Stoica, Clark Barrett, and Ying Sheng. Sglang: Efficient execution of structured language model programs. In Advances in Neural Information Processing Systems (NeurIPS), Vancouver, BC, Canada, 2024.
14