Beating vLLM and Llama.cpp on Gemma4-12B

Optimizing Kernels for RTX GPUs with Emmy Compiler

By Dmitry Trifonov, Slawomir Strumecki, Ivan Oleynikov•July 27, 2026

BenchmarksLLMGPU Performance

AI has demonstrated incredible prowess in coding tasks, implementing entire applications from scratch. However, when it comes to CUDA kernels, a seemingly obvious target (small, self-contained functions computing specific operations), we still rely almost exclusively on handwritten code.

Given enough attempts and benchmarking feedback, AI can produce an individual kernel competitive with hand-tuned libraries: Stanford's AI-generated kernels matched or beat expert PyTorch baselines on several FP32 operations (though the same run reached only 52% of torch.matmul and 9% of PyTorch's SDPA in FP16) NVIDIA ran DeepSeek-R1 in a closed verifier loop to generate working attention kernels, and Kevin-32B trains that loop with reinforcement learning. However, it doesn't scale: a large model needs hundreds of kernel shapes, each tuned per GPU, and single attempts still miss the mark: KernelBench finds even frontier reasoning models beat the PyTorch baseline in fewer than 20% of tasks. The generate-and-benchmark loop can take weeks for a single model/card, and as of today it still requires a human in the loop; the result transfers to neither the next model nor the next card.

What we need is an AI-driven compiler stack that can generate kernels automatically from a high-level description of the model, and optimize them for specific shapes and hardware. This is where Emmy comes in. Below, we present our results on Gemma 4 12B, a small open-weight model, running on RTX 5090 where Emmy-generated kernels achieved up to 1.6× speedup over cuBLASLt, a 1.32× geometric-mean speedup (geomean is used to prevent outliers from skewing measurements) over PyTorch eager, which also uses cuBLASLt under the hood, across all 254 deployed kernel shapes, and end-to-end serving that now beats stock vLLM on time-to-first-token at every long-context point we measure: −16% single-stream, −2% and −8% at four and eight concurrent streams, −11% on RAG-style long-input serving, and −10% on a saturated 64-request wave.

We cannot translate the full kernel speed-up to end-to-end numbers because of the integration costs with vLLM that cannot be eliminated without significant modifications to the latter. Integration via vLLM plug-in also costs us 0.3-1ms per token (2-5% of the performance). Decode kernels are fully memory-bound, so our kernel optimizations are not applicable, thus the decode is overall slightly slower than stock vLLM.

Those results demonstrate that the compiler approach can generate kernels that are competitive with hand-tuned libraries on both: popular and niche operations, including the automatic derivation of the best GEMM and FlashAttention schedules for one of the most popular local LLM models, effectively demonstrating that Emmy gets the best of both worlds: it scales to new hardware + kernel pairs like a compiler, while also applying ad-hoc optimizations for a given kernel on a concrete GPU like hand-tuned kernels.

Compiler Optimizations

Emmy is a compiler, in the same lineage as TVM with its Ansor auto-scheduler, Hidet — which plugs in directly as a torch.compile backend, and Mirage, which superoptimizes a graph built through its own API. All of them generate kernels from a high-level description of the computation rather than from hand-written CUDA.

The compiler approach offers the following advantages compared to hand-tuned libraries:

• The same optimization that wins on a specific kernel can be applied to optimize all similar kernels.
• A large number of kernels can be generated and tuned automatically, finding the best schedule for each kernel without manual intervention.
• Fusing a large number of operations into a single kernel, which reduces memory traffic and launch overhead (not demonstrated in this article, reserved for future optimizations).

The main advantage of the compiler-based approach is also the source of its biggest challenge. A compiler is limited to generic, uniform treatment of different kernels and GPU models, while the kernels it produces compete against more flexible hand-tuned kernels. Hand-tuned kernels can apply ad-hoc optimizations for a given kernel on a concrete GPU, which is a very strong baseline for common operations like GEMM and FlashAttention.

The conventional compiler-generated kernels benefit from hyper-specialization and can be tuned for specific shapes, but they still fall short of the performance of hand-tuned libraries. This gap is documented: Bolt measured Ansor reaching only 20% of cuBLAS throughput on FP16 GEMM on a Tesla T4, and attributes the shortfall to auto-tuners' inefficient use of the tensor cores, the non-FP32, compute-bound operators that dominate LLM inference. Each of these systems attacks it differently: Bolt closes the gap by templating its search on CUTLASS, Hidet replaces declarative schedule primitives with a task-mapping programming model that exposes the fine-grained mappings vendor kernels rely on, and Mirage superoptimizes across the kernel, block and thread levels.

To overcome this challenge, we have extended the schedule move set to cover a large number of known human-engineered optimizations. Since the space of the possible schedules covers the space of hand-tuned kernels, the compiler is able to find kernels that are on-par or better than handwritten ones.

Here are a few examples of the optimizations that compiler can apply to generate competitive kernels:

• Transport: a per-thread cp.async producer loop (Ampere-era), cp.async.bulk.tensor TMA copy (Hopper/Blackwell), or a plain vectorized ld.global
• Pipelining: single-buffer, double-buffer, or an N-stage ring pipeline, where each slab's fill is scheduled relative to the compute that last reads it; pipeline stages can be stored in either gmem, smem, or registers
• Reduction: a combination of warp-shuffle, smem tree, cross-CTA atomics or a separate kernel to finalize partials
• Schedule: warp-specialized (dedicated producer/consumer warps driving the async copy engines) or a uniform SIMT schedule where every warp does the same work
• CTA order: the rasterization of the CTA grid over the output tile: row-major, column-major, or a swizzled / tiled order that keeps the working set resident in L2 for longer
• Register tiling: the per-thread or per-warp reg_m × reg_n, the K-inner unrolling
• Bank conflicts: shared-memory padding or swizzle, hardware swizzle via TMA (Hopper/Blackwell)
• Tensor core operation support

