/xpu-kernels
Provides guidance for writing, optimizing, and benchmarking Triton kernels for Intel XPU GPUs (Battlemage/Arc Pro B50) using the Xe-Forge optimization framework. Includes an LLM-driven trial-loop workflow (analyze, validate, benchmark, profile, finalize), XPU-specific patterns
$ npx -y skills add huggingface/kernels --skill xpu-kernels --agent claude-codeHow it fires
How this skill gets triggered: by you, by Claude, or both.
- Fires itselfAuto-invocation. Claude auto-loads it when your prompt matches the work.Auto-invocation is when the right skill fires by itself at the right moment, driven by a FLOW.md router and a hook, instead of you invoking it by name. It is the difference between a skill being installed and a skill actually getting used.Read the full definition →
- You can call itInvoke it directly when you want it.
- Slash command
/xpu-kernels
Context preview
The summary Claude sees to decide when to auto-load this skill.
Provides guidance for writing, optimizing, and benchmarking Triton kernels for Intel XPU GPUs (Battlemage/Arc Pro B50) using the Xe-Forge optimization framework. Includes an LLM-driven trial-loop workflow (analyze, validate, benchmark, profile, finalize), XPU-specific patterns
SKILL.md
xpu-kernels.SKILL.mdname: xpu-kernels
description: "Provides guidance for writing, optimizing, and benchmarking Triton kernels for Intel XPU GPUs (Battlemage/Arc Pro B50) using the Xe-Forge optimization framework. Includes an LLM-driven trial-loop workflow (analyze, validate, benchmark, profile, finalize), XPU-specific patterns (tensor descriptors, GRF mode, tile swizzling), KernelBench fused kernels, and Flash Attention."
disable-model-invocation: false
user-invocable: true
allowed-tools: "Read, Grep, Glob, Bash"
argument-hint: "kernel type: gemm, reduction, flash-attention, optimize, benchmark, tensor-descriptors, xe-forge"
XPU Triton Kernels for Intel GPUs
This skill provides patterns and guidance for developing optimized Triton kernels targeting Intel XPU GPUs (Battlemage/Arc Pro B50). It integrates the [Xe-Forge](https://github.com/IntelLabs/Xe-Forge) optimization framework — an LLM-driven loop that transforms PyTorch code into fast Triton kernels.
Quick Start
Optimize a Kernel (Xe-Forge Workflow)
The full optimization workflow analyzes a PyTorch baseline, generates Triton kernel variants in a branching trial tree, benchmarks each on XPU hardware, and finalizes the best result.
# 1. Analyze the baseline
python scripts/analyze_kernel.py test_kernels/70_Gemm_Sigmoid_Scaling_ResidualAdd_pytorch.py
# 2. Initialize trial tracking
python scripts/trial_manager.py init 70_Gemm_Sigmoid test_kernels/70_Gemm_Sigmoid_Scaling_ResidualAdd_pytorch.py
# 3. Validate a generated kernel (no GPU needed)
python scripts/validate_triton.py my_kernel.py
# 4. Benchmark correctness + performance
python scripts/benchmark.py test_kernels/70_Gemm_Sigmoid_Scaling_ResidualAdd_pytorch.py my_kernel.py
# 5. Profile with VTune (optional)
python scripts/xpu_profiler.py my_kernel.py
# 6. Finalize best trial
python scripts/trial_manager.py finalize 70_Gemm_Sigmoid optimized_triton.py
Supported Hardware
| GPU | Architecture | XVEs | Mem BW | Key Feature | Verified | |-----|-------------|------|--------|-------------|:--------:| | **Battlemage G21 / Arc Pro B50** | Xe2 | 128 | ~500 GB/s | Tensor descriptors, GRF 256 | Yes |
> See the [Intel XPU Backend for Triton](https://github.com/intel/intel-xpu-backend-for-triton) for supported hardware.
When This Skill Applies
Use this skill when:
- Optimizing PyTorch operations into Triton kernels for **Intel XPU**
- Writing GEMM, fused kernels, reductions, or Flash Attention for Intel GPUs
- Running the **Xe-Forge optimization loop** (analyze → validate → benchmark → profile → finalize)
- Benchmarking kernel performance against PyTorch baseline on XPU
Xe-Forge Optimization Workflow
Transform PyTorch code into optimized Triton kernels for Intel XPU. Kernels must be numerically equivalent and faster than baseline.
Configuration — Read `config.yaml` first
At the start of every session, read `scripts/config.yaml`. It controls:
- **`max_trials`** — hard cap on optimization trials; always run all of them (use this instead of hardcoded "10")
- **`vtune_enabled`** — if `false`, skip ALL VTune profiling steps (Step 3.6 and profiler-related decisions)
- **`vtune_bin`** — path to the VTune binary (also settable via `VTUNE_BIN` env var)
Rules — Never Violate
1. **ONLY create** Triton kernel files (`test_kernels/*_triton.py` or trial files `t<trial_id>.py`). 2. **NEVER create** benchmark scripts, test scripts, helper utilities, or any other Python files. 3. **NEVER write custom scripts** to measure performance or test correctness — ONLY use `scripts/benchmark.py`. 4. If a tool fails, **STOP and report the error**. Do NOT work around it with custom scripts. 5. Generated kernels must be **self-contained** — all helper functions inline. 6. You **MUST run all `max_trials` trials** from `config.yaml`. Do NOT stop early due to plateau — LLM sampling can discover new ideas at any point. The only valid early stop is speedup > 5x.
Mandatory Tools
**CRITICAL — Single-XPU serialization**: There is only ONE XPU on this machine. You MUST NOT run multiple GPU workloads in parallel. `benchmark.py` and `xpu_profiler.py` must execute strictly one at a time — concurrent GPU jobs produce wrong results. CPU-only tools (`analyze_kernel.py`, `validate_triton.py`, `trial_manager.py`) are safe to parallelize with each other and with anything else.
| Tool | Command | Purpose | |------|---------|---------| | **Analyze** | `python scripts/analyze_kernel.py <file>` | Static analysis: operations, shapes, fusion opportunities | | **Validate** | `python scripts/validate_triton.py <file>` | Syntax + constraint checks before GPU time | | **Benchmark** | `python scripts/benchmark.py <baseline> <triton> [--triton-baseline] [--baseline-us <cached>]` | Correctness + performance via ai-bench | | **Profile** | `python scripts/xpu_profiler.py <file>` | VTune GPU hardware counters + recommendations | | **Init trials** | `python scripts/trial_manager.py init <kernel_name> <baseline_file> [--triton-baseline]` | Initialize trial tracking | | **Save trial** | `python scripts/trial_manager.py save <kernel_name> <file> [--parent <parent_id>] [--strategy "..."]` | Save trial to tree | | **Record result** | `python scripts/trial_manager.py result <kernel_name> <trial_id> --validation pass --correctness <pass\|fail> --speedup <float> --baseline_us <float> --triton_us <float>` | Record benchmark result | | **Check status** | `python scripts/trial_manager.py status <kernel_name>` | View trial tree | | **Best trial** | `python scripts/trial_manager.py best <kernel_name>` | Get best trial | | **Baseline time** | `python scripts/trial_manager.py baseline-us <kernel_name>` | Cached baseline time for `--baseline-us` | | **Finalize** | `python scripts/trial_manager.py finalize <kernel_name> <name>_triton.py` | Copy best trial to output |
Workflow Steps
Step 1: Analyze
- Read the baseline source file. Identify shapes, dtypes, operations, fusion opportunities.
- If baseline is PyTorch: ru
Read more
name: xpu-kernels description: "Provides guidance for writing, optimizing, and benchmarking Triton kernels for Intel XPU GPUs (Battlemage/Arc Pro B50) using the Xe-Forge optimization framework. Includes an LLM-driven trial-loop workflow (analyze, validate, benchmark, profile, finalize), XPU-specific patterns (tensor descriptors, GRF mode, tile swizzling), KernelBench fused kernels, and Flash Attention." disable-model-invocation: false user-invocable: true allowed-tools: "Read, Grep, Glob, Bash" argument-hint: "kernel type: gemm, reduction, flash-attention, optimize, benchmark, tensor-descriptors, xe-forge"
XPU Triton Kernels for Intel GPUs
This skill provides patterns and guidance for developing optimized Triton kernels targeting Intel XPU GPUs (Battlemage/Arc Pro B50). It integrates the [Xe-Forge](https://github.com/IntelLabs/Xe-Forge) optimization framework — an LLM-driven loop that transforms PyTorch code into fast Triton kernels.
Quick Start
Optimize a Kernel (Xe-Forge Workflow)
The full optimization workflow analyzes a PyTorch baseline, generates Triton kernel variants in a branching trial tree, benchmarks each on XPU hardware, and finalizes the best result.
# 1. Analyze the baseline python scripts/analyze_kernel.py test_kernels/70_Gemm_Sigmoid_Scaling_ResidualAdd_pytorch.py # 2. Initialize trial tracking python scripts/trial_manager.py init 70_Gemm_Sigmoid test_kernels/70_Gemm_Sigmoid_Scaling_ResidualAdd_pytorch.py # 3. Validate a generated kernel (no GPU needed) python scripts/validate_triton.py my_kernel.py # 4. Benchmark correctness + performance python scripts/benchmark.py test_kernels/70_Gemm_Sigmoid_Scaling_ResidualAdd_pytorch.py my_kernel.py # 5. Profile with VTune (optional) python scripts/xpu_profiler.py my_kernel.py # 6. Finalize best trial python scripts/trial_manager.py finalize 70_Gemm_Sigmoid optimized_triton.py
Supported Hardware
| GPU | Architecture | XVEs | Mem BW | Key Feature | Verified | |-----|-------------|------|--------|-------------|:--------:| | **Battlemage G21 / Arc Pro B50** | Xe2 | 128 | ~500 GB/s | Tensor descriptors, GRF 256 | Yes |
> See the [Intel XPU Backend for Triton](https://github.com/intel/intel-xpu-backend-for-triton) for supported hardware.
When This Skill Applies
Use this skill when:
- Optimizing PyTorch operations into Triton kernels for **Intel XPU**
- Writing GEMM, fused kernels, reductions, or Flash Attention for Intel GPUs
- Running the **Xe-Forge optimization loop** (analyze → validate → benchmark → profile → finalize)
- Benchmarking kernel performance against PyTorch baseline on XPU
Xe-Forge Optimization Workflow
Transform PyTorch code into optimized Triton kernels for Intel XPU. Kernels must be numerically equivalent and faster than baseline.
Configuration — Read `config.yaml` first
At the start of every session, read `scripts/config.yaml`. It controls:
- **`max_trials`** — hard cap on optimization trials; always run all of them (use this instead of hardcoded "10")
- **`vtune_enabled`** — if `false`, skip ALL VTune profiling steps (Step 3.6 and profiler-related decisions)
- **`vtune_bin`** — path to the VTune binary (also settable via `VTUNE_BIN` env var)
Rules — Never Violate
1. **ONLY create** Triton kernel files (`test_kernels/*_triton.py` or trial files `t<trial_id>.py`). 2. **NEVER create** benchmark scripts, test scripts, helper utilities, or any other Python files. 3. **NEVER write custom scripts** to measure performance or test correctness — ONLY use `scripts/benchmark.py`. 4. If a tool fails, **STOP and report the error**. Do NOT work around it with custom scripts. 5. Generated kernels must be **self-contained** — all helper functions inline. 6. You **MUST run all `max_trials` trials** from `config.yaml`. Do NOT stop early due to plateau — LLM sampling can discover new ideas at any point. The only valid early stop is speedup > 5x.
Mandatory Tools
**CRITICAL — Single-XPU serialization**: There is only ONE XPU on this machine. You MUST NOT run multiple GPU workloads in parallel. `benchmark.py` and `xpu_profiler.py` must execute strictly one at a time — concurrent GPU jobs produce wrong results. CPU-only tools (`analyze_kernel.py`, `validate_triton.py`, `trial_manager.py`) are safe to parallelize with each other and with anything else.
| Tool | Command | Purpose | |------|---------|---------| | **Analyze** | `python scripts/analyze_kernel.py <file>` | Static analysis: operations, shapes, fusion opportunities | | **Validate** | `python scripts/validate_triton.py <file>` | Syntax + constraint checks before GPU time | | **Benchmark** | `python scripts/benchmark.py <baseline> <triton> [--triton-baseline] [--baseline-us <cached>]` | Correctness + performance via ai-bench | | **Profile** | `python scripts/xpu_profiler.py <file>` | VTune GPU hardware counters + recommendations | | **Init trials** | `python scripts/trial_manager.py init <kernel_name> <baseline_file> [--triton-baseline]` | Initialize trial tracking | | **Save trial** | `python scripts/trial_manager.py save <kernel_name> <file> [--parent <parent_id>] [--strategy "..."]` | Save trial to tree | | **Record result** | `python scripts/trial_manager.py result <kernel_name> <trial_id> --validation pass --correctness <pass\|fail> --speedup <float> --baseline_us <float> --triton_us <float>` | Record benchmark result | | **Check status** | `python scripts/trial_manager.py status <kernel_name>` | View trial tree | | **Best trial** | `python scripts/trial_manager.py best <kernel_name>` | Get best trial | | **Baseline time** | `python scripts/trial_manager.py baseline-us <kernel_name>` | Cached baseline time for `--baseline-us` | | **Finalize** | `python scripts/trial_manager.py finalize <kernel_name> <name>_triton.py` | Copy best trial to output |
Workflow Steps
Step 1: Analyze
- Read the baseline source file. Identify shapes, dtypes, operations, fusion opportunities.
- If baseline is PyTorch: ru
The Kernel Hub allows Python libraries and applications to load compute kernels directly from the Hub.
Repo: huggingface/kernels
Other skills on kernels.
- /cpu-kernels
Provides guidance for writing, optimizing, and benchmarking C++ CPU kernels with SIMD intrinsics (AVX2/AVX512) for the Hugging Face kernels ecosystem. Includes a two-phase workflow: Phase 1 correctness (generic → AVX2) and Phase 2 performance exploration (AVX512 with branching
Open skill - /cuda-kernels
Provides guidance for writing and benchmarking optimized CUDA kernels for NVIDIA GPUs (H100, A100, T4) targeting HuggingFace diffusers and transformers libraries. Kernels must be kernel-builder/ABI3-compliant: no pybind11, no setup.py, TORCH_LIBRARY_EXPAND bindings only.
Open skill - /rocm-kernels
Provides guidance for writing and benchmarking optimized Triton kernels for AMD GPUs (MI355X, R9700) on ROCm, targeting HuggingFace diffusers (LTX-Video, SD3, FLUX) and transformers. Core kernels: RMSNorm, RoPE 3D, GEGLU, AdaLN. Includes XCD swizzle, autotune, diffusers
Open skill - /triton-kernels
| name | triton-kernels | | --- | --- | | description | Provides guidance for writing and benchmarking portable Triton kernels targeting NVIDIA and AMD GPUs. Covers core DSL patterns, @triton.autotune, numerics (fp16/bf16/fp8), masked loads, reductions, tiling, benchmarking
Open skill

