Xe-Forge: Optimizing Triton Kernels for Intel GPUs

• Xe-Forge is a hardware-aware LLM-driven system designed to optimize existing Triton kernels for Intel GPUs by preserving semantics while accelerating performance.
• It employs a multi-stage pipeline—including analysis, fusion, dtype fixes, and GPU-specific tuning—with real-device verification to ensure correctness and efficiency.
• The system significantly reduces repetitive manual tuning of deep learning kernels, achieving notable speedups and adapting to Intel Xe-class architectural constraints.

Xe-Forge is a hardware-aware, LLM-driven system for optimizing existing Triton kernels for Intel GPUs, rather than generating kernels from scratch. It is framed as a distinct “Triton-to-Triton” optimization problem: the input is already a correct Triton kernel, and the goal is to preserve semantics while making it fast on Intel GPU hardware. The system addresses the repetitive porting and tuning work required when adapting deep learning kernels to Intel Xe-class devices, including tile and warp selection, memory-access restructuring, pointer modernization, operator fusion, persistent-kernel conversion, and Intel-specific architectural constraints that are generally absent from generic LLM pretraining data (Spoczynski et al., 16 Apr 2026).

1. Definition and problem setting

Xe-Forge targets the manual optimization cycle that typically recurs across every kernel in a Triton code base: quantization and dtype cleanup, memory coalescing, tile and block-size tuning, architecture-specific workarounds, fusion, and autotuning. The paper presents this as a bottleneck for deployment on new accelerators, especially because many available LLM kernel heuristics are NVIDIA-centric, while Intel GPUs require different defaults and obey different backend constraints. In that setting, Xe-Forge takes as input a functionally correct Triton kernel, optionally a PyTorch reference implementation, and problem context such as input shapes, FLOP count, target dtype, benchmark specification, and target GPU configuration from runtime hardware query (Spoczynski et al., 16 Apr 2026).

The system’s output is an optimized Triton kernel, or the original code if optimization fails, subject to four conditions: the candidate must parse, satisfy structural constraints, produce outputs matching the original or reference within tolerance, and run faster than the original on real Intel GPU hardware. This makes Xe-Forge conservative by construction: it does not replace a working kernel with an unverified rewrite.

The Intel-specific difficulty is not limited to parameter retuning. The Triton path on Intel lowers through TTIR $\to$ TTGIR $\to$ LLVM $\to$ SPIR-V $\to$ Intel Graphics Compiler, and kernels must respect constraints involving warp counts, GRF mode, SLM capacity, tile-size interactions with register pressure, subgroup and workgroup behavior, and backend lowering limitations. The paper explicitly notes that generic LLM defaults such as num_warps=8 or 128x128 tiles can leave substantial performance on the table on Intel GPU, whereas Xe-Forge uses Intel-aware defaults such as 32 warps for many workloads, large GRF mode, shape-aware tile sizing, and GROUP_SIZE_M swizzling.

A practical implication is that Xe-Forge is not merely a code-rewriting assistant. It is a system for architecture-constrained search over semantics-preserving kernel variants, with validation performed on the target device itself.

2. Multi-stage pipeline and optimization stages

The Xe-Forge pipeline begins with an analysis stage and proceeds through up to nine optimization stages: algorithmic optimization, discovery, dtype fix, fusion, memory access, block pointers, persistent kernel, GPU-specific tuning, and autotuning. Rather than using a fixed sequence, an LLM-based planner chooses an ordered subset of stages based on detected issues, subject to hard dependencies; if planning fails, the system falls back to the sequence Algorithmic $\to$ Discovery $\to$ Dtype $\to$ Fusion $\to$ Memory $\to$ Block Pointers $\to$ Persistent Kernel $\to$0 GPU-Specific $\to$1 Autotuning (Spoczynski et al., 16 Apr 2026).

Stage	Role
Analysis	Produces a structured issue inventory
Algorithmic optimization	Reduces FLOPs or memory accesses while preserving semantics
Discovery	Handles open-ended opportunities outside predefined categories
Dtype fix	Removes unsafe or inefficient precision choices
Fusion	Combines launches and collapses operator chains
Memory access	Improves coalescing, masking, and data movement
Block pointers	Rewrites manual address arithmetic with tl.make_block_ptr
Persistent kernel	Converts relaunchable work into resident tile iteration
GPU-specific tuning	Applies Intel-Xe-specific warp, GRF, tile, and swizzling choices
Autotuning	Emits a curated @triton.autotune configuration grid

