GEAK v4: From Faster Kernels to Faster Workloads

Jul 23, 2026

Introduction

Kernel agents—LLM-driven systems that write, tune, and validate GPU kernels—are moving from research prototypes into practical performance engineering tools. Hand-writing a high-performance kernel for a new operator, dtype, model shape, or GPU generation still requires scarce expertise and repeated measurement on real hardware. A kernel agent turns this expert loop into an automated process: profile the target, form hypotheses, edit code, run correctness and performance tests, and retain what works.

GEAK, short for Generating Efficient AI-Centric Kernels, is the AMD agentic optimization framework for AMD Instinct™ GPUs. Given a kernel or a live model-serving stack such as SGLang or vLLM, GEAK searches for bottlenecks, generates and tunes implementations across backends such as Triton, TileLang, FlyDSL, HIP, and CK, and validates the result on the actual system.

GEAK v4 is a substantial redesign of GEAK v3. The earlier system was primarily a single-kernel optimizer. GEAK v4 keeps that capability but adds a system layer that optimizes whole-model serving throughput. Two design decisions make this possible. First, GEAK v4 uses Claude Code dynamic Workflows as a deterministic orchestration plane: JavaScript owns the control flow, including budget loops, parallel fan-out, verification, and stop conditions, while LLM agents are invoked only for structured technical judgment. Second, v4 adopts a fractal two-altitude architecture: the end-to-end workflow recursively calls the single-kernel workflow unchanged, allowing the system and kernel layers to evolve independently while preserving kernel-level rigor.

Key Takeaway

GEAK v4 upgrades GEAK from a single-kernel optimizer into an end-to-end GPU performance system. Instead of only asking whether one kernel can run faster, GEAK v4 asks whether the full SGLang/vLLM serving workload becomes faster.

The key design is a two-altitude workflow: the e2e layer profiles the real workload, ranks bottlenecks by Amdahl impact, and calls the kernel workflow only when a kernel can truly move model throughput. Deterministic JS Workflow controls the optimization process, while LLM agents focus on analysis, strategy, kernel authoring, and integration.

This redesign brings measurable performance gains across both levels: GEAK v4 improves single-kernel geomean speedup from 1.64x to 2.21x on Triton kernels and from 1.99x to 2.72x on HIP kernels, raises head-kernel geomean speedup to 1.33x, and achieves substantial end-to-end serving-throughput gains on real LLM workloads, including +60% on Qwen3.5-27B-FP8, +96% on Qwen3-14B-FP8, and +42.2% on Minimax-M3-MXFP8.

Why End-to-End Optimization Matters

The trap of local kernel wins

A faster kernel is not automatically a faster model. This observation is the central design lesson behind GEAK v4. A kernel can be measurably faster in isolation yet contributes almost nothing to model-level performance. If that kernel accounts for only a small fraction of GPU time, or if the dominant operator is already near the hardware limit, even a large local speedup may be invisible in end-to-end throughput.

The opposite is also true: the right kernel win can be transformative. When the optimized kernel sits on the workload's steady-state critical path, a single better implementation can lift whole-model throughput substantially. The difference is not only how fast the kernel is, but whether the workload actually spends time there. That answer requires profiling the real serving workload rather than relying on isolated microbenchmarks alone.

Amdahl-first triage

GEAK v4 makes Amdahl's law part of its control flow. Candidate work is ranked by a simple but operational priority score:

amdahl_priority = pct_gpu_time × achievable_speedup

A 5x improvement on a kernel that accounts for 2% of GPU time is less valuable than a 1.15x improvement on a GEMM that accounts for most of the decode path. GEAK v4 therefore spends its budget cheapest-first: configuration and backend sweeps before source changes, head-GEMM or attention bake-offs before broader editable-kernel loops, and recursive kernel authoring only for candidates that survive workload-level triage.

The workflow is designed to optimize where serving throughput actually lives. Isolated kernel speedup is treated as a hypothesis; end-to-end A/B measurement is the verdict.\

GEAK v3 versus GEAK v4

