All authors
Claude Skills by AMD-AGI
github.com/AMD-AGI23 skills0 installs59 views
- RooflineTurn the standardized Top-N into a per-kernel answer to: **"how much of this hardware's ceiling is this kernel already getting, and therefore how much is left to win?"** You (an agent) can execute this skill by reading this file. `roofline_tools.py` next to it holds mechanical primitives (peak-table parsing, counter parsing, unit math) — use them if they import, do the arithmetic yourself if they don't. **A broken helper must not disable this skill; a broken skill must not fail the run.** ---Votes: 0GitHub stars: 182
- ContributeContribute a human-authored, e2e-validated optimization recipe (an "expert skill") to GEAK — scaffold, fill, validate by scope, and open a PR.Votes: 0GitHub stars: 182
- Apply Flydsl Moe To Vllmvllm serve <model> -tp 4 --gpu-memory-utilization 0.9 \ --trust-remote-code --tool-call-parser kimi_k2 --enable-auto-tool-choice \ --reasoning-parser kimi_k2 --no-enable-prefix-caching --mm-encoder-tp-mode data ``` 4. **Verify startup (in order, no error)**: weight load → (silent in-place convert) → `Available KV cache` (no OOM) → `[flydsl-precompile] BEGIN exhaustive eager warmup ...` → `[flydsl-precompile] DONE eager warmup (N/N); all FlyDSL exes compiled+cached outside capture` (from step ...Votes: 0GitHub stars: 182
- Ensure FlydslThis skill makes `import flydsl, kernels.moe_gemm_2stage` resolve from ONE consistent tree, so the FlyDSL optimization skills (`flydsl_rewrite_quantized_moe` for authoring, `apply_flydsl_moe_to_vllm` for the e2e apply-back) can run. It is a **dependency skill**: other skills / the System Architect invoke it; it is not selected from the profile Top-N on its own.Votes: 0GitHub stars: 182
- Flydsl Decode Moe Stage1 BlkmapTrigger on the **problem signature, not a specific model**: a **grouped-GEMM MoE stage-1 (gate+up)** kernel at **decode / small-M** where repeated expert-weight HBM traffic is a **material** part of latency (the baseline may be mixed rather than purely memory-bound), the MoE token-sort pads each expert to a fixed sort-block granularity `B` (commonly 64 rows), and routing gives many experts **>=2 sort-blocks**. A per-sort-block compute tile then re-reads the same low-precision expert weight once Votes: 0GitHub stars: 182
- Flydsl Fp8 Blockscale GemmTrigger on the **problem signature, not a specific shape**: *any* **CK a8w8 fp8 block-scale** dense GEMM (`weight_block_size [128,128]`, arbitrary fp32 per-`[128,128]` block scales, bf16 out; the profiled hotspot name matches `kernel_gemm_xdl*blockscale` / `gemm_a8w8_blockscale`) being rewritten **CK → FlyDSL** on **MI355X / gfx950 (CDNA4, OCP fp8_e4m3)**, where the FlyDSL **native** fp8 scaled-MFMA path fails parity against the CK oracle (the symptom that blocks the port) and CK config-JSON / tVotes: 0GitHub stars: 182
- Flydsl Fp8 Gemm PlaybookAn fp8 **a8w8 block-scale** model (weight_block_size [128,128]) on MI300X where a dense GEMM head is the top GPU-time mass and the per-(N,K) Triton config-JSON overlay is already exhausted (config tiling gives only ~1.04–1.06x). Specifically the **K-heavy / narrow-N down-proj** (e.g. N=5120, K=17408, ~18% GPU on Qwen3.5-27B-FP8): config tuning is near its ceiling, so the remaining lever is a **code-level authored core**, not more config JSON.Votes: 0GitHub stars: 182
- Flydsl Fused Attention BackwardAn attention **backward** kernel is a top GPU-time entry in a training or fine-tuning run and the live path is a poor fit for the shape. The two situations that make authoring worthwhile:Votes: 0GitHub stars: 182
- Flydsl Prefill Moe Stage2 Fp8partialTrigger on the **problem signature, not a specific model**: a **grouped-GEMM MoE stage-2 (down-proj + top-k reduce)** at **prefill / large-M** on gfx950 that is **HBM-traffic-bound on the per-(token, top-k slot) partial tensor**. Down-proj writes `top_k` partial rows per token and the reduce sums them, so that partial is the largest stage-2 HBM stream — written once and read once. Applies where the down path runs the **non-accumulating** variant (each expert-slot's partial stored standalone, theVotes: 0GitHub stars: 182
- Flydsl Rewrite Quantized MoeA hot kernel that is an **int4 (W4A16 GPTQ/AWQ) or fp8 (blockscale) quantized GEMM / fused-MoE grouped GEMM** on MI300X (gfx942), where the live path is a Triton `@triton.jit` body (or a closed `.co`) and Triton body micro-opts have hit their ceiling. The Triton int4-unpack loop is latency-bound at ~15–20% HBM, so body-only tuning caps at ~1.1×; a full **FlyDSL rewrite** of the candidate reaches ~60–65% HBM at decode and ~3.6–6.2× over the Triton golden (speedup grows with M). This is the lever Votes: 0GitHub stars: 182
- Gluon AuthoringThe workflow is working a **Gluon** kernel on gfx942/gfx950 and needs two things this file supplies: the **Gluon API surface**, and the **do-not-write list** — the constructs and instruction choices that compile and then silently cost you.Votes: 0GitHub stars: 182
- Mla Tilelang To TritonThe MLA (multi-head latent attention) **decode** path is the bottleneck on a DeepSeek-class model on MI300X (gfx942) and the live implementation is a TileLang kernel. Port it to Triton to (a) get an editable, CUDA-graph-capturable core the workflow can further tune, and (b) close the gap to the hand-tuned aiter/asm path while staying author-friendly.Votes: 0GitHub stars: 182
- Env SetupGet 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.Votes: 0GitHub stars: 182
- Tuning AiterTune GEMM and fused ops through aiter on AMD Instinct — run the gradlib tuner across hipblaslt/triton/flydsl/asm/torch backends, write tuned config CSVs to the directory the runtime actually reads, and prove engagement with AITER_LOG_TUNED_CONFIG. Use when tuning ops dispatched through aiter in vLLM or SGLang.Votes: 0GitHub stars: 182
- Tuning CkTune GEMMs on the Composable Kernel path — race CK's compiled instances with ckProfiler, read the instance name to understand why one won, and route the result into a framework through aiter's per-op CK tuners. Use when the op dispatches through CK or when establishing what a CK path could reach.Votes: 0GitHub stars: 182
- Tuning CoreThe universal loop for tuning GEMM and other GPU ops on AMD Instinct — scope the search space, measure honestly, gate on correctness, and prove the tuned artifact actually engaged. Start here before any language- or framework-specific tuning skill.Votes: 0GitHub stars: 182
- Tuning FlydslTune FlyDSL kernels on AMD Instinct — build and prune a Config space, set the autotune key, and race FlyDSL against the other backends through aiter's libtype mechanism. Use when the kernel is written in FlyDSL or when deciding whether the FlyDSL path is competitive for a shape.Votes: 0GitHub stars: 182
- Tuning HipTune hand-written HIP kernels on AMD Instinct — choose launch geometry against the wave64/CU/LDS limits of the target, use rocprofv3 to get ground-truth kernel timings and occupancy, and confirm which kernel actually ran. Also the profiler reference for verifying engagement on any backend.Votes: 0GitHub stars: 182
- Tuning HipblasltTune GEMMs on the hipBLASLt path — capture the exact problem from a live workload, race all solutions with hipblaslt-bench, replay the winner by solution index, and avoid the bracket-number and hot-cache traps. Use when the op dispatches through hipBLASLt (torch matmul, aiter libtype=hipblaslt).Votes: 0GitHub stars: 182
- Tuning In SglangTune ops inside an SGLang serving deployment on AMD Instinct — capture real shapes, tune the Triton fused-MoE path into the correct triton_<version> config directory, deploy via SGLANG_MOE_CONFIG_DIR, and read the server log to distinguish an exact config hit from a silently degraded version fallback. Use when tuning for a live SGLang workload.Votes: 0GitHub stars: 182
- Tuning In VllmTune ops inside a vLLM serving deployment on AMD Instinct — capture the shapes the model actually issues, tune fused-MoE and GEMM paths, deploy configs where vLLM reads them via VLLM_TUNED_CONFIG_FOLDER, and confirm from the server log that the tuned config was loaded. Use when tuning for a live vLLM workload.Votes: 0GitHub stars: 182
- Tuning TritonTune Triton and Gluon GPU kernels on AMD Instinct — build a config space including the AMD-only knobs (waves_per_eu, matrix_instr_nonkdim, kpack), prune it before racing, set the autotune key correctly, and confirm the win exceeds the noise floor. Use when the kernel you are tuning is written in Triton or Gluon.Votes: 0GitHub stars: 182
- Run ReportBuild or read the report for a GEAK run — one self-contained page (report/geak_run_report_<model>.html) with every LLM API call's cost, tokens, billed time and elapsed time, arranged by phase, by workflow invocation and as a role tree of agents. Use when asked where a run's time or money went, which phase or agent dominates it, or which work is cheap enough to delegate to a smaller model.Votes: 0GitHub stars: 182