| name | env-setup |
| description | Get a container ready to tune GPU ops on AMD Instinct — audit which tuning tools are present, install the missing bench clients (hipblaslt-bench, ckProfiler, aiter tuners) at versions that match the runtime, and avoid the version-skew failures that silently produce useless tuned artifacts. Use before any tuning run in a vllm or sglang image. |
Environment setup for tuning
Serving images ship what is needed to run tuned kernels, not to produce them. The
tuned config files are there; the tools that write them are not. This skill closes that gap.
Budget: ~1 minute for ckProfiler, ~2 for aiter source, ~20 for hipblaslt-bench. Install only
what your task needs.
1. Audit first
Never assume. Run the inventory in a pristine container from the image you will tune in:
docker run --rm --entrypoint bash \
-v $PWD/audit_tools.sh:/tmp/a.sh:ro \
--device=/dev/kfd --device=/dev/dri --group-add video \
<image> /tmp/a.sh
audit_tools.sh prints OK / MISSING / INFO per tool and changes nothing. Grep for
MISSING to get your work list.
Two things that will bite you here:
--entrypoint bash — the vllm image's entrypoint is the vllm CLI. Without the
override your script is parsed as vllm arguments and you get
vllm: error: unrecognized arguments.- A container you have already worked in is not evidence. It carries your earlier
installs. Audit a fresh one.
image_tool_matrix.md records what the two target images shipped when measured, and which
image differences change your workflow.
2. Pass the GPUs in
--device=/dev/kfd --device=/dev/dri --group-add video
Without these, rocminfo fails and arch detection fails with it. The symptom surfaces far
from the cause: aiter raises ImportError: cannot import name 'dtypes' from 'aiter', which
reads like a broken install but is arch detection failing and disabling the CK/HIP op
registrations. When a GPU library throws a strange import error, run rocminfo before
debugging the library.
3. Install what is missing
| tool | how | cost | needed for |
|---|
| ckProfiler | apt-get install composablekernel-ckprofiler | ~1 min | CK instance selection → ckprofiler_install.md |
| aiter source | git clone + checkout the tag matching the installed wheel | ~2 min | aiter tuners (vllm image only) |
| hipblaslt-bench | source build, clients-only | ~20 min | hipBLASLt solution racing → hipblaslt_bench_build.md |
Do them in that order and stop when your task is covered. Only hipblaslt-bench requires a
build, and only because no package ships it; the clients-only path avoids regenerating
kernels and turns a multi-hour build into a coffee break. hipblaslt_bench_build.md has the
recipe plus each dependency mapped to the distinct stage it fails at — the build fails
three separate times on three separate missing packages, one of them at 98%.
4. Match versions to the runtime, not to HEAD
This is the rule that separates a tuning run that helps from one that appears to.
| tool | match it to | how to read it |
|---|
| hipBLASLt clients | the container's ROCm | cat /opt/rocm/.info/version → use branch release/rocm-rel-7.2 for 7.2.x |
| ckProfiler | the ROCm apt repo already in the image | take the candidate, do not pin across ROCm versions |
| aiter tuners | the installed aiter | python3 -c "import aiter._version as v; print(v.__version__)", then resolve to a tag with git ls-remote --tags origin |
A tuned artifact is a lookup-table entry keyed by shape, dtype and architecture. Produced by
a mismatched tool, it lands where the runtime never looks, or is looked up and rejected.
Neither failure raises an error. You get baseline performance and a report claiming a
win. This is why every tuning run in this skillset ends with
../tuning-core/engagement_verification.md rather than with a timing.
Corollary: never carry a tuned artifact between architectures. gfx942 and gfx950 differ in
CU count and FP8 dialect; the lookup keys include the arch for a reason.
The version string is not the compiler
Same architecture, same GPU, same kernel source, same config, same reported library version —
and 1.89x apart. Measured on one MI355X, aiter's Gluon MXFP4 GEMM at M512 N4096 K4096, both
images reporting triton 3.6.0:
| vllm image | sglang image |
|---|
| median | 0.0616 ms | 0.1163 ms |
| triton path | dist-packages/triton | /sgl-workspace/triton-custom |
| VGPRs | 254 | 256 (at the cap) |
| VGPRs spilled | 0 | 63 |
| scratch traffic | none | 55 stores + 55 reloads |
| MFMA count | 64 | 64 |
| LDS | 69 632 B | 69 632 B |
Identical instruction mix and identical LDS. The whole difference is that one compiler build
lands two registers over the 256-VGPR cap and spills 63 values into the inner loop. Both
report 3.6.0 because a vendored fork keeps the version string of the release it forked
from.
Three things follow, in increasing order of how much they cost:
- Read
triton.__file__, not triton.__version__, when recording what produced a result.
The version answers a question nobody was asking.
- Engagement verification will not catch this. The config is found, the lookup hits, the
intended kernel runs. Every check in
../tuning-core/engagement_verification.md passes.
The kernel is simply compiled worse, and nothing in the Python layer can see it. What sees
it is .vgpr_spill_count in the generated .amdgcn — dump it with TRITON_CACHE_DIR=/tmp/tc
and read the metadata, which is worth doing once per image for any kernel you care about.
- Tuning does not fully recover it. Sweeping the sglang image lifted that shape from
0.1163 to 0.0822 ms — a real 1.4x — and it is still slower than the vllm image untuned.
The compiler build sets a ceiling the config space cannot reach past. If you are comparing
a tuned result in one image against a baseline in another, you are measuring the images.
5. Confirm before tuning
Re-run audit_tools.sh. Everything your task needs should read OK, and gpu-arch should
name a real target (gfx942), not a rocminfo failed message. Then pin an idle GPU
(rocm-smi --showuse; export HIP_VISIBLE_DEVICES=<n>) and take a baseline with
../benchmark/run_case.py --smoke — if the harness cannot produce a sane number on a known
shape, nothing downstream of it is trustworthy.
Files
audit_tools.sh — read-only inventory; run it first and again at the endimage_tool_matrix.md — measured contents of both target images, and the aiter wheel-vs-source splitckprofiler_install.md — CK bench client: install, interface, and what it does and does not provehipblaslt_bench_build.md — the one real build, with its three failure modes
Next
Environment ready → ../tuning-core/SKILL.md for the loop itself, then the skill for your
source language or serving framework.
