tuning-in-vllm

Tune fused-MoE and GEMM kernels inside a vLLM serving deployment on AMD Instinct GPUs.

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

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve? Tuning kernels inside a live vLLM serving stack fails in non-obvious ways: config filenames depend on a device name that torch reports as empty, shipped configs rarely match MI355 hardware, and there is no signal that a tuned config was actually loaded. This Skill captures the real shapes a model issues, tunes the right kernels, deploys configs where vLLM looks, and proves engagement from the server log. ## Core Features & Use Cases - Shape capture from real traffic: Rank GEMM shapes by call count using HIPBLASLT_LOG_MASK and AITER_LOG_TUNED_CONFIG logs instead of tuning generated sweeps. - Two tuning surfaces: Tune dense GEMM through the aiter gradlib tuner and fused MoE through vLLM's benchmark_moe.py with its M-bucketed Triton config format. - Deployment and verification: Deploy via VLLM_TUNED_CONFIG_FOLDER and confirm from startup logs that vLLM loaded the tuned config rather than falling back to defaults. - Use Case: On an MI355 serving a MoE model, capture the issued shapes, tune the top fused-MoE configs, drop them into a folder pointed to by VLLM_TUNED_CONFIG_FOLDER, and A/B serving throughput by toggling the variable. ## Quick Start Use the tuning-in-vllm skill to capture the GEMM shapes my vLLM server issues on my MI300X, tune the fused-MoE configs, and verify from the startup log that the tuned config was loaded.

Frequently Asked Questions about tuning-in-vllm

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

FAQPage Schema
How do I tune fused MoE kernels in vLLM on AMD GPUs?

Run the benchmark_moe.py script shipped in the vLLM image with the --tune flag against your model. It produces a JSON config keyed by M bucket with Triton parameters like BLOCK_SIZE_M, num_warps, and waves_per_eu, which you deploy via VLLM_TUNED_CONFIG_FOLDER.

How do I find which GEMM shapes my vLLM model actually issues?

Set HIPBLASLT_LOG_MASK=32 before starting vllm serve and drive representative traffic. The log emits runnable hipblaslt-bench commands per shape; count and rank them with sort and uniq to tune the most frequent shapes first.

Why is my vLLM tuned MoE config not being loaded?

The most common cause is building the filename from torch.cuda.get_device_name(), which returns an empty string on ROCm. vLLM uses current_platform.get_device_name() instead, so the filename must match names like AMD_Instinct_MI355_OAM exactly.

Does VLLM_TUNED_CONFIG_FOLDER replace the shipped vLLM configs?

No. vLLM checks the folder pointed to by VLLM_TUNED_CONFIG_FOLDER first, then falls back to the package's fused_moe/configs directory, so both are searched. This differs from sglang, where the equivalent variable replaces the package root entirely.

How do I verify vLLM picked up my tuned config?

Grep the server startup log for 'Using configuration from' versus 'Using default MoE config'. The miss warning prints the exact filename vLLM wanted in both search locations, making filename mismatches immediately visible.

Can I reuse sglang tuned configs in vLLM on the same GPU?

No. The same silicon reports different device names across frameworks, for example AMD Instinct MI355X in sglang versus AMD Instinct MI355 OAM in vLLM, so config filenames do not transfer between the two.