Dimension	GEAK v3	GEAK v4
Scope	Single-kernel optimization.	Kernel optimization plus whole-model end-to-end serving throughput.
Agent backend	mini-SWE agent.	Claude Code SDK/CLI with dynamic Workflows.
Orchestration	Deterministic Python orchestration.	Deterministic JavaScript Workflow controlling budget, fan-out, verification, and stop logic.
Architecture	Orchestrator plus generic workers.	Director -> Architect -> specialist roles, recursively wrapping the kernel optimizer.
Verification	Kernel benchmark.	Kernel benchmark plus end-to-end serving benchmark.
Knowledge base	RAG-style knowledge base and reverse knowledge.	perf_knowledge: structured kernel optimization knowledge across operators, backends, GPU generations, dtypes, and regimes.
Evolution	Knowledge carried through reverse-learning mechanisms.	Learned cards, workflow memory, and cross-run evidence for faster convergence.
Depth control	Max rounds, parallelism, and time limit.	Fast/default/deep modes, wall-clock limits, agent budget, and multi-kernel/multi-backend parallelism.
Optimization horizon	Short single-kernel runs, roughly hours.	Longer system runs, up to multi-hour or multi-day search depending on mode and budget.

The redesign shifts both scope and control. GEAK v3 is a strong single-kernel optimizer bounded by per-kernel benchmarks. GEAK v4 preserves that loop, but adds a system objective, a deterministic workflow control plane, specialist roles, recursive reuse of the kernel layer, and an end-to-end validation gate.

GEAK v4 Architecture

Kernel Workflow: a hierarchy of specialists

The single-kernel workflow, kernel_workflow, optimizes one AMD GPU kernel implemented in Triton, HIP, CK, FlyDSL, or another supported source path. Its design separates deterministic control from agent judgment. JavaScript owns the loop, fan-out, budget, verification, and state transitions; agents perform analysis, implementation, benchmarking, and structured technical reporting.

The pipeline is:

		Setup -> [Author] -> Analyze -> Benchmark -> Profile
      -> LOOP { Plan -> Optimize || Verify -> [Integrate] -> Commit -> Re-profile -> Update memory }
      -> Report -> Validate

Figure 1: GEAK v4 Kernel Workflow

The roles are deliberately narrow. The Director creates the isolated workspace and performs final independent validation from the true baseline. The TechLead owns strategy, choosing a small number of orthogonal directions per round and maintaining cross-round memory; it does not edit kernels or run benchmarks directly. Specialist Engineers explore algorithms, memory, compute, and host-runtime directions in separate workspaces. A deep-explore engineer can attempt broader roofline-targeted rewrites when incremental tuning is no longer enough.

Verification is the trust anchor. An engineer's self-reported speedup is never accepted directly. A separate Verify Engineer re-applies the patch in a clean workspace and re-measures it against the true baseline. Only patches that are correct, verified, and faster survive. The Integrator then attempts to combine compatible wins; the round winner must improve the cumulative best before it is committed as the new baseline for the next round.

This architecture encodes a simple rule: measurement is the only source of truth, and no single agent can fake a win. The oracle is immutable, source history is excluded from workspaces, build caches are isolated, GPU access is serialized, and failed or timed-out agents degrade safely rather than blocking the run.

Director: prepares the workspace, protects the true baseline, and performs final independent validation.
TechLead: plans 2-3 orthogonal directions per round, manages the insight blackboard, and decides when to stop or continue.
Specialist Engineers: explore algorithms, memory, compute, and host_runtime directions in parallel and in isolation.
Deep-explore Engineer: attempts broader rewrites that combine multiple levers when the expected ceiling justifies the cost.
Verify Engineer: re-measures claimed patches in a clean workspace against the true baseline.
Integrator: merges compatible verified patches and keeps only combinations that improve the best individual result.

E2E Workflow: Modules, Modes, and Learning

The end-to-end workflow, e2e_workflow, is the system layer. It owns the model server, benchmark harness, warm-server profile, Amdahl triage, cheap configuration sweeps, head-kernel search, editable-kernel milestone loop, reversible overlays, and final end-to-end validation.

The high-level flow is:

		Setup/Preflight -> Profile -> Strategize -> ConfigSweep -> HeadKernel -> Milestone loop -> Finalize -> Report -> Validate

Figure 2: GEAK v4 End-to-End Workflow

Every serving number—baseline, sweep result, integration candidate, and final validation—passes through a backend-agnostic benchmark harness. The harness owns the server lifecycle, serving-GPU mutex, untimed warmup, timed repeats, and median/spread reporting. A backend adapter, such as sglang.sh or vllm.sh, lets the same workflow drive different serving engines.

The modules divide responsibilities cleanly:

E2E Director: builds the isolated environment, runs preflight checks, records the true warm-server baseline, and performs final arbitration.
System Architect: performs Amdahl triage, routes profiled kernels into optimization tracks, plans milestones, and curates learned knowledge.
Profiler: traces a warm server under the target workload and produces a standardized Top-N profile table.
Config Tuner: sweeps service-level switches first, parity-checks each candidate, compounds accepted wins and triggers re-profiling when the bottleneck shifts.
Kernel Extractor: captures real shapes and reference I/O into an immutable standalone task for the kernel workflow.
Op Benchmarker: evaluates candidate backends for the highest-impact head operations.
E2E Integrator: reintegrates accepted changes as reversible overlays and runs the end-to-end gate with warm A/B, engagement proof, and output parity.

