Why real-time inference wants a megakernel

In early 2025, during YC’s Winter batch, we tried to make a photograph hold a live conversation.

The product was called Dollyglot. The simplest description is Character.AI with real-time video: a character began as a single photograph, and a generative model had to keep producing the next video frame as the conversation unfolded.

That last part changed the engineering problem. An offline video model can take a minute to generate a few seconds of footage and still be useful. A live character cannot. If one frame is late, the character freezes. If frames are repeatedly late, the conversation feels broken. And when many characters share one GPU, the amount of work increases while the time available for the next frame does not.

What forced us to keep digging, however, was the cost. We could generate the video in real time, but serving it was too expensive: each GPU could sustain too few simultaneous characters. We did not merely need one session to run faster. We needed density: many more live sessions on the same GPU, without any of them missing their frame deadlines. For us, density was not an abstract performance metric. It determined whether the product could be economically viable.

To understand where that time goes, it helps to know what an inference engine normally asks a GPU to do. The GPU is used as an accelerator: it performs the large matrix multiplications and other parallel operations that dominate a model’s arithmetic, but it does not usually run the whole service by itself. The CPU is called the host, and the GPU is the device. The host accepts requests, groups them into batches, manages their state, decides what work is ready and dispatches that work to the device. The device executes the requested operations as GPU programs called kernels.

Modern engines such as vLLM and SGLang make this coordination remarkably efficient. They batch requests continuously, manage GPU memory and use techniques such as fused kernels and CUDA Graphs. Their exact architectures differ, but they preserve the familiar division of labour: a host-side engine schedules model execution, and GPU workers accelerate it.

That division is sensible. The host is flexible and good at managing changing requests; the device is exceptionally good at dense parallel arithmetic. But Dollyglot exposed a less obvious consequence: a machine can execute every individual calculation quickly and still execute the complete recurring program inefficiently. Time can disappear between calculations: ending one job, making its results available, deciding what runs next and starting the next job.

The question I kept returning to was: what if the individual operation was the wrong unit to optimize? Most inference software presents a model to the GPU as a sequence of separate jobs. What would happen if we instead gave the GPU the recurring program as a whole, including the schedule connecting those jobs, and kept that program alive between updates?

Following that question eventually led us to build the Wave Persistent Kernel (WPK): a compiler and runtime that turns a recurring model update into one persistent GPU program. It also led from Dollyglot to .wave. This article is my attempt to explain the idea from the ground up. I will first describe how an ordinary model reaches the GPU, then look at the techniques CUDA already gives us, and finally follow one update through WPK. The broader question is useful even if you never work on a live model: when does a convenient software boundary become the thing limiting the system?

The unit of execution

I find the easiest way to understand a megakernel is to start with a smaller question: what does the GPU consider to be one job?

An ML framework describes a model as a graph of tensor operations. A linear layer, a normalization, an attention operation and an activation appear as distinct nodes connected by tensors. This is a useful representation for writing and transforming a model. It does not dictate how the model should run on a GPU.

The usual unit of execution on a GPU is a kernel: one GPU program launched over a grid of threads. A framework or inference runtime translates the model graph into a sequence of these kernels. One kernel may perform a matrix multiplication, another may normalize its output, and another may apply an activation.

A megakernel moves the boundary. Instead of making each operation, or each small fused group of operations, a separate GPU program, it places a substantial part of the model’s execution inside one kernel. That kernel must now do work that used to be implicit in the sequence of launches: select or follow an execution schedule, coordinate groups of workers, respect dependencies and manage intermediate storage.

There is no universal operation count at which a kernel becomes “mega.” The useful definition is structural:

A megakernel is a single GPU kernel that owns the execution of a multi-operation program, including coordination between operations that would ordinarily be separate kernel launches.

This is different from a persistent kernel. A persistent kernel remains resident and waits for more work instead of exiting after one invocation. A kernel can be persistent without executing an entire model, for example a resident queue consumer that performs one type of task. A megakernel can execute a whole model step and then exit. The two properties are independent, although recurring real-time workloads often benefit from combining them.