Through a combination of these moves, the compiler can generate kernels that are competitive with hand-tuned libraries on the specific shapes and hardware, including the automatic derivation of the best GEMM and FlashAttention schedules.

Below we describe the optimizations that led to the performance gains on the RTX 5090 and 4090 specifically on the Gemma4 12B model.

TMA Transport for Matmul and Flash Kernels (Blackwell)

We've realized that cuBLAS and Flash Attention kernels are still using the same cp.async transport on consumer Blackwell dies (in fact most cuBLAS GEMM kernels on consumer Blackwell, the FP16 tensorop path included, are forward-ported Ampere-era cutlass_80_* kernels), even though the hardware has a better option: the Tensor Memory Accelerator.

TMA wins on the 5090 because the copy engine does the work the threads would otherwise do. The actual load is a single cp.async.bulk.tensor that one elected thread fires and the hardware runs asynchronously against an mbarrier. Every other thread in the block issues no address arithmetic at all — contrast the cp.async path, where all 256 threads cooperatively compute per-element global addresses — and the TMA's swizzle drops shared-memory bank conflicts for free.

The down_proj (60.4 GFLOP in 286.2 µs = 211 TFLOPS) and the fused gate_up (120.8 GFLOP in 561.7 µs = 215 TFLOPS) are already running at the ~209 TFLOPS FP16/FP32-accumulate peak. There is no bubble left for a cheaper transport to fill. On the other hand, o_proj and q_proj sit at 94% and 91% of that roof, and taking the per-thread address arithmetic and st.shared traffic off the critical path recovers most of the remainder.

Hybrid FP16/FP32 Accumulation

The default matmul path on the production stack is using FP16 tensor cores with FP32 accumulation. However, on consumer dies (both Ada and Blackwell), the FP16-input/FP32-accumulate HMMA runs at exactly half the rate of FP16-input/FP16-accumulate (see NVIDIA's spec tables list 330.3 vs 165.2 peak tensor TFLOPS on the RTX 4090 and 419 vs 209.5 on the RTX 5090). EMMY_FAST_MATH=1 switches the compiler to the fast atom (mma_m16n8k16_f16_f16), which runs the mma chain at the 2X FP16 rate, but keeps accuracy in check by promoting the FP16 partials into the FP32 registers and thus doing global accumulation accurately in FP32. So you get the FP16 tensor-core speed with an FP32 accumulation instead of paying the FP32-accumulate tax on every single mma.

This is a known trick in the literature: DeepSeek's DeepGEMM uses the same periodic-promotion scheme to rescue the FP8 tensor cores' limited-precision accumulation on Hopper (DeepSeek-V3 technical report §3.3.2, promoting every N_C = 128 K-elements), building on CUTLASS's generic fp8_accumulation.hpp; Ootomo & Yokota, extending Markidis et al., recover full FP32 accuracy from FP16 tensor cores by splitting the inputs into hi/lo FP16 pairs and moving the accumulator update off the tensor core onto the FP32 CUDA cores; and Blanchard et al. give the rounding-error analysis of blocked mixed-precision FMA that this scheme is an instance of, in which the number of accumulator roundings drops from K to K/b. What production stacks ship, though, is the unprotected form: PyTorch's allow_fp16_accumulation, added in 2.7, does all FP16 GEMM accumulation in FP16 "at the cost of numerical precision", with the error growth in K shown in the table below. Getting the rate and the accuracy needs a kernel that promotes mid-loop — demonstrated by hand for a single 4090 shape, but not exposed by any mainstream library across the hundreds of matmul variants an LLM needs. The Emmy compiler generates all of them automatically, and does so in a way that is competitive with cuBLAS.

Here is how this looks in PTX:

// (1) Standard: FP16 inputs, FP32 accumulate — the accurate default: 4 f32 accumulator regs
asm("mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {%0,%1,%2,%3}, ... ;");

// (2) FAST_MATH: FP16 inputs, FP16 accumulate — ~2x the mma-chain rate: C/D are 2 packed f16x2 regs.
asm("mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 {%0,%1}, ... ;");

// (3) The safety net: every 64 K-elements, fold the f16x2 partials (h) into the f32 shadow (c), rezero.
__device__ void emmy_mma_promote_f16acc(float* c, unsigned* h) {
    for (int i = 0; i < 2; ++i) {
        float lo, hi;
        asm("{.reg .b16 lo, hi; mov.b32 {lo,hi}, %2;"   // unpack the two f16 lanes
            "cvt.f32.f16 %0, lo; cvt.f32.f16 %1, hi;}"
            : "=f"(lo), "=f"(hi) : "r"(h[i]));
        c[2*i] += lo; c[2*i+1] += hi;                   // add into the f32 shadow
        h[i] = 0u;                                      // reset the f16 accumulator
    }
}

The three generated pure PTX helpers are used in the kernel below. The emmy_ldmatrix_x4 function loads an A fragment out of shared memory, emmy_ldmatrix_x4_trans_pair loads two canonical-B fragments in a single ldmatrix, and emmy_mma_m16n8k16_f16_f16 is the full-rate FP16-accumulate atom whose partials emmy_mma_promote_f16acc (above) periodically folds into the FP32 accumulators:

static __device__ __forceinline__
void emmy_ldmatrix_x4(unsigned* r, const void* smem) {
    unsigned addr = __cvta_generic_to_shared(smem);
    asm volatile("ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%0, %1, %2, %3}, [%4];\n"
                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3]) : "r"(addr));
}

