Kernel Kombat: How Multi-Agent LLMs Squeeze 1.32× More From Your GPUs

GPU bills have a charming way of turning “just one more model deployment” into a finance meeting.

For companies running large language model serving stacks, the problem is rarely that nobody knows GPUs matter. Everyone knows. The harder problem is that performance bottlenecks often live inside kernels most executives will never see: attention merges, normalization fusions, activation multiplications, tiny pieces of code called millions or billions of times until “small inefficiency” becomes “why is the infrastructure budget wearing a crown?”

The paper behind Astra is interesting because it does not ask an LLM to invent a full CUDA implementation from a PyTorch sketch. That is the more theatrical version of the problem. Instead, it asks something closer to real production engineering: given existing CUDA kernels from SGLang, can a multi-agent LLM workflow optimize them while preserving correctness?¹

The reported answer is yes, within a narrow but useful scope: Astra improves three SGLang kernels by an average of 1.32× on NVIDIA H100 GPUs, using OpenAI o4-mini in a zero-shot prompting setup. A single-agent baseline, given the same tools and the same five optimization rounds, averages only 1.08×.

That headline is tidy. Too tidy, actually. The important story is not “LLMs are now CUDA experts.” Calm down, procurement. The important story is that GPU kernel optimization already has a natural workflow: test, profile, plan, code, repeat. Astra turns that workflow into four specialized agents and shows that, for several real serving kernels, the structure of the workflow matters as much as the model sitting inside it.

The real product is not code generation; it is the optimization loop

Astra is built around four agents:

That division looks almost too obvious until you compare it with the single-agent version. In the single-agent baseline, one model instance is responsible for everything: making tests, profiling, planning, and rewriting code. In a business slide, that sounds efficient. In actual systems work, it is how one weak step contaminates the rest of the loop.

The paper’s clearest example is Kernel 1, merge_attn_states_lse. Under the single-agent setup, it slows down to 0.73×, while the multi-agent version reaches 1.26×. The authors attribute the single-agent slowdown to unrepresentative test inputs generated during test construction, which biased profiling and led the optimization process in the wrong direction.

That is not a side note. It is the central lesson.

When an LLM optimizes code, the model is not just writing code. It is also deciding what evidence counts. If the test cases are unrepresentative, profiling becomes theatre. If profiling is misleading, planning becomes confident nonsense. If planning is wrong, code generation merely implements the error with brackets.

Astra’s mechanism is therefore less magical and more industrial: isolate the failure modes. Testing should not be an afterthought. Profiling should not be a vibe. Planning should not be buried inside one long prompt that also asks the model to produce valid CUDA.

The system is still simple. It runs for five optimization rounds. It uses the OpenAI Agents SDK. The model is o4-mini. The kernels are manually extracted from SGLang into standalone versions, optimized, then manually patched back and validated against the original framework implementation. This is not an end-to-end autonomous compiler. It is a structured experiment in turning feedback loops into agent roles.

That distinction matters because it makes the paper more useful, not less. Companies do not need a religion about multi-agent systems. They need a way to ask which parts of their engineering workflow deserve separate accountability.

The 1.32× result comes from old CUDA wisdom, not alien intelligence

The three target kernels are small but meaningful pieces of an LLM serving stack:

There is a quiet but important pattern here: the optimized kernels are longer. Lines of code rise from an average of 110 to 184, a 64% increase. Kernel 1 grows by 87%. Kernel 2 grows by 50%. Kernel 3 grows by 59%.

So the “optimization” is not simplification in the ordinary software-engineering sense. It is specialization. The code becomes more elaborate because high-performance CUDA often means expressing hardware-aware intent explicitly: fewer redundant expensive instructions, better memory access, less synchronization, more register-level work, and careful use of CUDA intrinsics.

The paper’s case studies are useful because they show what Astra actually changed.

Kernel 1 improves by moving expensive math out of the hot loop

In merge_attn_states_lse, the baseline repeatedly computes scalar mixing weights inside an inner loop over output dimensions. That means operations such as max, exponentials, normalization, and division are repeated for every element, even though the scalar scores do not change across the vector.

Astra’s optimized version hoists those loop-invariant computations outside the inner loop. The weights are computed once per output vector. The inner loop is reduced to memory loads, multiply-adds, and a store.

That is not glamorous. It is also exactly the kind of thing that matters.

The mechanism is straightforward:

1. Identify scalar values that do not depend on the loop index.
2. Compute them once before the loop.
3. Leave the loop with only per-element work.

