Making an existing vLLM FlashAttention-3 forward kernel faster is one of the harder jobs in GPU operator optimization: beat parity on 26 configurations without breaking correctness. One paper ran a controlled comparison. On one side sat an autonomous agent built exactly along the prevailing recipe, loaded with every tool that recipe prescribes. On the other sat a minimal harness steered by questions from a human. The tool-laden agent stalled at 0.48x. The question-driven one reached 1.13x.

The fully equipped version: a whole pipeline for the agent

The equipped side is called Arsenal, a baseline the authors built themselves along the line of evolutionary agentic search (AVO, KDA). Every step fed the model a bundle: one of 16 optimizations selected by the NCU bottleneck label, a parsed performance diagnosis, instruction-mix deltas from disassembly, BM25-retrieved documentation sections from the PTX ISA and CUDA guides, hints mined from reference libraries like DeepGEMM, ThunderKittens and CUTLASS (1.6 GB on disk), 15 seeded knowledge notes, causal memory of recent steps, plus sanitizer runs and forbidden-symbol checks. The median prompt was 22.6K characters, 37% of it retrieved or pre-fetched material.

The outcome: 133 model-driven steps, 660 dollars of API cost, and 44.7 hours of a 62.4-hour run before it reached its best result. Best geometric mean was 0.48x (the abstract rounds this to 0.50x), the worst configuration sat at 0.24x, 0 of the 26 configurations reached parity, and the goal was never met. The champion hit 0.48x at step 35 and then stopped improving for the remaining 98 steps. Thirty-seven branches were opened along the way, 36 of them from that same step, and 3 human rollbacks discarded 117 steps.

The dividing line: 82% versus 26%

The paper lands on two numbers. For the equipped agent, 37% of the prompt was pre-fetched reference material, yet 82% of the tool-result text it read was spent re-reading its own kernel source. The question-driven agent had nothing pushed to it and spent 26% on its own kernels.

Tool usage actually went the other way. The question-driven side ran the profiler 429 times on raw output, did 372 SASS/PTX inspections, made 377 builds against 937 benchmark runs, and spawned 48 subagents. The equipped side ran one parsed profiler dump per step, did 92 documentation lookups, and spawned 0 subagents. Across 16 sessions there were 11,532 tool calls (8,791 in parent sessions, 2,741 in subagents), roughly 71 tool calls per human message.

The paper's own observation is blunt: the questions did not reduce tool use. They aimed it.

What those 17 questions look like

Of the 162 human messages, 58 were directives, 44 questions, 30 approvals, 19 corrections and 11 goals; 63 were phrased interrogatively. The questions fall into eight categories: measurement check 13, bound differencing 10, benchmark validity 7, gap triage 7, agent process 6, challenge a limit 5, strategy options 5 and hypothesis check 3.

The paper spotlights four cases, because the shape of the question matters more than the count:

  • "What configs are bringing down the geomean?" The agent believed tuning the hd128 paths was the priority. A per-configuration table showed hd96, hd192 and hd256 sitting on the scalar fallback at 0.01x. With the fallback removed, the geomean went from 0.44x to 0.96x in four hours.
  • "Is this ceiling a property of the hardware, or of our design? Where does the time go outside the kernel?" The agent had declared split-KV decode a "genuine ceiling within the harness rules". A host-side profile found a per-call tensor fill of 2-5 us outside the kernel. Removing it took split-KV from 0.87x to 0.97x and the geomean from 1.02x to 1.13x.
  • "Do our scenarios cover the shapes and features the operator is actually called with?" The agent had treated 40 of 40 correct scenarios as done. Porting the operator's full test grid showed 52 of 58 correct, with sink and left-pad outputs off by 495x and 103x. After six fixes, 58 of 58 were correct and the geomean on the new grid was 1.10x.
  • "Why do you think the goal was achieved? We are still well below parity." The Stop-hook judge had already accepted "goal met". The agent conceded it had stopped after one effective lever and declared victory prematurely. fp8 hd256 decode went from 0.20 and 0.34 to 0.66 and 0.77.

None of these questions carried an answer. The human wrote no kernel code. Of the 63 interrogative messages, only 17 mentioned an attention concept and none prescribed a CUDA construct. The agent wrote all 13.8K lines of the two kernels.

What actually shipped

fa3_hopper.cu (3.6K lines) and fa4_blackwell.cu (10.2K lines) were released under Apache 2.0, together with the evaluator, the PTX probes used to validate each Blackwell primitive in isolation, the question bank and every trace. fa3_hopper.cu is correct on all 182 Hopper scenarios, runs at 1.16x FA3 and 1.10x the fastest available implementation including cuDNN, with 135 of 182 at or above parity. fa4_blackwell.cu is correct on all 154 runnable Blackwell scenarios, runs at 1.46x FA4 and 1.07x the fastest available implementation, with 98 of 124 at or above parity and a worst case of 0.78x; on production shapes from SOL-ExecBench it reaches 1.10x FA4.

The authors' own discounts deserve equal billing. Excluding the fp8-KV decode claim shapes moves the Blackwell all-set geomeans to 1.37, 1.32 and 1.08x. The edge sets consist of short kernels, so part of the 1.5x to 2x there is upstream dispatch overhead. On Hopper the losses are fp8 dense prefill, softcapped prefill, head dimension 96 and long sliding windows, at 0.34x to 0.79x. On Blackwell, fp8 inputs and dense head dimensions of 192 and above are not built yet, which is 28 of the 182 scenarios. Neither kernel covers the backward pass.

How to read the comparison

The paper states its limitations plainly. Each arm is a single run, with different front-ends (the equipped agent ran Claude Code headless; the question-driven side mixed Claude Code and Cursor), and a steered arm is being compared against an unsteered one. KernelSeek had already built a kernel for FA3's 40-scenario test grid in the two preceding days, so its 0.009x starting point understates its head start. And of its 68.0 hours, 26.7 were an idle gap.

More importantly, the paper names the experiment it has not run: whether the question bank can replace the human. The design is fixed at four arms (the equipped loop, the bare agent with harness alone, the bare agent with a questioner drawing on the bank, and a human questioner) with the model and budget held constant. It is written down, not executed. The 72 bank entries mined from public sources have not yet steered an agent either.

And do not read the result as "tool-rich pipelines do not work." Other papers do get results out of them on other tasks: AVO reports 7 days of continuous autonomous evolution on B200 attention, outperforming cuDNN by up to 3.5% and FlashAttention-4 by up to 10.5%; Cursor and NVIDIA report a 38% geomean speedup across 235 Blackwell kernels in 3 weeks with a multi-agent system. The tasks and evaluators differ, so these are not directly comparable. What they do show is that this paper measured one specific failure, not a general verdict.

So what

If your agent is stuck on a performance target, resist the urge to add another retrieval source or knowledge store. First check what share of what it reads is its own output: for the equipped agent it was 82%. Then ask whether there is a question that could move it to a layer it has not examined, or force it to defend a limit it has already declared. The paper's conclusion is direct: in deep kernel optimization the scarce resource is not information but direction.

All data and the evaluator are at arXiv:2611.03517; the multi-agent system referenced above comes from the public Cursor and NVIDIA report.