// x4.trans: two col-adjacent canonical-B fragments in one ldmatrix,
//   landing directly as b0[0..1] / b1[0..1].
static __device__ __forceinline__
void emmy_ldmatrix_x4_trans_pair(unsigned* b0, unsigned* b1, const void* smem) {
    unsigned addr = __cvta_generic_to_shared(smem);
    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 "
                 "{%0, %1, %2, %3}, [%4];\n"
                 : "=r"(b0[0]), "=r"(b0[1]), "=r"(b1[0]), "=r"(b1[1]) : "r"(addr));
}

// FP16-accumulate HMMA — c/d are 2 packed b32 regs
//   (the same element map as the four f32 regs, pair-packed).
static __device__ __forceinline__
void emmy_mma_m16n8k16_f16_f16(unsigned* d,
                               const unsigned* a,
                               const unsigned* b,
                               const unsigned* c) {
    asm volatile("mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 "
                 "{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\n"
                 : "=r"(d[0]), "=r"(d[1])
                 : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]),
                   "r"(b[0]), "r"(b[1]), "r"(c[0]), "r"(c[1]));
}

Here is the actual generated mlp_down kernel (K=15360) trimmed to the K-loop:

for (int ks = 0; ks < 7680; ks += 64) {  // K-slab loop (64-stride = promote cadence)
    // TMA: one thread fires the hardware bulk copy...
    if (elected_thread) {  
        mbarrier_arrive_expect_tx(&mbar[next], 49152);
        // ...of the next A slab, and the next B slab (every other thread idles)
        cp_async_bulk_tensor_2d(&a_smem[next], desc_a, ...);
        cp_async_bulk_tensor_3d(&b_smem[next], desc_b, ...);
    }
    mbarrier_wait_parity(&mbar[cur], phase);  // wait on this slab's descriptor copy

    // 4 atom-K steps inside the slab
    for (int ki = 0; ki < 64; ki += 16) {  
        emmy_ldmatrix_x4(a, &a_smem[swizzle(...)]);  // A/B fragments out of shared memory
        emmy_ldmatrix_x4_trans_pair(b, &b_smem[swizzle(...)]);
        // FP16-accumulate HMMA into the f16 partials `ch`
        emmy_mma_m16n8k16_f16_f16(ch, a, b, ch);
    }
    emmy_mma_promote_f16acc(c, ch);  // fold f16 partials -> f32 shadow `c`, then rezero
}
// `c` holds the FP32 result; a g2k split-K finalize sums the two partials in FP32.

The measured payoff across the Gemma projections (RTX 5090, seq_len 512, µs):

Two things stand out:

• Even in the accurate FP32-accumulate mode, Emmy matches or beats cuBLAS on every shape (1.00–1.26×). The compiler is competitive with NVIDIA's hand-tuned library on the honest path.
• FAST_MATH then adds 1.24–1.55× on top of Emmy's own FP32 lane (1.34–1.61× measured against cuBLAS, which also folds in the TMA win from the previous section), purely by trading the FP32-accumulate rate for FP16-accumulate + periodic promotion. That 1.24–1.55× is below the 2× mma-chain ratio because the epilogue and memory traffic don't speed up, but it is a large, free win on the compute-bound projections.

We measured the error of C = A@B with FP16 inputs drawn from N(0,1), comparing each accumulation strategy against an FP64 reference over the identical FP16-rounded operands (so this isolates accumulation error, not input rounding). Each strategy is modeled at the hardware's rounding cadence: one m16n8k16 HMMA computes a 16-element dot internally in full precision and rounds once into the accumulator, so the running sum rounds once per atom, not per element. The generator is scripts/bench_accum_error.py in the emmy repo. The number is relative L2 error, ‖C − ref‖ / ‖ref‖, swept over accumulation depth K:

emmy bench experiments/gemma-4-12B/accum_error --local

Three observations:

• FP32 accumulation is essentially exact: 9.2e-8 at K=256, drifting only to 8.1e-7 at K=32768, three orders of magnitude below the pure-FP16 path at every depth.
• Pure FP16, the naive fast path, starts an order of magnitude worse and grows with K: by down_proj's K=15360 it is at 4.5e-3 relative L2 (worst-element error hits ~3.7%), because every one of 15360 running-sum roundings piles up.
• Hybrid accumulation breaks that growth: because the FP16 error can never accumulate past one 64-element chunk before being folded into FP32, it sits at a flat ~3.3e-4 regardless of K, inside FP16's own representational noise.

Below you can also find a correctness evaluation via lm-eval on GSM8K, comparing the vllm, llama.cpp, Emmy FP32 accumulation, and hybrid accumulation. Since the hybrid accumulation error is within the FP16 representational noise, the results are identical to the FP32 accumulation. Emmy's scores land a touch higher than the other two, though within one standard error.

Flash Attention Optimization

The previous article, Learning FlashAttention the Hard Way, took a standard attention shape (8 heads, seq 4096, head_dim 64) and walked the compiler's generic moveset all the way to FlashAttention-2 parity — 206 µs against FA-2's 206 µs, with the autotuner's winning geometry (128 query rows, 4 warps, 2 m-atoms per warp) landing on Dao's hand-picked constant set (FA-3 and FA-4 need Hopper/Blackwell datacenter silicon and cannot run on consumer dies). Gemma 4's attention is a bigger shape (head_dim 256, 16 query heads, causal), but the compiler arrives at FA-2's exact kernel a second time, then pulls ahead by employing the hybrid FP16/FP32 accumulation trick on the P·V matmul.

Profiling the SDPA call at this shape shows a flash_fwd_splitkv_kernel<256, 64, 64, 4> followed by a flash_fwd_splitkv_combine_kernel: a 64-row query block, 64-key block, 4-warp FA-2 kernel, single-buffered, running the split-KV (flash-decoding) variant. The split is an occupancy patch: this shape launches only 128 CTAs against the 5090's 170 SMs, so FA-2 cuts each query block's K/V range across the idle SMs and merges the partial softmax results in the combine kernel.