The analysis stage is non-mutating. It produces typed issues, a severity score from 1–5, a textual description, a suggested fix, and an estimated speedup range. The issue taxonomy contains more than 30 categories. Representative examples include float64 accumulator, chainable pointwise ops, strided column loads, Hand-computed offsets, Relaunchable reduction, and Intel-specific issues such as num_warps != 32, no GRF annotation, and missing GROUP_SIZE_M.

The later stages are specialized. Algorithmic optimization handles mathematically equivalent restructurings such as common-subexpression elimination, loop-invariant hoisting, or tree reductions instead of serial accumulation. Discovery exists for cases that do not fit the predefined taxonomy; the paper gives the example of rewriting sum(x @ W.T, dim=1) as x @ W.sum(dim=0) to eliminate an $\to$2 GEMM, and also mentions caching weight statistics that would otherwise be recomputed each forward pass. Dtype fix converts inefficient types such as float64 to float32, uses float16 or bfloat16 where safe, and removes redundant casts. Fusion targets multi-launch chains, post-reduction pointwise operations, or beneficial replacement of vendor-library calls with fused Triton.

The remaining stages focus on implementation form. Memory access reorders loads and stores for coalescing, adds boundary masks to tl.load, reduces host-device synchronization, and lowers register pressure. Block pointers convert hand-computed offsets into tl.make_block_ptr with correct shape, strides, order, and boundary_check handling. Persistent kernel conversion makes a fixed number of thread blocks iterate across multiple tiles instead of launching one block per tile. GPU-specific tuning applies Intel-Xe-specific choices such as 32 warps, GRF mode selection, shape-aware tile sizing, GROUP_SIZE_M swizzling, cached packed transposes, and avoidance of repacking weights in the forward path. Autotuning then emits up to 12 curated configurations rather than brute-force enumeration.

Two structural features are important for the pipeline’s behavior. First, issue-driven skip logic means that stages with no mapped issues are omitted. Second, the kernel is re-analyzed after each stage, so indirectly resolved issues can disappear and newly introduced issues can be detected before later transformations.

3. CoVeR verification loop and knowledge architecture

At the core of each stage is CoVeR, the “Chain-of-Verification-and-Refinement” agent. CoVeR is implemented as a DSPy module that maintains a trajectory recording thoughts, tool invocations, observations, and error feedback. At each iteration, the LLM receives the task inputs together with the full trajectory, proposes a candidate rewrite, invokes verification tools, and either terminates on success or appends the new observations to the trajectory for another repair attempt. The experimental setting uses a maximum of $\to$3 iterations per stage (Spoczynski et al., 16 Apr 2026).

The principal tool is compile_and_verify, which imposes a four-level gate:

1. Syntax: parse the candidate with ast.parse.
2. Structure: check for required Triton imports, a @triton.jit kernel, a Model wrapper class, and hardware constraints such as power-of-two warp counts and block dimensions.
3. Correctness: instantiate original and candidate kernels with identical weights and inputs and compare outputs with torch.allclose.
4. Performance: benchmark optimized and original kernels on the target Intel GPU and accept only if the candidate is faster.

If all iterations fail, a fallback Chain-of-Thought extractor produces a best-effort candidate, which is independently verified; if that also fails, the stage returns the original code unchanged. This suggests a design bias toward guaranteed semantic safety rather than unconditional code mutation.

The knowledge base is a second central component. It is encoded as machine-readable YAML and contains 84 entries across six files: gpu_optimizations, memory_patterns, fusion_patterns, persistent_kernel, correctness, and dtype_optimizations. It stores three kinds of artifacts: constraints, transformation patterns, and full-code examples. Critical constraints are always injected into prompts; stage-specific patterns and matching full-code examples are added selectively.

The constraint layer encodes rules that generic model knowledge often misses. Examples include the facts that num_warps must be a power of two, block dimensions must be powers of two and at most 256, boundary_check for block pointers must be a tuple of integer dimension indices, launch grids should be one-dimensional when tile swizzling is enabled, and batch offsets derived from tl.program_id must be cast to int64 before stride multiplication. The knowledge base also encodes Intel-specific GRF and SLM guidance: Intel supports large GRF mode with 256 registers per thread and small GRF mode with 128 registers per thread; large mode is preferred for compute-bound kernels; and the hardware query system uses 64 KB GRF capacity in large mode and 32 KB in small mode.