WPK is both: a complete model update executes inside one physical kernel, and the same cooperative grid remains alive across updates.

A minimal model of GPU execution

To see why this boundary matters, we need a small amount of CUDA vocabulary. I will introduce it through one deliberately simple execution path.

When the host launches a kernel, it creates a grid containing many thread blocks, also called CTAs. CUDA assigns those blocks to streaming multiprocessors, or SMs, the processing units that make up the GPU. Threads in the same block can exchange values through fast on-chip shared memory and synchronize with one another. Different blocks are deliberately more independent: CUDA normally makes no guarantee about the order in which they run.

This independence is one reason kernels scale well. A grid may contain far more blocks than the GPU can hold at once, and CUDA schedules new blocks as resources become available. It is also why the end of a kernel is such a convenient boundary. Once all blocks have exited, the next dependent kernel can safely consume their results.

The boundary is convenient, but it has consequences:

1. The next kernel must be submitted and scheduled.
2. Values that must survive the boundary need a lifetime outside the registers and shared memory of the previous block, often in global memory.
3. The boundary waits for the whole grid, even if some workers finish earlier than others.
4. Opportunities to overlap the end of one operation with the beginning of another must be represented explicitly, or are lost.

The CUDA Programming Guide describes this execution and occupancy model in detail. It is worth being precise about what the boundary does not imply. It does not mean that every tensor is copied to the CPU, nor that all GPU memory is discarded. Model weights, activations and session state can remain in device memory across launches. The costs are device-side scheduling, synchronization and any intermediate global-memory traffic that the chosen kernels require, in addition to host submission overhead.

Let us follow a toy fragment of a model: a matrix multiplication, followed by a bias, an activation and a normalization. A literal implementation may assign each operation its own kernel.

The figure separates four possible costs:

• submission cost: preparing and issuing work;
• boundary cost: ending one grid and beginning another;
• materialization cost: storing an intermediate result and loading it again;
• schedule cost: deciding when and where ready work should run.

Different CUDA techniques attack different rows of this list.

Real time changes the objective

Dollyglot’s video frames made the deadline tangible. For the calculation below, I will use the speech-model cadence of the WPK workload described later. The modality is different, but the systems problem is the same.

Inference systems are usually optimized for throughput: complete as many requests, tokens or samples as possible per second. Throughput still matters in a real-time system, but it sits inside a deadline.

Suppose a speech model consumes 80 milliseconds of new audio and must complete its next step before the following 80 milliseconds arrive. For one conversation, the recurring model update has a period P = 80 ms. With N conversations on one GPU, the work grows with N, but P remains fixed.

One model step is 80 ms. A 160 ms transaction contains two consecutive steps. Other continuous models may use a different period or freshness constraint; the argument does not depend on this particular number.

The useful capacity is therefore not “the largest batch that eventually completes.” It is the largest number of simultaneous sessions whose updates consistently finish on time. We can write that as

The percentile q and tolerated miss rate ε depend on the product. The exact notation is less important than the conclusion: a real-time capacity number needs both a latency distribution and a deadline. Average throughput is not enough.

There are three reasons for this.

First, input arrives on an external cadence. The engine cannot wait indefinitely to form a larger batch because the future audio frame does not exist yet. Waiting consumes the same time budget needed to process the current frame.

Second, state survives between updates. Attention caches, convolution histories, recurrent states, codec state and conversation state belong to a session and evolve each time the clock advances.

Third, tail latency determines usable density. If the median update takes 52 ms but the 99th percentile takes 87 ms, the system cannot safely promise an 80 ms cadence just because the average looks comfortable. One slow update also reduces the time available for the next.

This does not make batching useless. Resident sessions can form a natural cohort and share model-weight reads. The constraint is that the engine cannot create more headroom by waiting for unknown future arrivals. It has to make the current cohort efficient now.

Five useful techniques and their boundaries

At this point, the obvious question is whether CUDA already solves the problem. It gives us several mechanisms that reduce overhead, move data less often, expose concurrency or isolate resources. Before introducing a megakernel, I want to give each of them its strongest case.