Three Optimization Modes

Mode	Scope	Search behavior	Best fit
fast	HeadKernel only.	Per-head extract-and-bake phases run in parallel under a time cap, then integrate serially on the serving slot.	Quick capture of the largest head-kernel opportunities.
default	ConfigSweep + HeadKernel + Milestone.	Serial, disciplined progression from cheap switches to kernel authoring, with one e2e gate per accepted change.	Standard full optimization.
deep	ConfigSweep + deep HeadKernel optimize + Milestone.	A cross-kernel x backend lane pool with state persistence, shared knowledge, patience control, reseeding, and convergence-based stopping.	Longer searches over many kernels, backends, and regimes

For each high-value kernel, the backend can change through three progressively more expensive routes: launch flags, environment variables, and source-level replacement. GEAK v4 does not simply choose the fastest default backend. It optimizes viable candidates, reintegrates them, and lets the end-to-end gate decide.

Recursion is the key mechanism. When e2e_workflow identifies a kernel worth authoring, it calls kernel_workflow as a nested workflow and passes an immutable extracted oracle, target language, shape and dtype information, regime requirements, and a budget. The kernel layer returns a verified patch. The system layer then overlays that patch into the live server and validates whether it actually moves serving throughput.

GEAK v4 learns at two levels. Run-scoped memory records directions, isolated speedups, e2e deltas, verdicts, and lessons, so the workflow avoids confirmed dead ends and keeps pressure on the current bottleneck. Persistent learned cards store validated principles keyed by kernel class, GPU architecture, and regime. These cards can add candidates to a future search, but they cannot prune the search or override measurement. On-box validation remains the judge.

perf_knowledge: a sourced kernel matrix and expert skills

Both workflows are supported by perf_knowledge, a sourced and machine-queryable knowledge base for AMD kernel optimization. It is designed as reference knowledge—facts, recipes, and candidate-generation guidance—not as an automatic decision maker. Its contract is conservative: it may help agents find better candidates faster, but it must never reduce a measured baseline or skip required validation.

The knowledge matrix is indexed across operators, backends, GPU generations, dtypes, and regimes. Operator coverage spans GEMM, attention, normalization, activation, positional encoding, sampling, reductions, data movement, quantization, MoE, and collectives. Backend coverage includes authoring languages such as Triton, FlyDSL, HIP, CK, ASM, and TileLang, as well as library or framework paths such as AITer, hipBLASLt, rocBLAS, MIOpen, PyTorch Inductor, RCCL, and related AMD paths.

On top of that matrix, expert_skills provides human-authored, scenario-specific optimization recipes. A skill carries a match selector, expected conditions, validation evidence, and measured outcomes. In the workflow, a skill is not treated as truth; it is treated as a high-quality candidate to reproduce first. It still enters through the same correctness, kernel-performance, and end-to-end gates.

Evaluation

The following measurements are from GEAK v4 evaluation runs on AMD Instinct MI300X-class systems using Opus 4.8 with single-stream parallelism unless otherwise specified. They should be read as workload-specific evidence rather than universal speedup guarantees.

Single-kernel results

Agent	Triton, 15 kernels	HIP, 16 kernels
GEAK v3.2	1.64x	1.99x
GEAK v4	2.21x	2.72x

On mixed HIP and Triton benchmarks, GEAK v4 raises the geometric-mean speedup on both language paths relative to GEAK v3.2. The improvement comes from the workflow structure—specialist roles, independent verification, and cumulative round-by-round learning—rather than from simply using a larger model.

Figure 3: Budget-Performance Scaling Law of GEAK v4 Kernel Workflow

The head-kernel benchmark shows a clear budget-performance scaling behavior in GEAK v4. As average optimization time increases, the kernel workflow continues to turn additional search budget into better verified kernels, improving geometric-mean speedup from 1.300x at 47 minutes to 1.328x at 70 minutes and 1.511x at 104 minutes. This highlights that GEAK v4 is not a one-shot code-generation agent: its multi-round workflow, specialist-agent exploration, independent verification, and cross-round learning allow performance to scale with optimization time.

Head-kernel benchmark

Agent	Geomean speedup
Claude Code, out of the box	1.09x
GEAK v3.2	1.09x
GEAK v4	1.33x