FA-2 holds one K slab and one V slab in shared memory: no double buffer, no ring. Q, K and V at head_dim 256 already sum to 96 KB, filling the 5090's ~100 KB per-CTA budget, so a second live slab of either won't fit. It gets its copy/compute overlap not from a second buffer but from when it refills the one it has, staggered to each operand's kill point:

// one mainloop iteration over a key block — a single sK slab and a single sV slab:
copy(gV → sV);          // this block's V streams in, overlapping the Q·Kᵀ below
gemm(acc_s, Q, K);      // Q·Kᵀ reads sK — K is now dead
cp_async_wait();        //   ↳ V has landed
copy(gK_next → sK);     // NEXT block's K prefetched here, under the softmax + P·V (neither reads K)
softmax_rescale_o(...);
gemm_rs(acc_o, P, V);   // P·V reads sV — V is now dead

Next K streams in underneath the softmax and P·V (which no longer touch K); next V underneath the following step's Q·K. Emmy arrives at the identical schedule.

FA-2 runs both of its matmuls in FP32. Under FAST_MATH Emmy runs the P·V matmul on the FP16-accumulate atom: 2× the tensor-core rate on consumer dies with the periodic FP32 promotion from the hybrid accumulation section, while keeping Q·Kᵀ and the softmax statistics (the running max and denominator) in FP32.

The scoreboard, causal seq 512 × 16 heads × head_dim 256 (µs):

vLLM Integration

We ship Emmy as a plugin for vLLM. Emmy compiles the model's layers into its own kernels, while vLLM still does the scheduling, batching, memory, the API. Attention stays with vLLM, because of its paged-attention implementation and the lack of a plugin interface for the split-KV variant. Everything around it runs on Emmy's kernels: the projections, the normalizations, and the feed-forward blocks that hold most of the weights.

A convenience emmy serve command invokes vLLM and passes the necessary plug-in flags; anything it doesn't recognize is passed through to vllm serve:

# an OpenAI-compatible server on Emmy's kernels
emmy serve google/gemma-4-12B --generate

# the same server on stock vLLM kernels — the baseline in this article's benchmarks
emmy serve google/gemma-4-12B --generate --stock

# start, benchmark, print results, shut down — one shot
emmy serve google/gemma-4-12B --generate --bench

Benchmarks

Model & Serving Parameters

Hardware & OS

End-to-end numbers are RTX 5090 (32 GB) only: 12B of FP16 weights (~24 GB) do not fit the RTX 4090's 24 GB alongside a KV cache, but we also report per-kernel numbers for both the 5090/4090.

Engines

Emmy ships as a vLLM plugin, so the identical vLLM version and serve flags are used, the plugin only swaps the model's kernels for Emmy-generated ones.

It is worth noting that vLLM, SGLang and cuBLAS for dense models largely ship the same cuBLAS kernels, so the difference between them is primarily scheduling, batching, and memory management. This axis is orthogonal to the kernel optimizations in this article, so generally speaking, comparing Emmy to just vLLM is enough to illustrate the kernel wins. The llama.cpp lane is included as a sanity check to confirm that vLLM baseline is not underperforming.

The reason we chose vLLM for integration is primarily because it is written in Python and has a plugin system, which allowed us to integrate Emmy without forking or patching the codebase. However, because we're not replacing the entire serving stack, the gains from kernel optimization haven't been fully transferred to the end-to-end benchmarks. vLLM still owns attention kernels and the integration glue, so the end-to-end gains are lower than the per-kernel gains.

Workloads

Each point in the tables below answers a different "who is this for" question, and together they cover the three regimes the serving literature and the local-LLM community actually benchmark — single-stream latency, batch throughput scaling, and long-input prefill:

Reproducing these Results

Every benchmark in this article lives as a self-contained experiment in the emmy repo under experiments/gemma-4-12B/, runnable with one command on a pre-allocated card or on a cloud VM. To run on a remote machine use --ssh user@address instead of --local. The experiments are self-contained: they download the model, build the kernels, and run the benchmarks.

# the three vLLM lanes (stock / emmy / emmy-fastmath) on the article's five workload points
# replace --local with --ssh user@address for remote run
emmy bench experiments/gemma-4-12B/serving_rtx5090 --local

# the llama.cpp lane (source build + F16 GGUF conversion, same client and flags)
emmy bench experiments/gemma-4-12B/serving_llamacpp_rtx5090 --local

# per-kernel golden-set benchmarks, both precision lanes
emmy bench experiments/gemma-4-12B/kernels_rtx5090 --local
EMMY_FAST_MATH=1 emmy bench experiments/gemma-4-12B/kernels_rtx5090 --local
emmy bench experiments/gemma-4-12B/kernels_rtx4090 --local
EMMY_FAST_MATH=1 emmy bench experiments/gemma-4-12B/kernels_rtx4090 --local

# the accumulation-error sweep behind the numerics table above
emmy bench experiments/gemma-4-12B/accum_error --local

For a fast, tokenless emmy boot there is also a prebuilt serving image with the compiled kernels, the model snapshot, and the execution plans baked in:

docker run --gpus all --ipc=host -p 8000:8000 cloudriftai/vllm-emmy-gemma4

RTX 5090

End-to-end