These techniques compose. A runtime may use graphs for submission, fused kernels for local operation chains, multiple streams for overlap and Green Contexts for isolation. They are not rival answers to the same question.

The short version is: CUDA Graphs change how work is submitted. Fusion changes the boundary between compatible operations. Green Contexts change which resources may execute the work. Persistence changes how long the program lives. A megakernel changes what the GPU considers to be the program. The rest of this section makes those distinctions precise.

CUDA Graphs: amortize submission

Suppose our model update launches one hundred kernels and then repeats the same sequence. In ordinary stream-based execution, the host and driver prepare each launch separately. A CUDA Graph records the workflow, including kernel launches, memory copies, events and their dependencies, then replays it with a single graph launch.

This can be extremely effective when host launch overhead is significant. NVIDIA’s CUDA Graphs documentation explains that much of the setup is paid during instantiation, so repeated execution requires much less host work. Presenting the complete dependency graph may also give CUDA optimization opportunities that are unavailable when work arrives one operation at a time.

The catch is that a graph node which launches a kernel is still a kernel launch. A graph with one hundred kernel nodes does not become one kernel. The operations retain separate grids, resource requirements and dependency boundaries. CUDA schedules each node once its dependencies permit it.

This distinction is easy to miss because “graph” refers both to the model’s logical graph and CUDA’s submission object. Capturing a model into a CUDA Graph changes how efficiently the existing execution plan is replayed. It does not, by itself, replace that plan with a single GPU program.

CUDA Graphs can be the correct endpoint when the remaining kernels are large enough to use the GPU efficiently, their intermediate traffic is inherent or already fused, and host submission is the dominant avoidable cost. A megakernel is interesting when the boundaries inside the captured graph remain a material part of the latency.

Kernel fusion: remove local boundaries

Let us return to the matrix multiplication from our toy example. Its bias and activation can often become part of the same kernel. This is kernel fusion. Three launches become one, and two trips through global memory may disappear.

The fused kernel can apply the bias and activation while values are still in registers or shared memory. It removes launches and avoids materializing some intermediates in global memory. For many models, good fusion is one of the highest-value compiler optimizations available.

Fusion is constrained by more than graph connectivity. Two operations may want different tilings, different block shapes or different distributions of work across SMs. A consumer may require a global view of an intermediate tensor that no one block can produce alone. Combining bodies may increase register or shared-memory use enough to reduce occupancy. Branches and joins require coordination. Large code bodies put pressure on instruction caches and become harder to specialize.

This suggests an important distinction:

Fusion tries to erase a boundary by making adjacent operations one local computation. A megakernel can keep a boundary as an explicit phase inside one larger GPU program.

A megakernel does not need to keep every intermediate in registers, and it should not attempt to fuse operations whose requirements are incompatible. It can write a necessary intermediate to global memory, coordinate the producing and consuming worker teams, and continue without ending the physical kernel. Local fusion remains valuable inside the megakernel; it is not replaced by it.

Streams and dependent launches: expose overlap

What if two parts of the model do not depend on one another? CUDA streams allow independent kernels to execute concurrently when dependencies and resources permit. If one branch performs speech recognition while another prepares a generation step, putting them on separate streams may overlap useful work.

Hopper-class devices also support Programmatic Dependent Launch. A dependent kernel may begin its independent preamble before all work in its predecessor has completed, then synchronize before consuming the predecessor’s results. The mechanism can hide part of a boundary, including launch latency.

Concurrency cannot shorten a true dependency chain. Nor does declaring two kernels concurrent ensure that they will overlap: they may compete for the same SMs, registers, shared memory, bandwidth or work queues. NVIDIA describes dependent-launch overlap as opportunistic rather than guaranteed.

Streams therefore improve an execution plan by exposing concurrency to the runtime. A megakernel owns the schedule that maps operations and coordination onto workers inside one kernel. That extra control is useful only when the program has enough predictable structure to plan it well.

Green Contexts: partition resources