This is classic compiler territory. The business relevance is that production kernels often contain hand-written compromises, inherited assumptions, or performance bugs that survive because nobody wants to reopen the CUDA cave. Astra shows that an LLM workflow can rediscover and apply this kind of transformation in a constrained setting.

The boundary is equally important: this works when the agent can correctly infer the invariant and preserve numerical behavior. If the transformation changes precision characteristics or silently assumes shape properties not covered by tests, the speedup becomes a liability wearing running shoes.

Kernel 2 improves by replacing shared-memory-only reduction with warp-level work

The second kernel, fused_add_rmsnorm, contains a block-level reduction that dominates runtime. The baseline uses a tree-based reduction in shared memory. That is already much better than a naïve global-memory reduction, but it still pays synchronization costs and progressively disables threads as the reduction proceeds.

Astra’s version first performs an intra-warp reduction using __shfl_down_sync. This keeps partial sums in registers within a warp. Only after that does the kernel use shared memory for the remaining inter-warp aggregation.

The difference is not cosmetic. Shared memory is fast, but registers are faster, and synchronization is not free. By doing more work inside the warp before involving shared memory, the optimized kernel reduces memory traffic and synchronization overhead.

This is the point where the multi-agent design begins to look less like a software architecture fad and more like a practical scaffold. The optimization requires recognizing the dominant runtime pattern, mapping it to a hardware-level tactic, changing the code, then validating the result. A single prompt can try to do all that, but the paper’s comparison suggests that separation improves reliability when the kernel is complex enough for test quality and profiling quality to matter.

The result for Kernel 2 is a 1.25× speedup. Not revolutionary. Very billable.

Kernel 3 gets the largest gain from vectorized FP16 access and fast math

The biggest speedup comes from silu_and_mul, which improves by 1.46×. The paper identifies two main mechanisms.

First, the optimized kernel uses vectorized memory access with __half2. Instead of loading individual FP16 values one at a time, it groups two contiguous half-precision values and loads them together. That reduces the number of memory operations and improves effective bandwidth.

Second, it replaces the baseline SiLU computation using standard expf and division with a faster CUDA-intrinsic version using __expf, __frcp_rn, and __fmul_rn. In plain English: use a fast exponential, replace division with reciprocal-plus-multiply, and keep the arithmetic pipeline busier.

Again, this is not evidence that the model invented a new numerical method. It is evidence that a structured LLM loop can apply known low-level tactics to a specific kernel and produce a correct improvement under the authors’ tests.

That “under the tests” phrase is not decorative caution. It is the hinge. Fast-math intrinsics can alter numerical behavior. The paper reports that the optimized kernels are correct under manually constructed validation against the original SGLang implementations, but production adoption would still need broader tolerance-aware testing, real traffic shapes, and service-level monitoring.

The experiment table is main evidence; the shape table is a boundary test

The paper includes several pieces of evidence, and they do different jobs. Mixing them together would make the result look either stronger or weaker than it really is.

The shape table is especially useful because it prevents a lazy reading of the average. Kernel 3 is highly stable, showing roughly 1.47× to 1.50× across the four reported shapes. Kernel 1 is more variable: 1.46×, 1.57×, 1.00×, and 1.14× across its four shapes. Kernel 2 also varies, from 1.33× down to 1.07×.

That variation matters for deployment. A 1.32× geometric-mean speedup does not mean every production request gets one-third faster. It means the tested kernels, across the tested shapes, improve on average. If your traffic mix leans heavily toward shapes where gains are small, your service-level improvement will be smaller. If the optimized kernels sit on a critical hot path, the system-level gain may still be worth chasing. If they do not, congratulations, you have optimized a footnote.

Astra is not presented as shape-specific autotuning. The authors explicitly contrast it with tensor compiler approaches that tune for particular shapes. That makes the results more interesting in one way and less complete in another. The system aims for generally useful kernel improvements, but production teams still need shape telemetry to decide where general improvement is enough and where specialized tuning is required.

The single-agent baseline fails in the most production-relevant way

The average comparison is simple: multi-agent Astra reaches 1.32×; single-agent reaches 1.08×. Both generate correct kernels in the reported tests.

But the average hides the more instructive detail:

This is exactly the pattern one would expect if agent specialization is most useful when task complexity creates coupled errors. For a simple kernel, one agent can manage. For a more complex kernel, test generation, profiling, planning, and code editing interfere with each other if handled in one overloaded loop.