The head-kernel results are especially relevant to serving workloads because they target operators that dominate real inference profiles. GEAK v4 delivers a materially larger gain on these already-tuned kernels, indicating that its hierarchy and verification loop are effective precisely where local improvements are harder and more valuable.

End-to-end serving throughput

Model	Configuration	Hardware / framework	Bottleneck kernel	Optimization	E2E uplift
Qwen3.5-27B-FP8	TP1, ISL/OSL 1K/1K, conc 64	MI300X, SGLang	gemm_a8w8_blockscale_kernel	Triton -> CK	+60%
Qwen3-14B-FP8	TP1, ISL/OSL 1K/1K, conc 64	MI300X, SGLang	gemm_a8w8_blockscale_kernel	Triton -> FlyDSL	+96%
Minimax-M3-MXFP8	TP8, ISL/OSL 8K/1K, conc 64	MI355X, vLLM	Sparse attention kernel	Triton -> Triton	+42.2%
Kimi K2.6 INT4	TP1, ISL/OSL 1K/1K, conc 64	MI300, SGLang	fused_moe_kernel_gptq_awq	Triton -> Triton	+25.4%
DSV4 Pro FP4	TP1, ISL/OSL 1K/1K, conc 64	MI355, vLLM	MLA	TileLang -> TileLang; TileLang -> Triton	+47%; +108%

These end-to-end runs demonstrate the core thesis of GEAK v4. The largest gains come from kernels the system writes or replaces after profiling the live server, not from isolated microbenchmark wins alone. Because the workflow triages by workload impact, rewrites the bottleneck path, overlays changes reversibly, and validates warm-server output parity, kernel-level improvements translate into real serving throughput improvements.

Case Study: MiniMax M3 Sparse Attention — GQA Prefill and Decode Index Kernels

MiniMax M3 uses sparse attention: an indexer scores 128-token KV blocks, selects top-k blocks, and the GQA attention kernel computes attention only over those selected blocks. GEAK v4 introduced two main optimizations:

CDNA-optimized sparse GQA prefill sub-tiling
Vectorized decode index-score GEMV for the single-query case

CDNA-Optimized Sparse GQA Prefill Sub-Tiling

The baseline sparse GQA prefill kernel processes each selected 128-token sparse KV block as one full tile. For every selected block, it loads the full K tile, computes tl.dot(q, k), applies online softmax, loads the corresponding V tile, and computes tl.dot(p, v).

This is portable, but not ideal for AMD CDNA. On gfx942 and gfx950, the full 128-token block can create an inefficient MFMA/LDS shape. GEAK v4 splits each 128-token sparse KV block into smaller SUB_K sub-tiles, so each QK/PV operation is better matched to CDNA MFMA execution. On AMD CDNA GPUs, this tile is not always the best MFMA/LDS shape.

Baseline

		k = tl.load(
    kv_cache_ptr
    + page * stride_kv_blk
    + 0 * stride_kv_kv
    + off_n[None, :] * stride_kv_pos
    + pid_kh * stride_kv_h
    + off_d[:, None] * stride_kv_d,
)

qk += tl.dot(q, k) * sm_scale_log2e
p = tl.exp2(qk - m_ij[:, None])

v = tl.load(...)
acc_o += tl.dot(p.to(v.dtype), v)

GEAK v4 Optimized

		NUM_SUB = BLOCK_SIZE_K // SUB_K

for sub_i in range(NUM_SUB):
    off_sub = tl.arange(0, SUB_K) + sub_i * SUB_K

    k_sub = tl.load(...)
    qk_sub = tl.zeros((BLOCK_SIZE_Q, BLOCK_SIZE_H, SUB_K), dtype=tl.float32)
    qk_sub += tl.dot(q, k_sub) * sm_scale_log2e

    p_sub = tl.exp2(qk_sub - m_ij[:, None])

    v_sub = tl.load(...)
    acc_o += tl.dot(p_sub.to(v_sub.dtype), v_sub)

GEAK v4 keeps the same online-softmax recurrence, updating m_i, lse_i, and acc_o across sub-tiles, making the result numerically equivalent apart from normal flash-attention reassociation.

Vectorized Decode Index-Score GEMV

The decode index-score kernel scores visible 128-token index-K blocks and produces block scores for the top-k selection stage. In the baseline, the kernel always uses tl.dot(k, q, out_dtype=tl.float32), even when the decode shape degenerates into a single-column GEMV.

For the common decode case:

[128, D] × [D, 1] -> [128, 1]