Green Contexts solve a different problem: interference. A Green Context is a lightweight CUDA execution context associated with selected GPU resources. It can partition SMs and work queues so that submitted work is constrained to its provisioned resources. A latency-sensitive workload might reserve one partition while another workload uses the remainder.

This addresses interference and quality of service, not the internal execution of a model. The kernels submitted into a Green Context do not need to change. Consequently, their launch boundaries, intermediate tensors and dependency graph do not change either.

The distinction is simple: a Green Context answers which SMs may run this work? A megakernel answers what program runs inside one physical kernel?

Resource partitioning can make latency more predictable, especially when unrelated workloads share a GPU. It can also leave capacity unused if a partition temporarily lacks ready work. The current Green Contexts documentation also cautions that disjoint SM partitions alone do not guarantee concurrent execution or forward progress, because other resources can still couple the workloads.

For a recurring model, a compiler may create worker teams within a megakernel and statically assign model phases to them. This resembles resource partitioning in spirit, but it operates at a different level: the teams participate in one coordinated program rather than hosting independent streams of kernel launches.

Persistent kernels: keep the program alive

The final mechanism changes time rather than submission, fusion or ownership. A persistent kernel launches a bounded set of blocks intended to remain resident. Instead of completing one finite grid and exiting, its workers wait for commands, epochs or queue entries, process them and wait again.

Persistence removes repeated launch and teardown from the recurring path. It may also preserve control state in registers or shared memory, although large model weights and session histories still live in device memory. Cooperative launch and grid-wide coordination make it possible to build persistent pipelines whose blocks synchronize safely, subject to residency constraints described in CUDA’s grid synchronization documentation.

Persistence defines lifetime, not scope. A persistent task dispatcher that selects small operations from a queue may still reproduce much of a framework runtime on the device. Conversely, a non-persistent megakernel can execute a complete model step in one launch and then exit. Continuous inference makes the combination natural: the multi-operation program repeats, so keeping it alive avoids rebuilding its execution context every time the external clock advances.

What the megakernel changes

This is the point where the name “megakernel” can be misleading. The interesting property is not that the CUDA file is very large. It is that control of the multi-operation program moves inside one grid.

In conventional execution, the host runtime and CUDA kernel boundaries form the outer control flow of the model. In a megakernel, that outer control flow becomes device code.

Instead of:

# The host owns the model's control flow.
for operation in model_graph:
    cuda.launch(operation.kernel)

the conceptual structure becomes:

# The host launches once.
cuda.launch(model_program)

# The GPU now owns the model's control flow.
for phase in compiled_schedule:
    wait_for(phase.dependencies)
    run(phase, workers=phase.team)
    signal(phase.consumers)

I have left most of the difficult work out of this pseudocode. The second form is not automatically faster. It replaces mechanisms supplied by CUDA, including kernel scheduling and completion boundaries, with mechanisms the megakernel must implement correctly. The program now needs answers to several questions:

• Which workers execute each operation?
• When may an operation begin?
• How do producers make their results visible to consumers?
• Which values stay local, which enter global memory, and when may storage be reused?
• How does the program avoid deadlock if workers wait on one another?
• How are errors, shutdown and observability handled while the grid remains resident?

A hand-written megakernel can answer these questions for one fixed model. A megakernel compiler tries to derive the answers from a model representation, a hardware target and the workload constraints.

Recent systems make different choices here. Hazy Research’s Llama-1B megakernel uses one forward-pass kernel to reduce gaps between memory-bound operations and pipeline weight loads across instruction boundaries. Mirage Persistent Kernel explores compiling tensor programs, including multi-GPU inference, into a persistent megakernel runtime. I read these systems as evidence for a broader design space, not as a rule that every model should become one kernel.

Why continuous inference is unusually well suited

The natural follow-up question is: if moving the program into one kernel is useful, why not do it for every model?

The answer is that megakernels trade generality for control. This is worthwhile when the workload gives the compiler enough structure in advance. Continuous inference happens to reveal a great deal of structure before the next input arrives.

The graph repeats

