xpu-kernels

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 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
 

Supported Hardware

See the 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.

Workflow Steps

Step 1: Analyze

• Read the baseline source file. Identify shapes, dtypes, operations, fusion opportunities.
• If baseline is PyTorch: run python scripts/analyze_kernel.py <pytorch_file>.
• If baseline is Triton (--triton-baseline): skip analyze_kernel.py (it only supports PyTorch). Read the Triton file directly.
• Read relevant knowledge base files: start with references/correctness.yaml and references/xpu_optimizations.yaml.
• Read references/implementation_reference.md for templates and the Model class pattern.

Step 2: Initialize

Step 3: Trial Loop (always run all max_trials from config.yaml)

For each trial:

1. Write kernel — start from templates or modify previous trial. See references/implementation_reference.md.
2. Validate — python scripts/validate_triton.py <triton_file> (fix until passing; doesn't count as a trial).
3. Save — python scripts/trial_manager.py save <kernel_name> <triton_file> --parent <parent_id> --strategy "description". Omit --parent for the first trial (t0).
4. Benchmark (MANDATORY every trial):
   • Trial t0: python scripts/benchmark.py <baseline_file> <triton_file> [--triton-baseline] (measures both baseline and triton).
   • Trials t1+: Get cached baseline via python scripts/trial_manager.py baseline-us <kernel_name>, then run python scripts/benchmark.py <baseline_file> <triton_file> [--triton-baseline] --baseline-us <cached_value> (skips baseline perf, saves time).
   • After finalize: Re-run benchmark.py without --baseline-us for final accurate comparison.
5. Record — python scripts/trial_manager.py result <kernel_name> <trial_id> --validation pass --correctness <pass|fail> --speedup <float> --baseline_us <float> --triton_us <float> (runtimes from benchmark output).
6. Profile (MANDATORY after t1, if vtune_enabled is true in config.yaml) — Run python scripts/xpu_profiler.py <triton_file> after your first benchmarked trial. Use its output to guide subsequent trial strategies. Run again if speedup plateaus after 2+ additional trials. Skip this step entirely if vtune_enabled is false.
7. Decide next action (use profiler output from step 6 to inform decisions):
   • Speedup > 5x → stop (excellent), finalize
   • Speedup improved → continue on this branch, try next optimization level
   • Speedup regressed → branch back to best trial, try different strategy
   • Correctness failed → fix on same branch
   • Profiler says low occupancy (if vtune_enabled) → increase tile sizes, check references/xpu_optimizations.yaml
   • Profiler says overhead kernels dominate (if vtune_enabled) → pre-pack to bf16, see references/optimization_levels.yaml
   • Plateau → do NOT stop. Try a fundamentally different approach (different algorithm, tiling, fusion strategy). LLM sampling can discover new ideas.
   • See references/optimization_strategies.md for the full "try harder" decision tree

Step 4: Finalize

Reference Docs — Read During Step 1

Existing Baselines Are Naive

The test_kernels/*.py Triton files (non-pytorch) are unoptimized baselines. They use manual pointer arithmetic, lack autotune, and miss XPU optimizations. Do NOT copy their patterns. Use references/implementation_reference.md instead.

Core XPU Kernel Patterns

Tensor Descriptors (Preferred on XPU)

Tensor descriptors produce better address generation and memory access codegen than block pointers on Intel XPU.

GRF Mode '256'

Use the large register file for compute-heavy kernels:

Register-Spill-Aware Autotune Pruning

For autotuned XPU kernels, prune configs that are likely to spill even after large-GRF recompilation. This avoids slow double-compile/benchmark cycles where the Intel backend reports spills, recompiles in large GRF mode, and still spills.

Use an early_config_prune that estimates per-thread GRF pressure from the live accumulator and operand tiles:

Guidelines:

• Apply this only on XPU; leave other backends' configs unchanged.
• Read BM, BN, and BK from either autotune config kwargs or launch args.
• Count accumulator tiles explicitly: gate/up fused kernels often have 2 accumulators, down/simple GEMM usually has 1.
• Count weight tiles explicitly: gate/up fused kernels often load 2 weight tiles, down/simple GEMM usually loads 1.
• Use num_stages for pipelined MMA paths; use 1 for scalar/non-pipelined paths.
• Keep a never-empty fallback, such as the config with the lowest estimated register pressure.
• Treat the model as a coarse filter for catastrophic spillers, not a perfect predictor; leave borderline configs for real benchmarking.

This complements grf_mode='256': large GRF increases the register budget, while pruning removes configs that are still too register-heavy to be useful.

For an implementation pattern, see references/xpu_optimizations.yaml (xpu_register_spill_autotune_pruning).

Tile Swizzling

Use 1D grid with GROUP_SIZE_M for L2 locality:

bf16 Inputs with fp32 Accumulation

Critical XPU Constraints

• NO default values for @triton.autotune meta-parameters in kernel signature
• 1D grid when using tile swizzling (GROUP_SIZE_M)
• boundary_check uses dimension indices (0, 1), not booleans
• Cast batch indices to int64 before stride multiplication
• Prefer tensor descriptors over block pointers for all new XPU kernels
• Do NOT mix block pointer and tensor descriptor APIs on same operation
• Pre-zero output buffers when using atomic accumulation
• Model class must be compatible with ai-bench (nn.Module with nn.Linear)
• Match get_inputs(), get_init_inputs(), and module-level constants from *_pytorch.py

Full constraint list: correctness.yaml

Performance Results

Measured on Intel Battlemage G21 / Arc Pro B50 (128 XVEs). All runtimes are median of benchmark trials.

KernelBench Level 2 — Fused Kernels (bf16)

Speedup is vs. PyTorch eager baseline. Includes GEMM+Sigmoid+Scaling, GEMM+GELU+Softmax, Conv+BatchNorm+ReLU, and other fused patterns.

Flash Attention Forward (fp16)

Baseline is the flash attention kernel from the Intel XPU Triton backend; speedup is vs. that kernel across multiple sequence lengths.

Full results: see the Xe-Forge repository.

Common Issues

Project Structure

See Also

Xe-Forge Tools

• analyze_kernel.py — Static analysis of PyTorch reference
• validate_triton.py — Pre-benchmark constraint checks
• benchmark.py — Correctness + performance measurement
• xpu_profiler.py — VTune GPU hardware counters
• trial_manager.py — Branching trial tree management

XPU Optimization References

• correctness.yaml — Critical constraints
• xpu_optimizations.yaml — Tensor descriptors, GRF, swizzling
• optimization_strategies.md — Strategy reference
• optimization_levels.yaml — Progressive L1-L5 levels
• implementation_reference.md — Code templates

Other References

• kernelbench-classification.md — KernelBench operator taxonomy

External Resources

• Xe-Forge Repository
• AI-Bench — Benchmark harness for correctness + performance
• Intel XPU Backend for Triton
• Triton Language Guide