amd-agi/geak/env-setup

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

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.

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:

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 end
• image_tool_matrix.md — measured contents of both target images, and the aiter wheel-vs-source split
• ckprofiler_install.md — CK bench client: install, interface, and what it does and does not prove
• hipblaslt_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.