The pattern layer provides before/after transformations with rationale, expected speedup range, and applicability tags. The most populated stages are memory access with 21 patterns and GPU-specific tuning with 20 patterns. The example layer contains nine complete unoptimized-to-optimized kernel pairs, including cases with kernel fusion, block pointers, large tiles, 32 warps, grf_mode='256', and GROUP_SIZE_M swizzling.

4. Hardware model, shape-aware tuning, and validation methodology

Xe-Forge queries the target GPU at runtime using torch.xpu.get_device_properties() and falls back to xpu-smi JSON parsing when needed. Retrieved properties include EU count, subslice count, slice count, maximum compute units, workgroup size, subgroup size, global memory capacity, local memory or SLM capacity, and support flags for FP16, BF16, and FP64. The detected GPU family then selects architecture-specific defaults (Spoczynski et al., 16 Apr 2026).

The shape-aware tuning logic uses problem dimensions such as $\to$4, $\to$5, and $\to$6. Tile sizes are clamped to the nearest power of two not exceeding the dimension; tall-skinny matrices favor larger BLOCK_M, while short-wide matrices favor larger BLOCK_N. The paper gives the tile-memory estimate as

$\to$7

If this estimate exceeds GRF capacity, BLOCK_K is reduced. The same logic supports SLM-bounded kernels such as Flash Attention and persistent kernels.

Several Intel-specific heuristics are built into this stage. Warp counts must be powers of two, and valid values explicitly include 1, 2, 4, 8, 16, and 32. The paper states that Intel GPU usually achieves peak occupancy at 32 warps for most workloads. GROUP_SIZE_M is set to 1 when fewer than 16 tiles make swizzling unhelpful; otherwise it targets roughly four groups per SM. GRF mode and pipeline stages scale with tile size and with the $\to$8 dimension.

Correctness and performance validation are equally hardware grounded. Correctness uses seeded initialization, explicit state_dict copying, cloned inputs, execution under torch.no_grad(), and torch.allclose with tolerances rtol = 10^{-2} and atol = 10^{-5}. NaNs are rejected immediately, unexpected Infs are rejected, and mismatch reports include maximum absolute difference, mean difference, maximum relative difference, and counts of mismatched elements.

Performance measurement uses AI Bench timing on the actual Intel GPU, with a warmup loop, a 256 MB L2 cache flush buffer, preallocated torch.Event pairs, a dummy 1024x1024 matmul before timing, synchronized execution, and trimmed means after removing minimum and maximum observations. The paper gives the throughput formulas as

$\to$9

and

$\to$0

where $\to$1 is mean execution time in microseconds.

A plausible implication is that Xe-Forge’s notion of optimization is inseparable from runtime compilation and execution on the deployment target. It does not accept profiler-free or simulator-only evidence as final proof of improvement.

5. Experimental evaluation and empirical performance

Xe-Forge is evaluated on 97 Level-2 KernelBench kernels and on Flash Attention, using an Intel Arc Pro B70 with 32 Xe2 cores, 256 XMX engines, 32 GB GDDR6, and 608 GB/s bandwidth. The software stack uses the Intel Triton GPU backend targeting SPIR-V, PyTorch 2 with torch.xpu, and a torch.compile baseline via TorchInductor. In the benchmark setting, the starting Triton kernels are obtained through KernelFalcon or KernelAgent, which underscores that Xe-Forge optimizes functionally correct Triton rather than generating it (Spoczynski et al., 16 Apr 2026).

The 97 kernels span seven families: 19 GEMM, 17 MatMul, 1 BMM, 21 Conv2D, 12 Conv3D, 10 ConvTranspose2D, and 17 ConvTranspose3D. Against PyTorch eager, Xe-Forge achieves a 1.17x geometric mean speedup, with 67% of kernels improving; nine kernels exceed 5x, and the largest reported speedup is 82x on Matmul_Min_Subtract. The system reports 100% correctness across all benchmarked optimized kernels under the stated tolerances.