Despite up to 1.6× per-kernel speedups, the wins don't translate completely to end-to-end serving, and the residual is plugin-boundary cost. The plugin runs inside vLLM's layer loop, so every one of the 48 decoder layers is cut at the attention seam (Emmy computes q/k/v, hands them to vLLM's attention, and resumes with the output projection), and Emmy's norm/stat/scale glue runs as separate kernels where torch.compile fuses them: ~1100 launches per decode step against stock's ~650, worth the 0.3–1.0 ms/token. On a step that mixes a prefill chunk with decode tokens the plugin does run its chunk program and its decode program as two passes over disjoint rows where stock composes one fused varlen batch. Removing the seam and the glue means owning attention and batch composition too, which is a fork of the serving stack rather than a plugin.

Output token throughput (tok/s):

# the three vLLM lanes
emmy bench experiments/gemma-4-12B/serving_rtx5090 --local
# the llama.cpp lane   
emmy bench experiments/gemma-4-12B/serving_llamacpp_rtx5090 --local

¹ llama.cpp's 64-slot configuration OOMs: it allocates the full ~15 GB KV cache up front for 64×640-token slots. Unlike vLLM's paged allocator, it doesn't exploit Gemma's sliding-window layers here, and that doesn't fit alongside 23 GB of FP16 weights.

We have also tested long sequences (16K, 24K, 32K, 64K), but omitted them for brevity: the results follow the 4K/4K benchmark point.

Median latencies, TTFT / TPOT (ms), from the same runs:

emmy bench experiments/gemma-4-12B/serving_rtx5090 --local
emmy bench experiments/gemma-4-12B/serving_llamacpp_rtx5090 --local

² The c=64 TTFT is measured as a single 64-request wave (--num-prompts 64): every request's prefill is admitted in one wave, so the median measures prefill under saturation. (At --num-prompts 256 the median lands on second- and third-wave requests and measures queue drain — TPOT×256 plus admission — not prefill; the per-engine mean/median inversions make that visible.) The TPOT half of the row is from the --num-prompts 256 steady-state run.

The correctness check: 200 GSM8K questions, few-shot prompts, same seed (lm-eval 0.4.12, strict exact-match):

All four configurations land within one standard error of each other: FAST_MATH's hybrid accumulation does not degrade task quality, matching the kernel-level error analysis (its ~3.3×10⁻⁴ relative error sits inside FP16's own representational noise).

Speculative Decoding (MTP) Smoke-test

Gemma 4 ships an official speculative-decoding companion model (google/gemma-4-12B-it-assistant): a 4-layer drafter that shares the target's embedding and KV cache and proposes several tokens per step, which the target then verifies in one pass. vLLM supports it out of the box. In this section, we verify that Emmy running Gemma 4 with MTP and does not exhibit any surprising performance or quality degradation. The table below shows throughput (tok/s) of Emmy running Gemma 4 with MTP on inputs of random tokens. The assistant model is executed with stock vLLM means, i.e. Emmy optimizations for Gemma4 4 do not cover the assistant model.

emmy bench experiments/gemma-4-12B/serving_mtp_rtx5090 --local

(Plain-number cells are from the uniform re-measurement — all lanes under the main tables' serving settings, GPU memory utilization 0.96 and Emmy's per-workload knobs, via the emmy bench command above. The todo cells still show earlier mixed-provenance values and are being replaced as the same run finishes.)

We reran the identical GSM8K evaluation with the MTP drafter enabled as correctness check (google/gemma-4-12B-it-assistant, 3 speculated tokens per step).

emmy bench experiments/gemma-4-12B/gsm8k_mtp_rtx5090 --local

Disclaimer. This section only serves a smoke-test showing that MTP works, and Emmy's performance gains are similar. The absolute performance numbers are not directly comparable. The performance of the speculative decoding (compared to direct Gemma 4 runs above) depends on the problem and the domain-specific prompt you give it. One prompt may lead to an output with a lot of redundancy that the assistant model can frequently predict, while another one may be out of assistant's reach.

Per Kernel

The chart below covers the full tuning-record catalog: every Gemma 4 12B kernel shape deployed at every serving width, including decode buckets 1/8/32/64, prefill chunks 256/2048/4096, and symbolic variants. Kernels are measured against PyTorch eager and torch.compile, which delegates matmuls to cuBLAS, attention to FA-2, and uses Inductor for other kernels. Each shape carries three normalized bars: the default FP32-accumulate Emmy deploy pick, the same deploy under EMMY_FAST_MATH=1, and torch.compile.

emmy bench experiments/gemma-4-12B/kernels_rtx5090 --local
EMMY_FAST_MATH=1 emmy bench experiments/gemma-4-12B/kernels_rtx5090 --local

The 254 cases were measured live at -O3 on 2026-07-27 in both precision lanes, at the same revision as the serving numbers above — the catalog now includes the 2048-token chunk-quantum tier's shapes. (Four additional catalog-only fused-form cases fail the harness's fp16 accuracy check in isolation — the m64/m2048 fused norm→linear megakernels — and carry no Emmy bar; serving deploys their cut spellings, which are measured above as the bare matmul rows.)

The worst case in both lanes is the documented research-class residual rather than a tuning gap: attention.hd512, the 8 global layers' single-KV-head attention at head_dim 512 — a flash-codegen shape (the hd256 sliding-window attention, by contrast, beats torch SDPA on both cards). The standard lane's next floor is the wide M≥2048 gate/up projections: consumer dies run FP16-input/FP32-accumulate HMMA at half rate, so no FP32-accumulate schedule can reach cuBLAS's FP16-accumulate numbers there — the FAST_MATH lane is the fix (the same shapes run at 1.2–1.6× in it). Everything else at or below 0.7× is a launch-bound sub-30-µs norm fusion where torch.compile's single fused kernel wins a microbenchmark that the serving graph never sees in isolation.

RTX 4090

Gemma 4 12B in FP16 (~24 GB of weights) does not fit in the RTX 4090's 24 GB alongside a KV cache, so we only report per-kernel benchmarks for the 4090. The RTX 4090 (Ada, sm_89) also predates TMA, so the transport optimization from the Blackwell section doesn't apply here: all kernels use the cp.async transport.

