An inference request moves forward while GPU resources wait
Your server is running a trained model to answer a request: this is inference. It prepares the next token, a piece of text predicted by the model. It reads weights, the parameters learned during training, then performs matrix calculations and consults results retained for the context. If the model spans several GPUs, it also exchanges intermediate results between them. The request is progressing, but each stage uses a different mix of hardware resources.
A matrix multiplication may primarily need compute capacity while another operation mostly depends on memory bandwidth. A transfer may delay one calculation even though independent calculations are ready. The engineering task is to find that independent work and give it resources without compromising correctness. A megakernel combines multiple operations in one GPU program and controls their transitions: it can start a ready portion without waiting for an entire preceding operation to finish.
Flapping Airplanes’ funding brings this execution model into a wider discussion. The lab was co-founded by Ben Spector, one of ThunderKittens’ authors, together with Asher Spector and Aidan Smith. It works on both GPU execution and learning methods. This article examines its announcements against documented mechanisms and develops illustrative examples. LeCompute has not tested the company’s system.
What the October 2026 funding announcement establishes
On October 21, 2026, Index Ventures announced Flapping Airplanes’ $1 billion Series A. The company identifies Index and Lux Capital as co-leads of this financing round. Its initial funding, made public at the January 28 launch, amounted to $180 million. These are announced financing amounts, not a statement that all the proceeds will be spent on GPUs. Index’s October announcement, January 28 launch, company’s initial funding announcement.
NVIDIA appears among the supporters thanked in Flapping Airplanes’ Series A announcement. The statement neither quantifies an individual contribution from NVIDIA nor describes the terms of its support. The announcements reviewed concern financing for the lab; they do not announce NVIDIA acquiring ThunderKittens.
Bloomberg separately reported a $6 billion valuation, including the new financing. We attribute that figure to its October 21 reporting; it is not confirmed in the company and Index announcements we reviewed. A valuation describes investment economics, not revenue or a measured performance improvement. Rachel Metz’s original Bloomberg report.
The people involved also need careful attribution. The 2024 arXiv version of the ThunderKittens paper lists Ben Spector, Simran Arora, Aaryan Singhal, Daniel Fu and Christopher Ré. Asher Spector and Aidan Smith are not in that author list. Flapping Airplanes and ThunderKittens are linked through researchers and prior work, but the company should not be described as a new name for the entire academic project. Original ThunderKittens paper.
ThunderKittens starts with tiles, not whole applications
A kernel is a function whose code runs in parallel across many GPU threads, individual execution paths. To multiply large matrices, a program partitions the output into smaller blocks, called tiles. A group of threads loads the inputs for a tile, computes and writes the results. The tile layout determines which hardware operations the implementation can use effectively.
ThunderKittens provides tile primitives embedded in CUDA C++. The aim is to give kernel authors a more manageable vocabulary while preserving hardware control. Its public repository is MIT-licensed and describes a framework for NVIDIA GPUs. It operates within CUDA rather than offering an execution platform independent of NVIDIA’s instructions. Repository and license.
Tensor Cores Hardware units specialized in low-precision matrix multiplication, introduced by NVIDIA with Volta (2017). Each generation adds supported formats: FP16 → FP8 (Hopper) → FP6/FP4 (Blackwell). They run the bulk of the inference compute. , the units specialized for matrix operations, are only part of the path. The program must feed them, arrange memory and make results available when the next operation needs them. A concise primitive makes that work easier to express; its existence does not prove that a complete application uses the GPU more efficiently.
The program runs on SM Streaming Multiprocessor. An independent compute block of an NVIDIA GPU, holding its own execution units, registers, shared memory and Tensor Cores. An H100 SXM5 has 132 of them. SM occupancy is the key metric for how saturated a GPU is. , the multiprocessors that contain the GPU’s compute resources. Their registers hold working values for threads; shared memory lets threads in the same block exchange data. A block is a group of threads assigned to one SM, and its threads are organized into warps of 32. Some warps can prepare transfers while others compute or store results. A correct pipeline, in which these stages overlap, must track ready data and reusable storage. CUDA kernel programming model.
That distinction leads naturally toward megakernels. Even after each operation has been optimized, the transitions between operations still matter. Readers familiar with FlashAttention and matrix-throughput bottlenecks will recognize the pattern: speeding up one component makes previously smaller costs more visible.
Follow the request instead of one utilization number
Consider a request that has completed prefill The opening phase of LLM inference: prompt tokens are processed in parallel to build the context state. That reuse raises arithmetic intensity and can make the phase compute-bound; the exact regime depends on the model, batch, context and backend. , the stage that processes its prompt and builds the context state. During decode The autoregressive generation phase of an LLM: one token is produced at a time while rereading weights and the relevant context state. At low batch sizes, that traffic can make decode memory-bandwidth-bound. Its share of total cost still depends on input and output lengths, batching and cache reuse. , it generates tokens sequentially. Its resource requirements depend on the weights it must reread, the context it must consult and the other requests the server can batch with it.
HBM High Bandwidth Memory. Stacked DRAM integrated into the accelerator package through a very wide interface. It provides several TB/s, although the exact packaging topology depends on the product. is high-bandwidth memory packaged close to the GPU processor. It is different from thread registers and shared memory, which provide working space for smaller portions of an operation. Keeping a GPU program active does not imply that its matrix units are continuously computing or that memory is delivering its maximum bandwidth.
NVIDIA’s Nsight Systems profiler displays timelines of kernels and transfers. With GPU metric collection enabled and supported, it also exposes SM, Tensor Core and memory activity. Examining these views together helps distinguish useful execution from delays between stages. One average can hide an application alternating between very different bottlenecks. Nsight Systems analysis guide.
The practical question for our request is whether the GPU can load inputs for independent work while another result is being stored. Dependencies and available working space determine the answer. If both tasks need all the same memory bandwidth at once, scheduling them concurrently does not create additional bandwidth.
Kernel fusion, CUDA Graphs and megakernels change different things
Local kernel fusion combines operations in one kernel. An implementation might transform a value and consume it immediately, avoiding an intermediate array that a second kernel would read. Potential savings come from avoided traffic and launches. They must be weighed against the resources required by the larger kernel. NVIDIA’s explanation of kernel fusion.
CUDA Graphs packages tasks and dependencies into a reusable object the host can submit. Those tasks may remain separate kernels. NVIDIA also supports launching graphs from the device, enabling some scheduling decisions without a CPU round trip. Describing CUDA as requiring a new host decision before every operation would miss mechanisms already available in the platform. Device graph launch.
| Approach | What it reorganizes | What to measure |
|---|---|---|
| Specialized kernel | One operation or a small group | Does its execution dominate the request? |
| Local fusion | Operations sharing intermediate values | Which traffic disappears, and what resource pressure appears? |
| CUDA Graphs | Submission and task dependencies | How much delay comes from launching work? |
| Persistent megakernel | Internal execution of multiple operations | Does overlap outweigh the coordination it adds? |
PyTorch’s torch.compile compiler offers another route by transforming the computation regions it can analyze. A graph break occurs when it encounters code it cannot include in that representation. With default settings, it runs that portion in Python before resuming tracing. These interruptions can limit optimization, so a meaningful comparison must give existing frameworks a suitable configuration. PyTorch graph-break documentation.
A persistent program takes responsibility for progress
A CUDA stream is a queue of operations ordered by the program. In the conventional case, kernels in one stream execute in that order; different streams can already overlap work when dependencies and hardware resources allow it. A kernel boundary does not necessarily stop the entire GPU. CUDA also provides mechanisms such as Programmatic Dependent Launch to start preparing a dependent kernel before its predecessor finishes completely. CUDA asynchronous execution and overlap mechanisms.
A persistent megakernel remains active to execute successive portions of work. It must distribute them, check when their inputs are ready, assign resources and signal when their results become available. This is scheduling: deciding which work can advance and where it should run. The runtime, the software organizing execution, takes responsibility for dependencies previously separated by kernel launches.
Hazy Research’s public work uses an on-GPU interpreter: it reads a CPU-prepared instruction sequence and runs the corresponding portions of computation. These instructions describe model operations; they do not replace the GPU’s hardware instruction set. In the May 2025 prototype, one kernel covers a model forward pass, the computation that prepares the next prediction, rather than an entire conversation. Host-side work and further passes remain necessary to generate subsequent tokens. The public research code illustrates this approach without establishing that Flapping Airplanes’ current system has the same implementation.
Choose between predetermined work and a shared work queue
A scheduler can prepare a fixed instruction sequence for each unit or distribute ready portions through a common queue. The first approach avoids some runtime decisions. The second lets a unit that finishes early obtain more work. These are possible design choices whose value depends on the cost of a scheduling decision relative to the work being assigned.
In our example, a fixed assignment might leave one unit waiting while another still has several long portions to complete. A shared queue can rebalance the workload, but querying and updating it also has a cost. For very short portions, that cost can exceed the delay it removes. A useful profile should therefore account for scheduling decisions and the dependencies they actually release.
The CPU also retains service responsibilities, including receiving requests, preparing inputs and handling outputs. Moving part of scheduling onto the GPU still leaves an end-to-end path to measure. A fast instruction generator does not shorten a model-loading delay or a request queue before execution.
Prepare the next operation without overwriting the current one
Return to our request. A compute unit has finished a projection, a matrix transformation of the model’s values, but still needs to store its output. It might begin loading the next inputs before that store completes. However, the two stages may need the same shared-memory locations. The runtime must know when each location becomes reusable, not merely when an instruction started.
An illustrative failure makes the distinction clear. If a loader overwrites a buffer while a storer is still reading it, the computation may run quickly and produce incorrect results. Reserving space for the next operation avoids that conflict but reduces working space for the current one. Memory allocation and pipelining must be designed together.
Make dependencies visible without waiting for unrelated work
A completion signal is insufficient if the code reading the result has no guarantee of seeing the corresponding data. Writing the result, publishing the signal and the subsequent read need the correct ordering. CUDA’s memory model defines those guarantees and their scope. Cooperative Groups can synchronize threads within a group, subject to its size and launch requirements; it does not provide an arbitrary barrier across all GPUs. Cooperative Groups, CUDA C++ memory model.
Progress creates another constraint. Imagine every currently resident block waiting for a producer that has not obtained resources to run. The waiting work prevents the producer from resolving the wait. This example shows why persistent execution requires a resource and scheduling argument, not just an atomic counter.
The public benchmarks establish specific results
In May 2025, Hazy Research examined Llama-3.2-1B in BF16, a 16-bit floating-point format, with one sequence, a 32-token prompt and 128 generated tokens. The test used no speculative decoding, where an auxiliary mechanism proposes multiple tokens before verification. Its published chart shows decoding throughput of approximately 1.5 times SGLang’s on H100 and B200. The baselines in that study already used CUDA Graphs and PyTorch compilation. Low-latency study and Figure 1.
Their September 2025 Llama-70B megakernel used eight H100s and was integrated into Tokasaurus, an engine that batches requests and manages their memory state. With 65,536 prompts and completion lengths sampled from the ShareGPT conversation dataset, reported total throughput increased from 19,170 to 23,468 tokens/s relative to SGLang, approximately 22.4%. That total adds processed input tokens and generated output tokens; it is not the speed of one conversation. The test measures completion of the full request set. This is an author-reported 2025 result, not a benchmark of the platform announced in October 2026. Protocol and Table 4.
That workload adds a problem our single-GPU request did not face: a GPU may wait for activations, the model’s intermediate results, produced elsewhere. The runtime must coordinate communication and computation. The GPU interconnect still has physical limits; a scheduler tries to place useful independent work within some of those waits.
The status of the implementation matters as much as the chart. The research branch’s README explicitly describes the release as research reference rather than a supported product. That changes the effort required to integrate, validate and maintain it. README at the inspected commit.
ThunderKittens 2.0 addresses a different comparison: whether selected B200 general matrix multiplication, or GEMM, kernels can match or exceed cuBLAS, NVIDIA’s linear algebra library. Its authors emphasize matching inputs, cache behavior and thermal conditions. Winning on those matrix multiplications alone does not establish an end-to-end serving speedup. ThunderKittens 2.0 methodology.
What Flapping Airplanes now says its framework can do
The October announcement describes a megakernel framework capable of complete training runs, serving inference and playing video games, including all three at once in a single kernel. The company links this capability to overlapping memory-bound inference with compute-bound training. These are Flapping Airplanes’ stated capabilities. Company’s Series A announcement.
To assess them, separate the evidence needed. A reproducible test running several activities would establish their coexistence under its conditions. Showing operational value would also require comparing completed work, result quality, memory consumption, energy and delays imposed on priority tasks. The announcement alone does not provide that comparative protocol.
Suppose a training job yields some resources to interactive requests. The arrangement is useful if the machine completes more work in a fixed window while meeting response targets. Higher average throughput is much less attractive if some requests wait far longer. Evaluation must include the latency distribution rather than just its mean.
Resource complementarity also needs testing. Training reads and writes memory even when compute is its limiting resource. Inference still computes when it is mainly waiting on memory. Combining them can introduce contention for caches, shared memory or execution resources. Their isolated profiles suggest an opportunity; the combined execution profile establishes how much of that opportunity survives.
Why NVIDIA’s support can make strategic sense
From a hardware supplier’s perspective, software that makes a chip useful for more tasks can increase its value. An application previously too expensive to deploy may become viable if the same equipment completes more work. Flapping Airplanes’ investors argue for efficiency unlocking new uses. We treat that as an economic hypothesis, not a disclosed explanation of NVIDIA’s decision.
The argument does not automatically predict more hardware sales. With a fixed amount of work, greater efficiency may initially reduce a customer’s equipment requirements. If new applications absorb the released capacity, total demand can instead rise. Prices, demand and deployment constraints determine which effect dominates.
There is also a platform connection. ThunderKittens is embedded in CUDA, so better applications for these GPUs may strengthen the value of that ecosystem. NVIDIA’s announced support establishes neither exclusivity nor a commitment by the lab to use only its chips. Those would require contractual information absent from the reviewed announcements.
Compute efficiency and data efficiency need separate evidence
Flapping Airplanes’ launch communication emphasizes data efficiency: the examples and information needed to acquire a capability. Hardware efficiency concerns the time, energy and machines required to execute a program. A model can improve along one axis without improving along the other.
A method that needs fewer examples but more computation per example does not necessarily cost less to train. Conversely, a faster runtime can execute the same algorithm and dataset without changing how the model learns. A useful comparison first fixes the required quality, then measures the resources needed to reach it.
There is a plausible connection. A new architecture may be awkward to express efficiently in existing tools. Better GPU control could lower the cost of experimenting with it, allowing researchers to explore more alternatives. That is a chain of reasoning, not evidence that a public inference speedup proves a new learning algorithm.
Memory pressure, changing inputs and operations constrain the approach
A program combining more operations may need to retain more intermediate values. High register or shared-memory use can reduce the number of blocks active on an SM at once. Register shortages can also cause values to spill into local memory, which is more expensive to access. More fusion can eventually make a program slower. CUDA resource and occupancy guidance.
Serving inputs also change. A batch of short conversations, one long-context request and a prefill-heavy workload do not call for identical scheduling. A tensor’s shape means its dimensions, such as the number of requests and their sequence lengths. A compiler can specialize code for particular dimensions and need to recompile when they change. PyTorch also offers dynamic dimensions to avoid some recompilations. A specialized runtime must account for preparation, adaptation and fallback to a less optimized path. Dynamic shapes documentation.
An operator also needs cancellation, priorities, error recovery and observability. Moving scheduling decisions into a persistent program moves some of those responsibilities with them. For example, throughput alone cannot tell an engineer whether an interactive service remains responsive while another workload shares the device.
This distinction affects hardware decisions too. A new scheduler does not remove model weights or context state from the memory budget. Readers evaluating a deployment should keep the LLM VRAM calculation separate from a claim that the same memory can be used more productively.
Translate throughput gains without overstating savings
Consider an assumption independent of Flapping Airplanes’ results: a system delivers 1.22 times its baseline throughput on the same work. Assume an unchanged hourly machine price, full use of the extra capacity and a measured cost proportional to execution time. Cost per unit of work then becomes 1 / 1.22, approximately 0.820 of the original. That is an 18% reduction in unit cost for a 22% throughput increase. The denominators differ.
Now assume a sequential task in which the accelerated portion accounts for only 30% of the original processing time, with the other 70% unchanged. Relative time becomes 0.70 + 0.30 / 1.22, approximately 0.946: the task takes 5.4% less time. Inverting that duration gives a speedup factor of 1.057, a 5.7% increase in processing speed. This calculation assumes the portions do not overlap; it does not predict every kind of GPU optimization.
The exercise gives a useful way to read an announcement: identify the accelerated portion and the metric being counted. A faster matrix multiplication, a faster model pass and a service that responds faster to users are different levels of evidence. Moving between them requires measurement rather than carrying the same percentage across every level.
A benchmark that could test the broader promise
We would begin with a recent, appropriately configured baseline and record its versions and options. Model, precision, inputs and quality requirements would match. Compilation and preparation would be separated from cold start, steady-state execution and shape changes. This is our proposed evaluation protocol, not a completed experiment.
For inference, we would measure time to first token, inter-token latency and throughput at several concurrency levels. Long requests and bursts of arrivals would remain in the workload. The KV cache The stored key and value vectors an LLM has already computed for every token it has processed. It avoids recomputing attention over the whole history, at the cost of a memory footprint that grows with the context length. , the stored attention state, would be included in memory and concurrency accounting. Why KV-cache capacity matters.
For training, the target would be time and energy to a defined quality, including checkpointing, interruptions and recovery. For combined training and inference, both workloads would be reported separately before and after sharing the GPU. Readers could see which activity benefits and which pays for the arrangement.
An execution trace would then help explain the result through launch delays, memory activity, computation and communication. Resource metrics would be related to end-to-end time. A gain that disappears with different inputs or cache conditions would remain part of the published result, rather than being omitted from the headline.
What a developer can use today
Start by locating the wait in a representative request. If an existing library, local fusion, improved batching or CUDA Graphs addresses the measured problem, integration effort belongs in the decision. A megakernel becomes a serious option when finer execution control can deliver something the current stack does not obtain and the team can maintain its correctness.
ThunderKittens and the public megakernel implementations provide concrete material to study: code, methods and identifiable test configurations. Flapping Airplanes aims to extend that control into a research program involving models and learning. The next decisive evidence will be a reproducible comparison of its system, with workload, quality and operational conditions stated alongside the result.
Sources and method
Scope. Research cutoff: October 31, 2026. LeCompute ran no GPU benchmark and did not test Flapping Airplanes’ platform. We distinguish announced facts, author-reported measurements, documented mechanisms and our own illustrative assumptions.
Financing and people. Index Ventures, January 28, 2026 launch, Flapping Airplanes’ initial announcement, Index’s October 21 Series A announcement, company’s Series A statement and October 28 interview. These establish what was announced; they are not a technical audit. No individual contribution from NVIDIA is quantified there. The valuation is separately attributed to Rachel Metz’s original Bloomberg report, October 21, rather than presented as a company confirmation. This edition retains the source funding amounts in USD; the French edition applies the site’s conventional indicative conversion of 1 USD = 0.92 EUR, referenced to mid-2026 rather than the financing date.
Public research. ThunderKittens, arXiv:2410.20399, October 27, 2024 version, ThunderKittens repository and MIT license, May 27, 2025 Llama-1B study, September 28, 2025 Llama-70B study, Megakernels code, inspected commit 91eaff2, February 19, 2026 ThunderKittens 2.0 release. Results describe those studies and their historical baselines. For Llama-1B, we use “approximately 1.5 times” after inspecting the original chart, rather than the strict “over 1.5 times” stated in the post: its H100 bars do not clearly substantiate that threshold. For Llama-70B, Table 4 counts input and output; (23,468 / 19,170 - 1) × 100 yields approximately 22.4%. Neither experiment compares October 2026 engine versions.
Execution and observation. NVIDIA: kernel fusion, device graph launch, asynchronous execution, Cooperative Groups, memory model, CUDA best practices, Nsight Systems analysis. PyTorch: graph breaks and dynamic shapes. These explain their own systems, not the lab’s private implementation.
Assumptions and arithmetic. Buffer reuse, resident blocks waiting for an unscheduled producer, shared-workload scenarios and the 1.22-times-throughput examples are LeCompute illustrations. Relative unit cost of 0.820 assumes constant hourly pricing and full use of the additional capacity. For the sequential task, 1 - (0.70 + 0.30 / 1.22) gives 5.4% less time; 1 / (0.70 + 0.30 / 1.22) - 1 gives a 5.7% increase in processing speed. These calculations do not predict commercial savings or energy reductions.
Cover. AI-generated illustration created for LeCompute on November 1, 2026 with OpenAI’s built-in image generation tool. Its silhouette is inspired by the airplane on Flapping Airplanes’ official website; a screenshot was used as a visual reference. This is an independent editorial interpretation, not an official brand asset. The green lines evoke airflow rather than a measured physical simulation. The creation prompt is retained with the site’s images.