The tensor values change on every update, but most of the operation graph does not. A speech model repeatedly runs the same perception stack, state update, language backbone, generation stack and codec. The compiler can study this graph ahead of time, group compatible operations, identify branches and joins, and produce a schedule for the recurring update.

An unpredictable collection of unrelated requests gives a compiler less information. A stable repeated graph gives it more.

State is meant to remain resident

Continuous models carry histories across updates. Moving those histories between host and device every period would be prohibitively expensive, so a sensible engine already keeps them in GPU memory. A persistent megakernel extends residency from data to execution: the state stays in its planned device layout while the workers that transform it remain available for the next epoch.

This does not mean the whole model sits in registers or shared memory. Those memories are far too small. Weights and large histories reside in HBM and must still be read. The opportunity is to remove unnecessary traffic, reuse storage according to known lifetimes, and overlap transfers and computation where the dependency graph permits.

The deadline rewards predictability

Throughput systems can tolerate some variation if aggregate work remains high. A real-time engine is judged by its slow updates. Static worker ownership, bounded data-dependent paths and a stable memory plan reduce the number of runtime decisions that can perturb latency.

This is not a guarantee that a megakernel has a narrow tail. Memory contention, clocks, input-dependent work and host scheduling still exist. It is a structural opportunity: fewer dynamic decisions on the critical path give the system fewer ways to surprise itself.

The cohort is known

A continuous service usually knows the maximum number of sessions assigned to a replica. It can compile for that capacity, reserve separate state for every slot and execute active sessions together. Inactive slots can be masked or given neutral input so that session arrivals do not require changing the execution topology.

This fixed shape would be wasteful for a sparse, bursty workload. It is useful when predictable service capacity is the goal. The compiler can optimize not for “whatever batch happens to arrive,” but for a declared operating point: a model, a device, a session capacity, a context limit and a period.

A concrete design: one WPK update

The implementation I know best is WPK. I will use one update through its current H100 backend to make the previous sections concrete. This is an example of the design, not the only possible megakernel architecture.

WPK compiles a continuous model into a persistent program specialized for an H100 and a fixed session capacity.

The compiler begins with a tensor-level graph containing operations, dependencies, state reads and writes, and estimates of compute and memory cost. Lightweight straight-line work can be attached to the beginning or end of a larger parallel operation. We call the resulting unit a phase: at most one parallel anchor plus compatible local work.

The scheduler assigns every phase to a fixed CTA team. On an H100, the available team hierarchy includes all 132 worker CTAs, halves, quarters and individual workers. It then emits an ordered stream of phases for every worker. Dependencies within one team are already expressed by that order. Only dependencies crossing team boundaries need a device-visible completion signal.

The H100 used here has 132 SMs. WPK launches one worker CTA per SM so the complete cooperative grid can remain resident and synchronize safely.

At runtime, WPK launches those 132 worker CTAs cooperatively. Each worker reads its compiled stream and remains inside the kernel until shutdown.

The host has not disappeared. It still owns networking, admission, jitter buffering, the external clock, slot lifecycle and output delivery. What moves to the GPU is the operation-by-operation execution loop. During one update, the GPU performs input conversion, the compiled model program and output conversion before publishing completion.

Inside the kernel, each worker walks a static list. Before a phase begins, workers wait only for completion signals corresponding to dependencies from another team. They execute the phase, join a phase barrier, and the final participant publishes any signal needed elsewhere. There is no general-purpose device queue choosing among arbitrary ready operations.

This is why I think “one launch” is an incomplete description of a megakernel. The important artifact is the compiled program inside the launch: worker assignment, ordering, synchronization and the memory plan. Without those pieces, a large kernel is merely large.

What a megakernel does not solve

I do not want to turn one useful architecture into a universal rule. A megakernel removes some overheads by taking responsibility for work that CUDA previously handled. That creates new limits of its own.

It cannot remove the model’s work

For a model update requiring F floating-point operations, B bytes of unavoidable memory traffic and a dependency critical path L, a rough physical lower bound is