Per Kernel

emmy bench experiments/gemma-4-12B/kernels_rtx4090 --local
EMMY_FAST_MATH=1 emmy bench experiments/gemma-4-12B/kernels_rtx4090 --local

The 139 cases were measured on 2026-07-27 and cover all of the 4090's recorded golden shapes; the 4090's tuning-record mirror of the later golden rounds (the m8/m2048 tiers above) is still pending, so these numbers are a floor, not the card's final word.

Conclusion & Future Work

The optimizations described above have allowed us to achieve 2-16% TTFT improvement on the Gemma4 12B on RTX 5090 thanks to the combination of compiler optimizations, hybrid accumulation strategies, and transport mechanisms (TMA). All kernels used for serving were generated fully automatically by Emmy, demonstrating that compiler-generated kernels can match or exceed hand-tuned libraries in performance.

The biggest limitations is the lack of quantization and sparsity support in the compiler, and the small number of benchmarks across models and GPUs. Since the compiler search is machine-learning-based, extending the GPU fleet and ensuring that we can generate optimal kernels for a wide range of models requires collecting a significant amount of performance data on each new GPU. This is an active area of research.

Additionally, integration with vLLM via a plugin costs 0.3-1 ms per token (1-5% of decode performance). A better integration with vLLM or other serving stacks could further improve end-to-end performance.

References

Tensor Compilers and Schedule-based Code Generation

1. Jonathan Ragan-Kelley, Connelly Barnes, Andrew Adams, Sylvain Paris, Frédo Durand, Saman Amarasinghe. Halide: A Language and Compiler for Optimizing Parallelism, Locality, and Recomputation in Image Processing Pipelines. PLDI 2013. The algorithm/schedule split the compiler's Loop IR / Tile IR boundary mirrors.
2. Tianqi Chen et al. TVM: An Automated End-to-End Optimizing Compiler for Deep Learning. OSDI 2018.
3. Lianmin Zheng et al. Ansor: Generating High-Performance Tensor Programs for Deep Learning. OSDI 2020. Auto-scheduling over a hierarchical search space, measured against vendor libraries.
4. Siyuan Feng et al. TensorIR: An Abstraction for Automatic Tensorized Program Optimization. ASPLOS 2023. The abstraction underpinning TVM's MetaSchedule.
5. Philippe Tillet, H. T. Kung, David Cox. Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations. MAPL 2019.
6. Yaoyao Ding et al. Hidet: Task-Mapping Programming Paradigm for Deep Learning Tensor Programs. ASPLOS 2023.
7. Lei Wang et al. TileLang: A Composable Tiled Programming Model for AI Systems. 2025.
8. Benjamin F. Spector, Simran Arora, Aaryan Singhal, Arjun Parthasarathy, Daniel Y. Fu, Christopher Ré. ThunderKittens: Simple, Fast, and Adorable Kernels. ICLR 2025.
9. Yuka Ikarashi, Gilbert Louis Bernstein, Alex Reinking, Hasan Genc, Jonathan Ragan-Kelley. Exocompilation for Productive Programming of Hardware Accelerators. PLDI 2022.
10. NVIDIA. CUTLASS: CUDA Templates for Linear Algebra Subroutines and Solvers (docs, CuTe layout algebra). The hand-written reference for the CTA-tiling, pipelining, and TMA patterns the compiler emits.
11. Jason Ansel et al. PyTorch 2: Faster Machine Learning Through Dynamic Python Bytecode Transformation and Graph Compilation. ASPLOS 2024. The torch.compile / TorchInductor stack this post benchmarks against — Dynamo graph capture plus an Inductor backend that fuses pointwise/reduction ops and, in its default mode, lowers large GEMMs to ATen/cuBLAS with no epilogue fusion. Under mode="max-autotune" Inductor can generate the GEMMs itself (Triton templates with epilogue fusion, plus CUTLASS/CuteDSL backends); we benchmark the default lane.

Compiler-generated Kernels vs. Hand-tuned Libraries

1. Jiarong Xing et al. Bolt: Bridging the Gap between Auto-tuners and Hardware-native Performance. MLSys 2022. Directly argues auto-tuners trail hand-optimized cuBLAS/cuDNN and closes the gap with CUTLASS-templated search; measures Ansor at 20% of cuBLAS on FP16 GEMM (Tesla T4), attributing the shortfall to inefficient tensor-core use.

LLM-generated Kernels

1. Anne Ouyang, Simon Guo, Simran Arora, Alex L. Zhang, William Hu, Christopher Ré, Azalia Mirhoseini. KernelBench: Can LLMs Write Efficient GPU Kernels?. ICML 2025 (PMLR 267:47356–47415). The standard benchmark for LLM kernel generation — out of the box, frontier reasoning models beat the PyTorch eager baseline in fewer than 20% of tasks; iterative refinement with execution and profiler feedback lifts that substantially.
2. Anne Ouyang, Azalia Mirhoseini, Percy Liang (Stanford CRFM). Surprisingly Fast AI-Generated Kernels We Didn't Mean to Publish (Yet). 2025. A search loop over LLM-generated kernels matching or beating expert PyTorch baselines on several FP32 operations on an L40S — while reaching only 52% of torch.matmul and 9% of PyTorch SDPA in FP16, which the post attributes to FP32 being the less-optimized baseline.
3. NVIDIA (Terry Chen, Bing Xu, Kirthi Devleker). Automating GPU Kernel Generation with DeepSeek-R1 and Inference-Time Scaling. 2025. A closed generate-verify loop (~15 min/problem on H100) producing numerically correct kernels for 100% of KernelBench Level-1 and 96% of Level-2 problems, and attention kernels 1.1–2.1× faster than PyTorch Flex Attention.
4. Carlo Baronio, Pietro Marsella, Ben Pan, Simon Guo, Silas Alberti. Kevin: Multi-Turn RL for Generating CUDA Kernels. 2025. Reinforcement learning over the multi-turn generate-and-benchmark loop.

