skills.rest
huggingface
huggingfaceOFFICIAL
Explore all skills by this author →
Website𝕏@huggingfaceGitHub

xpu-kernels

Generates and optimizes Triton kernels for Intel XPU GPUs through an iterative benchmark-driven workflow.

736|125|Updated Nov 29, 2024
Search Tags:#kernel optimization#pyTorch#GEMM#Triton#gpu benchmarking#vtune profiling#intel xpu
One-click install
npx skills add https://github.com/huggingface/kernels --skill xpu-kernels
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: xpu-kernels
Source: https://github.com/huggingface/kernels/tree/main/kernel-builder/skills/xpu-kernels
Command: npx skills add https://github.com/huggingface/kernels --skill xpu-kernels

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill requires torch, triton, yaml, ai-bench, and includes scripts (resource) and references (resource) components.

What problem does it solve?

Writing fast Triton kernels for Intel XPU GPUs requires deep knowledge of XPU-specific patterns like tensor descriptors, GRF mode, and tile swizzling, plus repeated manual benchmarking to find the best configuration. This Skill automates that loop: it analyzes a PyTorch baseline, generates and validates Triton kernel variants, benchmarks each on XPU hardware, profiles with VTune, and tracks a branching trial tree until the fastest correct kernel is found.

Core Features & Use Cases

  • Xe-Forge Optimization Loop: A structured analyze → validate → benchmark → profile → finalize workflow driven by CLI tools (analyze_kernel.py, validate_triton.py, benchmark.py, xpu_profiler.py, trial_manager.py) with a configurable max_trials budget.
  • XPU-Specific Kernel Patterns: A knowledge base covering tensor descriptors, GRF 256 mode, register-spill-aware autotune pruning, tile swizzling, bf16 inputs with fp32 accumulation, and fusion rules for GEMM, reductions, and Flash Attention.
  • Correctness and Performance Verification: Benchmarks kernels against PyTorch or Triton baselines via ai-bench, checking numerical equivalence and reporting per-variant speedups.
  • Use Case: Given a KernelBench-style fused kernel like GEMM+Sigmoid+Scaling+ResidualAdd in PyTorch, run the trial loop to produce a numerically equivalent Triton kernel that is measurably faster on an Intel Battlemage G21 / Arc Pro B50 GPU.

Quick Start

Optimize my PyTorch GEMM baseline file into a faster Triton kernel for Intel XPU using the analyze, validate, benchmark, and finalize workflow.

Frequently Asked Questions about xpu-kernels

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

FAQPage Schema
How do I optimize a PyTorch kernel into a Triton kernel for Intel XPU?▼

Run the Xe-Forge workflow: analyze the PyTorch baseline with analyze_kernel.py, initialize trial tracking with trial_manager.py, then iterate by writing Triton variants, validating with validate_triton.py, and benchmarking with benchmark.py. Finalize the best trial once speedup targets are met.

How do I benchmark a Triton kernel against a PyTorch baseline on XPU?▼

Use benchmark.py with the baseline file and Triton file as arguments. It checks numerical correctness via ai-bench and reports per-variant runtimes and speedup, with an option to reuse a cached baseline time via --baseline-us.

What Intel GPUs does Triton XPU kernel development support?▼

The workflow is verified on Intel Battlemage G21 / Arc Pro B50 (Xe2 architecture, 128 XVEs, roughly 500 GB/s memory bandwidth). Other hardware supported by the Intel XPU Backend for Triton may work but is not verified.

Why does my Triton kernel fail with conflicting meta-parameters on autotune?▼

Autotune meta-parameters like BLOCK_M must not have default values in the kernel signature when using @triton.autotune. Remove the defaults and let the autotune configs supply the values.

Should I use tensor descriptors or block pointers for XPU Triton kernels?▼

Tensor descriptors are preferred on Intel XPU because they produce better address generation and memory access codegen than block pointers. Do not mix both APIs on the same operation, and note that descriptor loads handle bounds internally without a boundary_check argument.

When should I stop the kernel optimization trial loop?▼

Run all trials defined by max_trials in config.yaml; the only valid early stop is a speedup greater than 5x. Do not stop on plateaus, since later trials can still discover better strategies.

Related Skills

Explore More Software Engineering Skills→
Just-itJust-it
33

reference-generation

Generate a reference PyTorch nn.Module from operator description JSON.

Community
Advanced
dongg622dongg622

mctriton

Compile and run Triton GPU kernels on曦云GPU systems.

Community
Intermediate
AscendAscend
31

triton-operator-precision-eval

Benchmark Triton operator precision against PyTorch references and generate MERE/MARE reports.

Official
Advanced
mindspore-aimindspore-ai
258

triton-cuda-attention

Implement Triton-CUDA attention kernels with online Softmax for transformer workloads.

Official
Advanced