The business translation is blunt: multi-agent orchestration is not automatically valuable because “more agents” sounds sophisticated. It is valuable when the workflow contains distinct forms of evidence that can corrupt one another. Kernel optimization has that property. So do many enterprise automation tasks, but not all of them. Four agents are not a strategy. Four separable failure modes might be.

What this directly shows, and what Cognaptus infers

The paper directly shows that Astra can optimize three manually extracted SGLang CUDA kernels on H100, using a five-round multi-agent loop powered by o4-mini, achieving 1.32× average speedup while passing manually constructed correctness tests against the original SGLang implementations.

It also directly shows that, in this setup, a single-agent version underperforms the multi-agent version on average, especially for the more complex kernel where weak test construction produced misleading profiling.

The paper further shows, through case analysis, that the improvements are explainable in CUDA terms: loop-invariant hoisting, warp-level reductions, vectorized half-precision loads, and fast-math intrinsics.

Cognaptus infers the following for business use: the near-term value of systems like Astra is not replacing compiler stacks or CUDA engineers. It is cheaper diagnosis and faster iteration on known hot kernels. In a serving organization with mature observability, a shape corpus, and a correctness harness, an Astra-like loop could become a practical assistant for performance engineering. It can propose candidate rewrites, test them, profile them, and hand engineers a narrower set of promising changes.

That is useful. It is not autonomous infrastructure magic. Thankfully. The cloud bill was already dramatic enough.

The production boundary is manual integration, not model intelligence

The strongest limitation is not hidden. Astra depends on manual pre-processing and post-processing.

Before optimization, the authors manually extract and simplify SGLang kernels into standalone versions. After optimization, they manually patch the generated kernels back into SGLang and validate them against the original framework implementation. The paper notes that these steps are non-trivial to automate because modern serving frameworks have complex internal dependencies.

This boundary changes how a company should interpret the result.

If you have a small number of high-value kernels, manual harness construction is reasonable. Build standalone tests. Capture production-like shapes. Run warm-ups. Use repeated timing. Validate against the framework. Canary the result. Monitor p95 and p99 latency after deployment. This is performance engineering, not astrology.

If you want to optimize hundreds of kernels across multiple frameworks and hardware targets, Astra is not yet a turnkey answer. The manual extraction and reintegration burden becomes the product. The paper’s future-work direction—more automation, possibly with human-in-the-loop guidance—is exactly where the next practical bottleneck sits.

There is also a hardware boundary. The experiments run on NVIDIA H100 GPUs. A rewrite that improves H100 behavior may not transfer cleanly to other GPUs, other CUDA versions, or different memory/computation balances. CUDA optimization is not a moral truth. It is a negotiation with a particular machine.

Finally, correctness remains empirical. The paper uses manually designed test cases and compares against original SGLang outputs. That is far stronger than trusting agent-generated tests alone. It is still finite validation. For production, especially with fast-math changes, teams need tolerance design, adversarial shape coverage, numerical drift monitoring, and rollback.

A practical adoption frame for GPU-heavy teams

For engineering leaders, the question is not whether to “adopt Astra.” The paper is a research prototype. The better question is whether your organization has the ingredients that made Astra useful.

The most sensible near-term use is a hybrid workflow. Let compilers, libraries, and framework defaults do the broad work. Then point agentic optimization loops at the small set of kernels that remain expensive, stable enough to test, and valuable enough to justify engineering attention.

That is not as exciting as “AI replaces GPU programmers.” It is also much more likely to survive contact with production.

The broader lesson: specialization beats heroic prompting when feedback is expensive

Astra’s contribution is not that an LLM can write CUDA. We already knew LLMs could produce code that occasionally compiles and sometimes even behaves. Truly, the bar is majestic.

The contribution is that kernel optimization benefits from decomposing the work into roles that match the real engineering loop. Testing, profiling, planning, and coding are different cognitive tasks. They consume different evidence. They fail differently. Astra’s results suggest that giving those tasks separate agents can improve outcomes, especially when complexity makes a single all-purpose agent vulnerable to cascading mistakes.

The 1.32× number is useful. The mechanism behind it is more useful. It tells us where agentic AI may actually earn its keep: not by hallucinating brilliance from a blank page, but by participating in disciplined loops where feedback is concrete, errors are measurable, and improvements can be validated.

For GPU-heavy businesses, that is the sober opportunity. Agentic systems may not replace the performance engineer. They may, however, make the performance engineer less lonely in the CUDA swamp. Frankly, that is already a decent start.

Cognaptus: Automate the Present, Incubate the Future.