tuning-hipblaslt

Tune hipBLASLt GEMM kernels by racing solutions and replaying winners by solution index.

178|52|Updated Jul 30, 2025
One-click install
npx skills add https://github.com/AMD-AGI/GEAK --skill tuning-hipblaslt-amd-agi
Or copy as Structured Prompt for Agent
Please help me install this Agent Skill.
Skill: tuning-hipblaslt
Source: https://github.com/AMD-AGI/GEAK/tree/main/perf_knowledge/expert_skills/tuning/tuning-hipblaslt
Command: npx skills add https://github.com/AMD-AGI/GEAK --skill tuning-hipblaslt-amd-agi

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve? GEMM operations dispatched through hipBLASLt on AMD Instinct GPUs often run slower than they could because the library's default heuristic does not always pick the fastest pre-compiled solution, and naive benchmarking with hipblaslt-bench produces inflated or misleading numbers through cache reuse, clock boost, and misread solution identifiers. ## Core Features & Use Cases - Exact problem capture: Log every matmul from a live workload with HIPBLASLT_LOG_MASK=32 and replay the emitted hipblaslt-bench commands verbatim, avoiding the M/N transpose trap. - Full solution racing: Run --algo_method all with --print_kernel_info to race every pre-compiled solution and record the true solution index rather than the unstable bracket number. - Measurement hygiene: Use --rotating buffers, standalone index replay, and interleaved A/B comparison against the heuristic default to separate real gains from race-inflated numbers, including gfx950-specific guidance where MX formats cannot be raced at all. - Use Case: A serving engineer notices torch matmuls dominating latency on an MI300X, captures the exact GEMM shapes from the workload, races all 1231 solutions, and deploys the winning solution index through the aiter tuner CSV for a measured throughput gain. ## Quick Start Capture the exact GEMM problems from my running PyTorch workload using HIPBLASLT_LOG_MASK, then race all hipBLASLt solutions for the hottest shape and tell me the winning solution index to deploy.

Frequently Asked Questions about tuning-hipblaslt

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I tune hipBLASLt GEMM performance on AMD GPUs?

Capture the exact problem from your workload with HIPBLASLT_LOG_MASK=32, then run hipblaslt-bench with --algo_method all and --print_kernel_info to race every solution. Replay the winner by its solution index and deploy it through aiter, TunableOp, or the hipBLASLt extension API.

How to find the hipBLASLt solution index from hipblaslt-bench output?

Run hipblaslt-bench with --print_kernel_info, which prints a --Solution index: line for the winner. The bracket number prefixing each result row is only an enumeration position that drifts between runs and cannot be used with --algo_method index.

Why does hipblaslt-bench report higher TFLOPS than my workload achieves?

Racing many solutions back-to-back keeps clocks boosted and caches warm, inflating results by up to 20% on gfx950. Replay the winning solution standalone by index and A/B it against the heuristic default to get the honest number.

Does hipblaslt-bench --algo_method all work with MXFP4 or MXFP8?

No, MX data types reject --algo_method all entirely because MX GEMMs route through RocRoller rather than Tensile solutions. Use --algo_method heuristic --requested_solution -1 to enumerate candidates, but note no solution index is printed, so alternatives cannot be deployed.

Is racing hipBLASLt solutions worth it on MI355X (gfx950)?

Measured gains over the default heuristic range from -2% to +14% on gfx950, far below the 19-40% the race output claims. The gfx950 default uses Origami analytical selection, so always interleave-replay both candidates before adopting a winner.

Why does my hipBLASLt benchmark overestimate small GEMM shapes?

Repeated iterations reuse the same A and B buffers, which stay resident in cache and inflate results by up to 43% on gfx942 for cache-resident shapes. Pass --rotating with a working set larger than last-level cache to measure realistic memory traffic.