This is an optimistic bound: real kernels do not sustain every published peak simultaneously, and scheduling, barriers, instruction issue, cache effects and imbalance all add cost. A megakernel attacks part of the gap above the bound. It cannot make a bandwidth-bound model stop reading its weights and state, nor can it parallelize a true dependency.

It also still needs excellent operator implementations. A slow matrix multiplication inside a megakernel remains a slow matrix multiplication.

It makes resource management harder

One long-lived kernel must accommodate operations with different register, shared-memory and occupancy needs. Resources reserved for the kernel may limit how many blocks can reside on an SM. A design that keeps the entire GPU occupied can also make coexistence with unrelated workloads difficult.

This is one reason Green Contexts, resource partitioning and hybrid execution remain relevant. A system might reserve a partition for a latency-sensitive persistent program while using the remaining GPU for other work, or keep large irregular stages as conventional kernels. “Megakernel” need not mean “everything, at any cost.”

It prefers bounded control flow

Static schedules work best when shapes, paths and resource needs are predictable. Mixture-of-experts routing, variable-length loops, speculative branches and dynamic admission introduce choices that may not be known at compile time. They can be bounded, masked, compiled into several variants or handled by a device scheduler, but each option spends either work, memory or control overhead.

It increases compiler and verification burden

Kernel boundaries provide well-tested synchronization and memory-visibility semantics. Removing them means rebuilding the necessary coordination explicitly. A compiler must prevent races and deadlocks, validate storage reuse, and preserve numerical behavior through fusion and reordering.

Debugging is also less forgiving. A fault in a short kernel identifies one operation and releases the device. A fault in a persistent whole-model program may stop completion publication and leave the host waiting unless the runtime monitors the CUDA stream and records enough execution state to locate the failure.

It is hardware-specific

Schedules, team sizes, shared-memory budgets and fast operator implementations depend on the GPU architecture. The model representation may be portable, but the compiled artifact is not automatically so. A serious megakernel system therefore needs retargetable compiler machinery rather than one heroic CUDA file.

Choosing the right endpoint

The techniques in this article form a toolbox, not a maturity ladder. A smaller intervention is usually preferable when it removes the actual bottleneck.

Use CUDA Graphs when repeated host submission is the main bottleneck and the captured kernels are otherwise efficient. Use fusion when adjacent operations share a compatible mapping and intermediate traffic is avoidable. Use streams or dependent launches when independent work exists and resource usage permits overlap. Use Green Contexts when interference or resource isolation is the problem. Use a persistent kernel when the same device workers should repeatedly consume work without relaunching.

When I now look at a workload, I consider a megakernel only when several conditions coincide:

• the workload repeats a substantial, mostly stable operation graph;
• many individual kernels are short enough that boundaries matter;
• state remains on the GPU across updates;
• runtime scheduling decisions can be moved to compile time or bounded tightly;
• latency and tail predictability matter more than supporting arbitrary shapes;
• the performance benefit justifies a specialized compiler and verification effort.

Continuous models often satisfy all six. That is why the match is deeper than launch overhead. The application’s clock exposes a recurring program, persistent state gives the program continuity, and the deadline rewards control over its complete execution path.

The lesson I took from Dollyglot was not that real-time video required a one-off GPU trick. It was that the natural unit of a workload can be larger than the unit exposed by its software stack. Conventional inference asks CUDA to run the next operation. A megakernel gives the GPU a multi-operation program for the next update. WPK additionally keeps its workers alive across updates. For continuous inference, that update, rather than the individual tensor operation or the isolated request, is often the useful unit to optimize.

Further reading

• NVIDIA, Writing SIMT Kernels
• NVIDIA, CUDA Graphs
• NVIDIA, Programmatic Dependent Launch and Synchronization
• NVIDIA, Green Contexts
• NVIDIA, Cooperative Groups and grid synchronization
• Spector et al., Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B
• Spector et al., ThunderKittens: Simple, Fast, and Adorable AI Kernels
• Cheng et al., Mirage Persistent Kernel: A Compiler and Runtime for Mega-Kernelizing Tensor Programs