FlashAttention

1. Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, Christopher Ré. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. NeurIPS 2022.
2. Tri Dao. FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning. ICLR 2024. The kernel our attention schedule ties on consumer dies.
3. Jay Shah, Ganesh Bikshandi, Ying Zhang, Vijay Thakkar, Pradeep Ramani, Tri Dao. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision. NeurIPS 2024. The Hopper (sm_90a) rung — WGMMA plus TMA warp specialization, neither of which consumer dies implement.
4. Ted Zadouri, Markus Hoehnerbach, Jay Shah, Timmy Liu, Vijay Thakkar, Tri Dao. FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling. 2026. The datacenter-Blackwell rung, built on tcgen05.mma and TMEM — sm_100 instructions that consumer Blackwell does not have.
5. Tri Dao, Daniel Haziza, Francisco Massa, Grigory Sizov. Flash-Decoding for long-context inference. 2023. The split-KV variant profiled here: FlashAttention per K/V chunk, one extra log-sum-exp scalar per row per split, then a log-sum-exp-rescaled merge in the combine kernel.
6. Dao-AILab. flash-attention. The source the tile constants, kernel-traits template arguments, split heuristic and mainloop listing in this section are read from — none of that appears in the papers above.

Megakernels and Superoptimization

1. Mengdi Wu et al. Mirage: A Multi-Level Superoptimizer for Tensor Programs. OSDI 2025. Superoptimization over μGraphs at the kernel, block and thread levels — the search-space ancestor of Emmy's moveset. (The whole-model megakernel is the same group's later MPK system, not this paper.)
2. Benjamin Spector, Jordan Juravsky, Stuart Sul, Owen Dugan, Dylan Lim, Dan Fu, Simran Arora, Christopher Ré. Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B. Hazy Research, 2025. The whole-model megakernel approach this post deliberately stops short of.

GPU Architecture, Transport, and Warp Specialization

1. NVIDIA. NVIDIA Hopper Architecture In-Depth. 2022. Introduces the Tensor Memory Accelerator. Note this is a datacenter part — for TMA on consumer Blackwell the normative sources are the Programming Guide and PTX ISA below, not any architecture document.
2. NVIDIA. NVIDIA RTX Blackwell GPU Architecture, v1.1. 2025. The GB202 die this post measures: 170 SMs, 1792 GB/s, and the Appendix A tensor-core table (419 vs 209.5 peak FP16 TFLOPS with FP16 vs FP32 accumulate).
3. NVIDIA. NVIDIA Ada GPU Architecture, v2.02. 2022. The AD102 counterpart: 330.3 vs 165.2 peak FP16 TFLOPS — the same exact 2:1 FP16-vs-FP32-accumulate ratio on the RTX 4090.
4. NVIDIA. CUDA Programming Guide — Asynchronous Data Copies. API-level reference for TMA (tensor maps, mbarrier completion) and the older cp.async family; §4.11.2.2.5 is the source for TMA swizzle removing shared-memory bank conflicts.
5. NVIDIA. PTX ISA v9.3 — ldmatrix, mma, cp.async.bulk.tensor, elect.sync. The instruction-level reference; cp.async.bulk.tensor requires sm_90 or higher, which is why sm_120 has base TMA but not Hopper's multicast or cluster features.
6. Michael Bauer, Henry Cook, Brucek Khailany. CudaDMA: Optimizing GPU Memory Bandwidth via Warp Specialization. SC 2011. The origin of the dedicated producer/consumer warp split in the moveset above.

GEMM Work Decomposition

1. Muhammad Osama, Duane Merrill, Cris Cecka, Michael Garland, John D. Owens. Stream-K: Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU. Poster, PPoPP 2023 (the arXiv version is the full treatment). The work-centric generalization of the split-K reduction the projections use: an even partition of the aggregate MAC iterations across CTAs, with a fix-up reduction for output tiles split across them. Attention's split-KV combine is a different operation — a log-sum-exp-rescaled merge of partial softmax states, not a plain sum.

Mixed-precision Numerics

