Bohan Hou

After our kernel-agent experiments, we had to confront an uncomfortable question: Do we still need DSLs and ML compilers when agents can do it all？ In short: We're still here! 😎 If you're curious why, here's our take: tvm.apache.org/2026/11/03/re… We've also rebuilt TIRx from the ground up as a shared DSL and compiler foundation for agentic kernel development.

2.94× KDA fwd, 6.84× KDA bwd, 2.59× MSA prefill, 3.99× MSA decode, 1.33× KDA decode, 1.71× MLA, 1.68× VSA — just by pressing Enter. After years of working on ML compilers and designing many abstractions, the most useful compiler thing in TIRx I’ve ever built might just be a PTX tablegen.

BTW, I noticed that the CuTe DSL 4.7 PR includes a static analysis tool for catching certain synchronization bugs. I expect we’ll see many more tools like this going forward. :)

Very interesting post. One optimistic thought: Models will keep getting better—GPT-6, 7, 8, whatever comes next. But if humans still need to review every GPU kernel they generate, the real bottleneck remains. I’ve seen PRs add a single PTX instruction like fence.proxy.async to fix bugs that only appear after hours or days on hundreds of GPUs. For example, FlashKDA relased last week: github.com/MoonshotAI/FlashK… I already run 3–5 agents in parallel. The slowest part is still me. I still have to read the code. So how far are we from a truly human-free GPU kernel development loop? And if we still need tooling around generated code—even at the PTX level—then abstractions haven’t disappeared. They’ve just moved from the programming interface into the verification infrastructure.

Don’t feel bad about causing infra folks pain—NVIDIA does that every 2–3 years anyway :)

Every time I fight ptxas over why it put some variable in R instead of UR, or why it lifted a CSE this time but not that time, I think back to someone telling me years ago that truly great programmers can write plain C and mentally compile it all the way down to the final asm.

“A compiler optimization that always works is just part of the programming model.” Many of you have probably watched @cHHillee 's Jane Street talk: janestreet.com/tech-talks/bu… The talk was given in 2025, but the question had already been on my mind since the summer of 2023, when I was interning on Meta’s PyTorch Compiler team, where I worked with Horace. At the time, my TensorIR paper had just been accepted to ASPLOS 2023. TensorIR gave us a schedule-oriented IR for tensorized program optimization, and much of my thinking then was shaped by Ampere Tensor Core kernels. But the target was already moving. Hopper made warp-specialized kernels, TMA-driven pipelines, new synchronization patterns, and FP8 much more central. Kernels like FlashAttention-2 were also pushing beyond what our existing schedule abstraction could express comfortably. Meanwhile, in the MLC project — a compiler-based ML serving engine targeting both servers and mobile devices — TensorIR’s schedule system felt awkward when we tried to reach peak performance on quantized GeMV/GEMM kernels. During a Topgolf team event, Horace asked me a question: “How much time do you think you need to upgrade your compiler to Hopper? And what about the generation after Hopper?” That question has stayed with me ever since. Every time I thought about what my next project should be — auto warp specialization, auto pipelining, or another compiler pass in that direction — I found it hard to fully convince myself. Not because those directions are unimportant, but because they still did not answer the question that bothered me most. If every new hardware generation requires another major compiler redesign before users can even express the best program, then maybe that is simply not the right path. For a while, I honestly believed the answer might just be: write CUDA C++, and at most build libraries on top so that we do not have to rewrite everything from scratch. Things become more interesting now that agents are starting to become capable of writing kernels with specific low-level instruction patterns. One might think stronger agents will make this problem go away: they can write more code, search more schedules, patch more compiler passes, and try more variants than humans can. Personally, I do not have a decisive answer. It is risky to make strong predictions about what agents will or will not be able to do. From my own experience, they are already reasonably useful for tuning some kernels in CUDA C++ and performing local rewrites that make the generated PTX/SASS look better. But when asked to design or orchestrate new warp-specialization plans, pipelined data movement, or synchronization protocols, they can still struggle to produce a kernel that is even correct — let alone one that does not hang. One experiment I tried was asking Codex to port DeepSeek’s MegaMoe kernel into TIRx. I expected the task to be relatively straightforward, since TIRx intentionally keeps the source structure close to hardware-native CUDA C++. In practice, the agent spent several days and still could not resolve a hang without my intervention. What eventually helped was not simply asking it to “try again,” but changing the feedback loop: I asked a reviewer agent to produce an audit file, line by line, comparing the generated kernel against the original and justifying that the two were doing the same thing. That workflow can help when an expert kernel already exists. But it is much less satisfying if the goal is to have agents design new kernels from scratch. For those curious, here is the experimental agent-generated TIRx port of MegaMoe (which I believe matches the performance of the initial release of DS’s kernels, but they have some optimizations later) that came out of this process: gist.github.com/spectrometer… It is not part of the initial release; it is mostly a record of the experiment. The result is not particularly polished, but it was useful for understanding what kind of feedback loop agents need when working on low-level kernels. Of course, a sufficiently high-level DSL can help when the abstraction already captures the kernel pattern. If the user can write GEMM(A, B) and the compiler deterministically lowers it to a good implementation, then life is easy. But that is exactly the hard part at the frontier: the pattern may not yet be mature enough to hide behind a stable abstraction. Until a strong enough model can solve this end-to-end, what we can do is improve the workflow around the model: give it more structured feedback and perhaps even lightweight verifiers that can check whether a generated kernel shape is valid with respect to synchronization. This connects back to the opening sentence: “a compiler optimization that always works is just part of the programming model.” For example, if a warp-specialization plan relies on mbarrier coordination, the compiler should at least be able to reason about whether that plan can introduce races, deadlocks, or violations of the intended producer-consumer protocol. That kind of feedback is hard to provide if the program is only raw CUDA C++ without a richer intermediate representation. Maybe the better question is not whether agents can write kernels, but what programming model gives humans, agents, and compilers enough shared structure to reason about kernels before the final benchmark.

before use/in use/after use

Making @AMD @amdradeon GPUs competitive for LLM inference! 130 toks/s of Llama 2 7B, 75 toks/s for 13B with ROCm 5.6 + 7900 XTX + 4 bit quantization 80% performance of Nvidia RTX 4090 See how we do this in detail and try out our Python packages here: blog.mlc.ai/2023/08/09/Makin…

Now available in AppStore! apps.apple.com/us/app/mlc-ch…

#Llama2 is running on iPhone, iPad📱natively with GPU acceleration. No internet connection is required. See IOS instructions to get the test flight app now: mlc.ai/mlc-llm/docs/get_star…

Running LLM natively on your 🤖@Android phone, following our release of the iOS app. With MLC-LLM and TVM Unity, we are able to optimize and deploy the model in 1 week! 6~7 toks/sec on Galaxy S23. Demo: mlc.ai/mlc-llm/#android Check out for details: mlc.ai/blog/2023/05/08/bring…

Can LLMs run natively on your iPhone📱? Our answer is yes, and we can do more! We are introducing MLC-LLM, an open framework that brings language models (LLMs) directly into a broad class of platforms (CUDA, Vulkan, Metal) with GPU acceleration! Demo: mlc.ai/mlc-llm/

Excited to share our #ASPLOS23 paper. “TensorIR: An Abstraction for Automatic Tensorized Program Optimization”. Come to our talk in Section 5C at Grand D! TensorIR lies at the core of the next-level generation of the TVM compiler @ApacheTVM, from @SCSatCMU @OctoML.