Performance is heterogeneous across families. GEMM and MatMul constitute the strongest cases: the conclusion reports a 1.28x geometric mean over best baseline for GEMM and 1.76x for MatMul. The paper also cites up to 62x on Gemm_Max_Subtract_GELU and notes that some “super-peak” TFLOPS values arise because algorithmic restructuring reduces effective computation while FLOP accounting still reflects the original benchmark specification. The single BMM-family kernel, BMM_InstanceNorm_Sum_ResidualAdd_Multiply, reaches 23x over best baseline, largely through aggressive operator fusion. Conv2D kernels usually fall between 1.0x and 2.2x, with only two mild regressions in the 0.6x–0.7x range. ConvTranspose2D mostly lies within 0.6x–1.4x of best baseline, with kernel 44 reaching 7.5x by eliminating GlobalAvgPool intermediate materialization. Conv3D is mostly bandwidth-bound and usually stays in the 0.9x–1.3x range, with one notable 0.6x regression. ConvTranspose3D spans 0.5x–1.2x, with lower-half regressions attributed to suboptimal memory-access patterns.

Flash Attention functions as a stress test because it combines online softmax accumulation, GRF-bounded tiling, SLM capacity management, and careful boundary masking for irregular shapes. Xe-Forge evaluates 16 configurations varying batch size, head count, sequence length, and head dimension, including Llama 3 8B and 70B, Mistral 7B, Falcon 40B, GPT-NeoX-20B, Qwen variants, Mixtral 8x7B, and long-context settings up to $\to$2. Every tested Flash Attention configuration improves, with speedups from 2x to 13.3x and no regressions. The longest contexts, $\to$3, achieve 9x–13.3x, while shorter contexts at $\to$4–$\to$5 achieve 2x–4.4x.

The roofline analysis separates compute-bound and bandwidth-bound regimes. GEMM and MatMul occupy the high-arithmetic-intensity region, where the largest gains come from algorithmic and fusion stages; Conv3D and ConvTranspose3D remain close to the bandwidth slope, where gains are structurally smaller and memory-access optimization becomes the main lever. For Flash Attention, original kernels sit at 5–15 TFLOPS, while optimized kernels rise to 40–90 TFLOPS, which the paper describes as much closer to roofline though still 1.5x–3x below peak.

6. Interpretation, limitations, and related naming

Xe-Forge is presented as a staged, conservative, hardware-in-the-loop optimizer whose novelty lies in the combination of issue-driven planning, stage decomposition, an Intel-specific knowledge base, and real-device verification. Its strongest cases arise when semantically equivalent transformations eliminate computation or intermediate materialization, while more modest or negative outcomes appear in bandwidth-bound kernels and in cases where memory access remains suboptimal (Spoczynski et al., 16 Apr 2026).

The paper is explicit about several limitations. Not all kernels improve, and some regressions remain in the 0.5x–0.8x range. Failure modes include novel algorithmic patterns absent from model knowledge, complex interactions across multiple kernels that the current per-kernel pipeline does not reason about, and compiler-level failures in the Intel backend where Triton constructs that may be valid elsewhere fail during SPIR-V or IGC lowering. The cost profile is also nontrivial: worst-case operation can require 40–50 LLM calls per kernel, while typical runs use 10–20 calls because many stages are skipped, with an estimated cost of about \$\to$62.00 per kernel.

The authors nevertheless argue that the hardware-specific parts of the architecture are modular. Porting to another accelerator would mainly require hardware detection and parameter mapping, a target-specific knowledge base, and updates to the GPU-specific issue taxonomy; the estimated porting effort is 2–3 days of engineering effort. They also suggest extension to other source-level kernel languages such as SYCL, though that would require a new issue taxonomy and new transformation patterns.

The name “Xe-Forge” can invite confusion because “FORGE” also appears as the title of an unrelated 2026 system for automated vulnerability assessment and detection engineering (Shaikh, 2 Jun 2026). In the usage established by “Xe-Forge: Multi-Stage LLM-Powered Kernel Optimization for Intel GPU,” however, Xe-Forge denotes the Intel-GPU kernel-optimization system alone. Within that scope, its defining claim is that structured domain knowledge plus hardware-in-the-loop verification can systematically reduce the repetitive manual effort that otherwise gates deployment of Triton kernels on Intel Xe-class accelerators.