1. Massimiliano Fasi, Nicholas J. Higham, Mantas Mikaitis, Srikara Pranesh. Numerical Behavior of NVIDIA Tensor Cores. PeerJ Computer Science 7:e330, 2021. Probes the tensor cores' actual arithmetic on V100/T4/A100: the internal accumulation is binary32 with round-toward-zero, partial products are held exactly, and the FP16-output mode rounds the accumulator down to binary16 at every step — the hardware behaviour the hybrid path works around.
2. Pierre Blanchard, Nicholas J. Higham, Florent Lopez, Théo Mary, Srikara Pranesh. Mixed Precision Block Fused Multiply-Add: Error Analysis and Application to GPU Tensor Cores. SIAM J. Sci. Comput. 42(3):C124–C141, 2020. Rounding-error analysis of blocked mixed-precision FMA: with block size b the error constant becomes γ_{K/b} rather than γ_K — the framework this scheme is an instance of. Their measured TC16-vs-TC32 comparison on a V100 is the gap between the FP16-accum and FP32-accum columns above.
3. DeepSeek-AI. DeepSeek-V3 Technical Report. 2024. §3.3.2 documents that Hopper's FP8 MMA accumulates at only ~14 mantissa bits, and mitigates it by promoting the tensor-core partials into FP32 registers on the CUDA cores every N_C = 128 K-elements — the same periodic-promotion structure as the hybrid path, applied to a hardware-truncated accumulator rather than a genuine FP16 one.
4. DeepSeek-AI. DeepGEMM (README as of the original release). The kernels shipping the above; DeepSeek in turn credit CUTLASS's fp8_accumulation.hpp, the generic "promote at an MMA interval, then zero" utility.
5. Hiroyuki Ootomo, Rio Yokota. Recovering Single Precision Accuracy from Tensor Cores while Surpassing the FP32 Theoretical Peak Performance. IJHPCA 36(4):475–491, 2022, doi:10.1177/10943420221090256. Recovers FP32 accuracy from FP16 tensor cores by splitting each input into hi/lo FP16 pairs (extending Markidis et al.) and moving the accumulator update off the tensor core onto the FP32 CUDA cores to escape its round-toward-zero — a different fix for a different problem (input precision, not accumulation depth), sharing only the "get the accumulate out of the mma" move.
6. Stefano Markidis, Steven Wei Der Chien, Erwin Laure, Ivy Bo Peng, Jeffrey S. Vetter. NVIDIA Tensor Core Programmability, Performance & Precision. AsHES @ IPDPSW 2018. The origin of the split/compensation line the above extends.
7. Sam Patterson. Improving FP16/16 matmul accuracy with two-stage accumulation. 2025. The nearest prior art to the hybrid path: the same FP16-accumulate mma with out-of-mma FP32 promotion, hand-written for a single shape on an RTX 4090.

Transformer Architecture: Normalization, Activation, Attention

1. Biao Zhang, Rico Sennrich. Root Mean Square Layer Normalization. NeurIPS 2019. The RMSNorm whose per-row statistic Emmy folds into the matmul prologue. Gemma 4 applies it both pre- and post-sub-layer, plus QK-norm, so a decoder layer carries more norms than a naive reading suggests.
2. Dan Hendrycks, Kevin Gimpel. Gaussian Error Linear Units (GELUs). 2016. The activation — in its tanh approximation — inside Gemma's GeGLU epilogue.
3. Noam Shazeer. GLU Variants Improve Transformer. 2020. The gate/up GLU structure (GeGLU here, not the Llama family's SwiGLU) behind the fused gate_up kernel — two accumulator channels sharing one A operand, hence N = 30720 = 2 × 15360.
4. Joshua Ainslie, James Lee-Thorp, Michiel de Jong, Yury Zemlyanskiy, Federico Lebrón, Sumit Sanghai. GQA: Training Generalized Multi-Query Transformer Models from Multi-Head Checkpoints. EMNLP 2023. Why the sliding layers' KV projections are half the width of q_proj (16 query / 8 KV heads at head_dim 256). The 8 global layers go further still — a single KV head at head_dim 512, with keys reused as values, which is why the kernel list has a k_proj_global and no v_proj_global.

Serving Engines and Benchmark Methodology

1. Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph E. Gonzalez, Hao Zhang, Ion Stoica. Efficient Memory Management for Large Language Model Serving with PagedAttention. SOSP 2023. The vLLM serving stack Emmy's generated kernels plug into.
2. ggml-org. llama.cpp. The third contender in the tables below, run against the F16 GGUF conversion of the same checkpoint.
3. Amey Agrawal et al. Taming Throughput-Latency Tradeoff in LLM Inference with Sarathi-Serve. OSDI 2024. Chunked prefill, listed in the contenders table.
4. vLLM. vllm bench serve CLI reference. The single external client driving every engine here — --ignore-eos, --seed, --num-prompts, --percentile-metrics.
5. MLCommons. MLPerf Inference Rules. The formal codification of the benchmark-hygiene list this section's setup follows.

Speculative Decoding

1. Yaniv Leviathan, Matan Kalman, Yossi Matias. Fast Inference from Transformers via Speculative Decoding. ICML 2023. The mechanism behind the MTP experiment.
2. Fabian Gloeckle, Badr Youbi Idrissi, Baptiste Rozière, David Lopez-Paz, Gabriel Synnaeve. Better & Faster Large Language Models via Multi-token Prediction. ICML 2024. The multi-token-prediction drafter family Gemma 4's assistant checkpoint belongs to.

Gemma

1. Gemma Team, Google DeepMind. Gemma 4 Technical Report. 2026. The model this post optimizes (google/gemma-4-12B; MTP drafter google/gemma-4-12B-it-assistant).
2. Gemma Team, Google DeepMind. Gemma 3 Technical Report. 2025. The source Gemma 4 defers to for its 5:1 local-to-global attention interleaving and 1024-token local span — the sliding / global layer split this post's kernel names refer to.

Companion CloudRift Writeups

1. Dmitry Trifonov. Learning FlashAttention the Hard Way — Part 2. The generic moveset walking one attention kernel (8 heads, seq 4096, head_dim 64) to FlashAttention-2 parity on the RTX 5090 — the result this post builds on.
2. Dmitry Trifonov. Modern GPU Matmul Optimization. Register tiling, swizzle, pipelining, TMA, warp specialization and split-K measured one move at a time — the same schedule moves, on the kernel that taught them.
3. Dmitry Trifonov. Surfacing a 60% performance bug in cuBLAS. The hand-tuned TMA SGEMM walk-through, the batched-FP32 dispatcher bug, and the cutlass_80_* Ampere forward-ports that most of cuBLAS's Blackwell compute paths — the FP16 tensorop HGEMM included — still run.
4. Dmitry Trifonov. A Principled ML Compiler Stack in 5,000 Lines of Python — Part 1: Torch IR → Loop IR, Part 2: Tile IR, scheduling loops onto a GPU, Part 3: autotuning, an SP-MCTS search over Tile-IR rewrites. The IR pipeline and the search loop behind the compiler.