amd-agi/geak/tuning-core
Claude Code and agent skill by amd-agi/geak
- community
Checking that the SKILL.md file is still in the repository.
The universal loop for tuning GEMM and other GPU ops on AMD Instinct — scope the search space, measure honestly, gate on correctness, and prove the tuned artifact actually engaged. Start here before any language- or framework-specific tuning skill.
Install
npx vskill@latest install amd-agi/geak/tuning-coreSKILL.md
Tuning core
This is the discipline every other tuning skill in this set specializes. Read it first. It is language-agnostic: the same six steps apply whether you are tuning a Triton kernel, racing hipBLASLt solutions, or deploying an aiter config into a live server.
The claim this skill exists to defend
A tuning run that reports a speedup has proven nothing until you have shown that theartifact it produced is the code the machine actually runs.
The dominant failure in GPU op tuning is not picking a bad config. It is picking a good config that never gets used — a lookup key that misses, a config file in a folder nothing reads, an env var the live dispatch path ignores. Nothing errors. The tuner prints a win. End-to-end performance does not move. If you take one thing from this skill, make it step 6.
The loop
Steps 1, 2, 4 and 6 are where tuning efforts fail. Step 3 is the part everyone thinks is the hard one; it is usually the most automated.
1. Scope
Answer these before running anything. Guessing here wastes the entire run.
- Which op, and which implementation is live right now? Not which one you expect —
which one dispatches. A framework may route the same nn.Linear through torch, a library heuristic, or a fused kernel depending on env vars and shape.
- Which shapes? Capture real ones if the workload exists (see the per-framework
skills). Otherwise generate a corpus spanning the regimes — see ../benchmark/.
- Which dtype, and which dialect of it? On gfx942, FP8 is FNUZ. On gfx950 it is OCP.
Same 8-bit layout, different exponent bias — a mismatch is silently wrong numerics, not a crash.
- What is the arithmetic intensity? A K-heavy 4096×4096×4096 GEMM and a decode-shaped
M=1 GEMV are different problems with different winning strategies. Compute-bound work wants tiles and MFMA scheduling; memory-bound work wants split-K, vectorization, and occupancy.
2. Baseline
Measure the existing path first, on the same box, in the same process shape you will use for the final comparison. A tuned number compared against a remembered or documented baseline is not a measurement.
Pin your GPUs. On a shared box, other tenants will move your numbers:
See measurement.md for how to time correctly. It is not as simple as wrapping a timer.
3. Search
Enumerate → prune → race. The details are per-tool, but the strategy is shared; see search_strategy.md. Two rules that hold everywhere:
- Prune before you race, not after. Search spaces are combinatorial. Constraining to
hardware-sane candidates first (tile shapes that match the MFMA instruction, workgroup counts that fill the CUs, configs that do not spill) cuts run time by an order of magnitude without costing you the winner.
- Change one variable at a time when hand-tuning. If you change three things and it
gets faster, you have learned nothing transferable.
4. Gate on correctness
Never accept a config on timing alone. A config that races ahead because it computes the wrong thing is the easiest way to ship a regression.
The gate must be a relative error measure. Absolute error is meaningless at scale — a bf16 GEMM at K=4096 accumulates ~1.0 absolute max error against an fp32 reference purely from dtype, with nothing wrong. See correctness_gates.md, which works that exact case.
The convention used across the AMD tuning tools is err_ratio < 0.05 against a higher-precision reference. Match it unless you have a reason not to.
5. Deploy
Write the artifact where the live path reads it. This is per-tool and per-framework, and it is where lookup-key mismatches are introduced. The recurring trap: a tuned entry is keyed on a tuple, and every field must match the live call exactly. A single wrong field — most often a bias flag or a device-name string — means a 100% lookup miss.
Tuned artifacts do not transfer across GPU architectures. Some lookup keys literally include the CU count (304 on MI300X, 256 on MI355), so a config tuned on one part is a guaranteed miss on the other. Re-tune per arch; never copy. When moving a whole tuning setup to a new part, arch_migration.md lists what has to be re-derived and, more importantly, which of it fails silently rather than erroring.
6. Verify engagement — then re-measure
Two distinct checks, both required.
(a) Did the artifact engage? Find the positive signal that the tuned path was taken. Every per-tool skill in this set ends with a concrete one. Examples of the form it takes:
If you cannot find a positive engagement signal for a path, that is the first thing to build, before tuning anything. Tuning without it is unfalsifiable.
Two traps that make a log-line check report failure on success, both covered in engagement_verification.md: the hit line may only be emitted under a debug flag you cannot always set, and hit counts are usually cached and so measure shape diversity rather than call frequency. Prefer kernel identity from a profile for anything you intend to publish.
(b) Did the number move where it matters? An isolated kernel speedup is not an end-to-end speedup. A kernel that is 3× faster but 2% of total GPU time buys you 1.3%. Check the op's share of total time before investing in it, and re-measure end-to-end after.
When to stop
Stopping rules matter as much as search:
- Gate met → stop. Define the target before you start.
- Three consecutive iterations under ~2% improvement → the config space has plateaued.
Further tile tuning is not the lever; a structural change is.
- Profiler shows the kernel near a hardware limit (MFMA utilization saturated, or at
achievable bandwidth for a memory-bound op) → stop, you are done.
- Register pressure sitting at an occupancy boundary → more config search will not
help; the kernel needs restructuring.
Anti-patterns
| Anti-pattern | Why it burns you | |
|---|---|---|
| Comparing backends without checking operand layout | A row-major vs column-major B operand changes throughput substantially. Two "GEMM benchmarks" at the same M/N/K can be different problems. | |
| Absolute error as the correctness gate | Scales with K and dtype; either passes everything or fails everything. | |
| Tuning on a shared, unpinned GPU | Another tenant's load lands directly in your numbers. | |
| Trusting a tuner's own reported speedup as the result | It measured its candidate in isolation, not your live path. | |
| Copying tuned configs across architectures | Different CU count, LDS budget, and instruction set. Often a silent lookup miss. | |
| Reporting a win without an engagement check | The most common way a "successful" tuning effort delivers zero. |
Where to go next
| You are tuning | Read | |
|---|---|---|
| A Triton or Gluon kernel you control | ../tuning-triton/ | |
| A FlyDSL kernel | ../tuning-flydsl/ | |
| A raw HIP kernel | ../tuning-hip/ | |
| CK / ckProfiler instances | ../tuning-ck/ | |
| hipBLASLt solution selection | ../tuning-hipblaslt/ | |
| aiter's per-shape config DBs | ../tuning-aiter/ | |
| A live vLLM server | ../tuning-in-vllm/ | |
| A live sglang server | ../tuning-in-sglang/ | |
| Missing tools in your container | ../env-setup/ |
Supporting detail in this skill: measurement.md, correctness_gates.md, search_strategy.md, engagement_verification.md, arch_migration.md, clocks_and_power.md, graph_captured_benchmarking.md.
If you are tuning a serving decode path, read `graph_captured_benchmarking.md` first. It governs both how you must time a kernel and whether your change can take effect at all, and it changes the noise floor every other page's advice is calibrated against.
Security
- Security scanPassed, 95 of 100Pattern checks on the skill's files.
- Model reviewPassed
Quality evaluation
- Task evaluationNot runNo evaluation yet. Run one in Skill Studio.
Source
- SKILL.md in the repositoryCheckingperf_knowledge/expert_skills/tuning/tuning-core/SKILL.md
Popularity
- Repository stars
- 182
- Repository forks
- 56
- 7-day trend score
- 26
Stars and forks are for the whole GitHub repository, not this skill alone. Not refreshed yet.
Works with
Any agent vskill installs to: 56 targets in all. These 39 have a page on this site.
- Claude Code
- Cursor
- GitHub Copilot
- Windsurf
- Codex
- Gemini CLI
- Cline
- Amp
- Roo Code
- Goose
- Continue
- Kilo Code
- Junie
- Trae
- OpenHands
- Qwen Code
- AI Agent
- Tabnine Chat
- Sourcegraph Cody
- Amazon Q
- Firebase Genkit
- Sweep
- Devin
- Phind
- Mentat
- Claude Desktop
- Aider
- GPT Pilot
- Plandex
- bolt.new
- v0
- Replit Agent
- Anse
- TypingMind
- BetterChatGPT
- LibreChat
- Warp AI
- Pieces
- Fig AI
Embed badge
[](https://verified-skill.com/skills/amd-agi/geak/tuning-core)