[Skill] Add liger-kernel-perf skill for Triton kernel performance optimization (#1185)

## Summary

Adds a new Claude Code skill (`liger-kernel-perf`) that optimizes
existing Liger Kernel Triton kernels through a 3-stage pipeline: Profile
-> Optimize -> Finalize. The skill automatically diagnoses bottlenecks,
generates multiple optimization variants, benchmarks them against the
full suite, and creates a PR with the winning variant -- all while
maintaining correctness.

Tested on `fused_add_rms_norm` (see PR #1187), which achieved up to 70%
backward speedup on H100 via register pressure reduction.

## What the skill does

- **Stage 1 (Profiler):** Runs baseline benchmarks, detects GPU
architecture (Ampere/Hopper/Blackwell), optionally runs NCU profiling,
classifies bottleneck (memory-bound vs compute-bound), produces
optimization profile with recommended strategy order
- **Stage 2 (Optimizer):** Autonomous optimization loop -- tries
parameter tuning first (manual sweep, NOT @triton.autotune), then
diagnosis-driven techniques. Each variant gets a versioned file + lab
notebook tracking hypothesis/changes/results/learnings. Stops on budget
exhaustion, diminishing returns, or target met
- **Stage 3 (Finalizer):** Applies winner in-place, runs full test suite
(hard gate), generates 3-way comparison plots (original vs optimized vs
baseline), updates benchmark CSV, creates PR with descriptive body

## Key design decisions

- **No @triton.autotune**: Incompatible with Liger's forward-backward
ctx coupling pattern and NPU backends. Uses manual parameter sweeps
instead.
- **Full benchmark suite every iteration**: No lightweight shortcuts --
ensures optimization is good across ALL input sizes, not just
cherry-picked ones.
- **Balanced guardrails**: Rejects variants that regress non-target
metric >5% or regress one pass >10% to improve the other.
- **Comment preservation**: All existing comments kept, explanatory
comments added for every optimization change.
- **Autonomous + interactive modes**: Runs end-to-end when told "just
optimize it", or pauses at human checkpoints between stages.

## Files (7 files, ~2,100 lines)

| File | Lines | Purpose |
|------|-------|---------|
| `SKILL.md` | 116 | Orchestration, input parsing, pipeline flow,
guardrails |
| `profiler.md` | 156 | Stage 1: baseline benchmarks, GPU detection,
bottleneck diagnosis |
| `optimizer.md` | 395 | Stage 2: optimization loop with accumulated
learning |
| `finalizer.md` | 417 | Stage 3: apply, test, plot, update CSV, create
PR, report |
| `optimization-strategies.md` | 794 | Technique catalog: parameter
tuning, memory-bound, compute-bound, architecture-specific |
| `templates/optimization-profile.md` | 140 | Cross-stage contract
between Profiler and Optimizer |
| `templates/variant-notes.md` | 70 | Per-variant lab notebook format
for learning accumulation |

## Test plan

- [x] Skill triggers correctly on "optimize the X kernel" prompts
- [x] Tested end-to-end on `fused_add_rms_norm` (PR #1187)
- [x] Verified against Claude Code skill best practices (conciseness,
progressive disclosure, SKILL.md under 500 lines)

---------

Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Author: vaibhavjindal

.claude/skills/liger-kernel-perf/SKILL.md

@@ -0,0 +1,116 @@
+---
+name: liger-kernel-perf
+description: "Optimizes the performance of existing Liger Kernel Triton kernels. Profiles kernels, diagnoses bottlenecks (memory-bound vs compute-bound), generates multiple optimization variants with benchmarking, and applies the best variant while maintaining correctness. Supports GPU architecture-specific optimization (Ampere, Hopper, Blackwell). Use when a user asks to optimize, speed up, tune, profile, or reduce memory of an existing Liger kernel."
+---
+
+# Liger Kernel Perf
+
+Optimizes existing Liger Kernel Triton kernels through a 3-stage pipeline: Profile, Optimize, Finalize. Supports interactive mode (human checkpoints between stages) and autonomous mode (runs end-to-end). NVIDIA GPUs only.
+
+## Mode Detection
+
+- **Interactive mode** (default): Human checkpoints between each stage
+- **Autonomous mode**: User says "just optimize it", "run without asking me", "optimize autonomously" → all stages run end-to-end, user sees only the final report
+
+## Input Parsing
+
+Extract from the user's request:
+
+| Field | Description | Default |
+|-------|-------------|---------|
+| `target_kernel` | Which kernel to optimize (e.g., "rms_norm", "cross_entropy") | **Required** |
+| `optimization_goal` | speed / memory / balanced | balanced |
+| `scope` | Specific pass (forward/backward), input regime, or general | general |
+| `target_gpu` | Ampere / Hopper / Blackwell / auto-detect | auto-detect |
+| `autonomy` | interactive / autonomous | interactive |
+| `max_variants` | Max optimization variants to try | 8 |
+| `target_metric` | Optional concrete target (e.g., "forward under 0.3ms at hidden_size=4096") | none |
+
+## Pre-Flight Validation
+
+Before starting the pipeline, validate:
+
+1. Kernel file exists: `src/liger_kernel/ops/{kernel}.py`
+2. Benchmark script exists: `benchmark/scripts/benchmark_{kernel}.py`
+3. Test file exists: `test/transformers/test_{kernel}.py`
+4. GPU is available and CUDA works
+5. Project is installed in dev mode (`pip install -e ".[dev]"`)
+
+If any validation fails, report clearly and stop.
+
+## Pipeline
+
+### Stage 1: Profile
+
+Spawn a **Profiler** agent (read [profiler.md](profiler.md)).
+
+The agent:
+1. Creates the workspace directory `optimization/{kernel}/`
+2. Copies the original kernel as a snapshot
+3. Runs baseline benchmarks using the existing benchmark script
+4. Detects GPU architecture (or uses user-specified target)
+5. Optionally runs NCU profiling (if `ncu` is available)
+6. Analyzes the kernel code (tier classification, patterns, optimization opportunities)
+7. Classifies the bottleneck: memory-bound vs compute-bound
+8. Produces an optimization profile with a recommended strategy order
+9. Saves profile to `optimization/{kernel}/profile.md`
+
+**Human checkpoint (interactive mode):** Present the optimization profile with bottleneck diagnosis and proposed strategy order. Confirm before proceeding.
+
+### Stage 2: Optimize
+
+Spawn an **Optimizer** agent (read [optimizer.md](optimizer.md)).
+
+The agent runs an autonomous optimization loop:
+
+1. Read the optimization profile and original kernel
+2. **Always try parameter tuning first** (BLOCK_SIZE, num_warps, num_stages manual sweep -- NOT @triton.autotune)
+3. Then apply diagnosis-driven techniques from [optimization-strategies.md](optimization-strategies.md)
+4. For each variant:
+   a. Generate the variant code → `optimization/{kernel}/{kernel}_vN.py`
+   b. Write the variant lab notebook → `optimization/{kernel}/{kernel}_vN_notes.md`
+   c. Run quick smoke test (single shape, float32, forward+backward) → discard on failure
+   d. Run the **full existing benchmark script** → `optimization/{kernel}/benchmarks/vN_results.csv`
+   e. Check guardrails (no catastrophic regressions)
+   f. Update the variant notes with actual results
+5. Read all prior variant notes before generating the next variant
+6. **Stop when:** budget exhausted, 2 consecutive variants with <1% improvement, or target metric met
+7. Produce a comparison table of ALL variants
+
+**Human checkpoint (interactive mode):** Present the comparison table across all variants. User approves the winner (or skill picks best if autonomous).
+
+### Stage 3: Finalize
+
+Spawn a **Finalizer** agent (read [finalizer.md](finalizer.md)).
+
+The agent:
+1. Applies the winning variant in-place to `src/liger_kernel/ops/{kernel}.py`
+2. Runs the full test suite: `python -m pytest test/transformers/test_{kernel}.py -xvs` (hard gate)
+3. Runs checkstyle: `make checkstyle` (auto-fix with `ruff check . --fix && ruff format .`)
+4. Generates 3-way comparison plots (original liger vs optimized liger vs huggingface baseline) using `benchmarks_visualizer.py`
+5. Generates the final optimization report → `optimization/{kernel}/report.md`
+6. Creates a PR with only the kernel code changes (no plots or optimization workspace files)
+7. Presents the before/after summary with plots
+
+**Human checkpoint (interactive mode):** Present the final report with before/after numbers, comparison plots, and test results.
+
+## Guardrails
+
+These apply to EVERY variant, regardless of mode:
+
+| Guardrail | Threshold | Action |
+|-----------|-----------|--------|
+| Non-target metric regression | >5% worse | Reject variant |
+| Cross-pass regression | >10% on one pass to marginally improve other | Reject variant |
+| Smoke test failure | Any correctness failure | Discard variant immediately |
+| Full test suite failure | Any | Do NOT apply winner, report failure, stop |
+| Checkstyle failure | Any | Auto-fix with ruff, retry once |
+
+## Reference Files
+
+- [profiler.md](profiler.md) -- Profiler Agent specification
+- [optimizer.md](optimizer.md) -- Optimizer Agent specification
+- [finalizer.md](finalizer.md) -- Finalizer Agent specification
+- [optimization-strategies.md](optimization-strategies.md) -- Catalog of optimization techniques
+- [templates/optimization-profile.md](templates/optimization-profile.md) -- Profiling output format (cross-stage contract)
+- [templates/variant-notes.md](templates/variant-notes.md) -- Per-variant lab notebook format