tuning-hip

Tune hand-written HIP kernels on AMD Instinct GPUs using launch geometry and rocprofv3 profiling.

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

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve? Hand-written HIP kernels have no autotuner, so developers must choose launch geometry and compile flags themselves and often rely on flawed timing loops or guesswork about device limits. This Skill provides a disciplined workflow for selecting block/grid sizes against wave64, CU, and LDS constraints, measuring kernels correctly, and using rocprofv3 to verify which kernel actually ran. ## Core Features & Use Cases - Device-aware launch geometry: Read wavefront size, CU count, and LDS limits from rocminfo for gfx942 and gfx950 targets, and size blocks as multiples of 64 with grids covering all CUs. - Compile-flag search space: Use --offload-arch, launch_bounds, and -Rpass-analysis=kernel-resource-usage to prune register- and occupancy-doomed variants before benchmarking. - Ground-truth profiling with rocprofv3: Generate kernel stats CSVs that rank kernels by cost, confirm a tuned kernel engaged, and expose hidden costs like transposes and buffer fills. - Use Case: You ported a GEMM kernel from an NVIDIA GPU to MI300X and it underperforms. Use this Skill to correct the wave64 block sizing, synchronize and repeat your timing loop, then confirm via rocprofv3 that the intended kernel dominates runtime. ## Quick Start Ask the AI to tune your HIP kernel for your AMD Instinct GPU by reading device constants, fixing the launch geometry, and verifying the result with a rocprofv3 kernel trace.

Frequently Asked Questions about tuning-hip

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

FAQPage Schema
How do I tune HIP kernel launch geometry on AMD GPUs?

Read device constants with rocminfo first: wavefronts are 64 threads, so block sizes should be multiples of 64, and grids should cover all CUs (304 on gfx942, 256 on gfx950). Check LDS per workgroup limits before sizing staged tiles.

How to use rocprofv3 to verify which kernel actually ran?

Run rocprofv3 with --kernel-trace --stats -f csv to produce a kernel stats CSV listing every executed kernel by name with call counts and durations. Seeing the tuned kernel's name in the trace, and the untuned one absent, is proof of engagement.

Why are my HIP kernel timings faster than hardware peak?

Kernel launches are asynchronous, so timing without hipDeviceSynchronize() or HIP events measures enqueue cost, not execution. The inflation is shape-dependent: small kernels can show 2.4x or worse error while large shapes hide the bug.

Does a HIP kernel compiled for gfx942 run on gfx950?

No, binaries are architecture-specific and --offload-arch makes this fail at build time. Tuned launch geometry also does not transfer: block sizes and LDS staging balanced for 304 CUs are unbalanced on 256, so re-tune geometry per architecture.

What compile flags matter when tuning HIP kernels?

Always set --offload-arch explicitly for your target. Use __launch_bounds__ to cap registers per thread for higher occupancy, and -Rpass-analysis=kernel-resource-usage to check registers, LDS, and occupancy at compile time before benchmarking variants.