using a full MFMA-style matrix tile can be inefficient. GEAK v4 adds a special fast path when BLOCK_SIZE_HQ == 1: instead of using tl.dot, it performs a vectorized FP32 multiply-and-reduce. For wider multi-head or speculative-decode tiles, it keeps the original tl.dot path.

Baseline

		k = tl.load(...)  # [N, D]

kq = tl.dot(k, q, out_dtype=tl.float32)  # [N, HQ]
score = tl.max(kq, axis=0)

GEAK v4 Optimized

		k = tl.load(...)  # [N, D]

if BLOCK_SIZE_HQ == 1:
    q_vec = tl.sum(q, axis=1).to(tl.float32)  # [D]
    kq = tl.sum(
        k.to(tl.float32) * q_vec[None, :],
        axis=1
    )[:, None]
else:
    kq = tl.dot(k, q, out_dtype=tl.float32)

score = tl.max(kq, axis=0)

This avoids unnecessary matrix-tile overhead for single-query decode while preserving FP32 accumulation. The optimized path is explicitly documented as numerically equivalent to tl.dot.

Results

GEAK v4 was evaluated on MiniMax-M3-MXFP8 serving workloads under TP8. For an 8k input / 1k output workload at concurrency 64, output throughput improved from 1183.8 to 1683.3 tokens/s, a +42.2% gain. Median TTFT improved from 1896 ms to 937 ms, a -50.6% reduction, and median TPOT improved from 51.4 ms to 36.2 ms, a -29.5% reduction. Accuracy on GSM8K remained effectively unchanged, with the reported deltas within the confidence range.

Conclusion

GEAK v4 is built around one principle: the unit of value is the model, not the kernel. A performance agent that optimizes serving throughput must first identify where the workload spends time, then decide whether the best lever is a configuration switch, a backend substitution, a library path, or a newly authored kernel.

Three design choices make this operational. Deterministic JavaScript Workflows provide reproducible control while reserving LLMs for structured technical judgment. A fractal two-altitude design lets the end-to-end layer recursively reuse the kernel workflow without weakening kernel-level verification. An evidence-gated learning loop captures what worked, where it worked, and where it failed, while never allowing remembered knowledge to override on-box measurement.

The result is an autonomous optimization system that turns AMD Instinct GPUs capability into software-visible performance faster. For AMD, GEAK v4 accelerates model bring-up and reduces dependence on scarce manual kernel expertise. For customers, it points toward faster, more cost-efficient inference on AMD GPUs without requiring every team to build its own dedicated kernel engineering organization.

GEAK v3 Foundation: https://rocm.blogs.amd.com/artificial-intelligence/kernel-optimization-agent/README.html
GEAK v4 Repository: https://github.com/AMD-AGI/GEAK/tree/v4.0.0

Footnotes

Endnote:

Testing conducted by AMD Performance labs on July 8, 2026.

System Configuration

AMD MI300X GPU system
Model: Microsoft C278A, CPU: 2x Intel(R) Xeon(R) Platinum 8480C, 56-Core (112 cores / 224 threads total)
NUMA: 2 NUMA nodes (1 per socket). NUMA auto-balancing disabled
Memory: 2.0 TiB (per-DIMM vendor/speed unavailable — needs root/dmidecode)
Disk: 3x NVMe (894 GB + 1.7 TB + 1.7 TB); total mounted FS 232 TiB (3.1 TiB used)
GPU: 8x AMD Instinct MI300X 192GB (gfx942, SKU M3000108)
Host OS: Ubuntu 22.04.5 LTS (kernel 5.15.0-70-generic)
System BIOS: C2789.5.BS.1C19.AG.2
Bios Vendor: American Megatrends International, LLC.
Host GPU Driver: ROCm 7.0.0 (amdgpu kernel module 6.16.13)
AMD MI355X GPU system
Model: Supermicro AS -4126GS-NMR-LCC, CPU: 2x AMD EPYC 9575F 64-Core (128 cores / 256 threads total)
NUMA: 2 NUMA nodes (1 per socket). NUMA auto-balancing disabled
Memory: 3.0 TiB (24x 128 GB Micron DDR5, 6400 MT/s, configured 6000 MT/s)
Disk: 2x NVMe 3.5 TB (Micron 7450) + LVM 3.5 TB + NFS 59 TB; total mounted FS 71 TiB (68 TiB used)
GPU: 8x AMD Instinct MI355X 288GB (gfx950, DID 0x75a3)
Host OS: Ubuntu 22.04.5 LTS (kernel 5.15.0-70-generic)
System BIOS: 1.4a
Bios Vendor: American Megatrends International, LLC.
Host GPU Driver: ROCm 7.2.0 (amdgpu kernel module 6.16.6)