Xe2 Esimd Gemm

Xe2 (Lunar Lake/LNL, Battlemage/BMG) ESIMD GEMM Skill

Specialized knowledge for authoring and optimizing SYCL ESIMD matrix-multiply kernels on Intel Xe2 (Lunar Lake/LNL, Battlemage/BMG) architecture. Reference files hold detail; this file holds the critical rules and workflow.

Quick-Reference Rules (must follow every time)

Hardware limits

• Max WG threads = 32 when doubleGRF is on (256 × 64-byte GRF = 16 KB/thread). Never set nd_range work-group size > 32. See references/hardware-constraints.md.
• Always compile with doubleGRF — it is mandatory for large tile kernels. Do not remove it.
• Barriers: every thread in a WG must execute the same number of barrier() calls. Unequal counts cause GPU hang.

Inner-loop rules

• No if, no ?: inside the K-loop body. Move all runtime conditionals to the host or split the loop into phases.
• Use index += increments instead of recomputing full expressions each iteration. Saves XVE ALU ops.
• More XVE + XMX parallelism lowers frequency. Minimize XVE ops inside the K-loop; DPAS dominates and XVE pressure throttles the clock.

API namespace

• Use sycl::ext::intel::experimental::esimd (xesimd) for lsc_load_2d, lsc_prefetch_2d, config_2d_mem_access, xmx::dpas.
• lsc_prefetch_2d does not accept a payload object (config_2d_mem_access) — use the inline 6-argument form: lsc_prefetch_2d<T,BW,BH,N,L1H,L2H>(ptr, surfW, surfH, surfPitch, x, y).
• lsc_load_2d supports the payload object API (config_2d_mem_access) with set_x() / set_y().

Payload CSE (critical optimization)

• Construct config_2d_mem_access once outside the K-loop with static fields (ptr, surfW, surfH, surfPitch). Only call set_x() / set_y() inside the loop. This eliminates per-call descriptor rebuilds and dramatically reduces XVE activity.

DPAS / VNNI

• See references/kernel-patterns.md for tile sizes, VNNI packing, double-buffer pattern, and dpas call signature.

Performance testing

• Use different buffers (cache-busting) for memory-bound kernels; rotate through an array of input buffer pairs each iteration.
• Random non-zero init — avoid all-zero inputs; they hide NaN/denormal bugs and give unrealistically fast cache results.
• Warmup: run ≥ 20 iterations before timing.
• Iteration count: 100 iters minimum; 1000 iters preferred for stable measurement on compute-bound kernels.
• Sleep before timing (5 s) to let GPU reach steady frequency before measurement begins.
• See references/perf-testing.md for boilerplate.

Correctness testing

• Check for NaN in output before computing relative error (NaN denominator silently passes threshold checks).
• Use an FP32 CPU reference; compare half-precision GPU output with rel_rms < 0.5% as pass criterion.
• Test at M=N=K=256 (fast) before benchmarking at 4096.

Tile walk order

• Hilbert / Z-order (Morton) walk of work-group indices can improve L2 reuse only when not all WGs run concurrently. On BMG with a 4096×4096 problem all 256 WGs run simultaneously — walk order has no effect.
• Z-order walk: wg_m = z_deinterleave_even(wg_id), wg_n = z_deinterleave_odd(wg_id). Keep index compute light (bit ops only).

Compile command

icpx <src>.cpp -o <out>.exe \
  -fsycl -fsycl-targets=spir64_gen \
  -Xs "-device bmg -options -doubleGRF"

Do not use -fsycl-targets=intel_gpu_bmg_g21 (wrong target string). Do not use -O3 or -doubleGRF as top-level flags — they are ignored or warn.

SPIR-V linker errors

• Cause: kernel contains a runtime branch (if) whose condition can be evaluated on the host. Move the branch outside the SYCL kernel (template or host-side dispatch) to eliminate the SPIR-V link-time issue.

Workflow

1. Write kernel following inner-loop rules above.
2. Compile with the command above; check for spill warnings (warning: ... spilled ... bytes). Spill > 0 usually means GRF budget exceeded — reduce tile size or payload count.
3. Correctness test at small size (M=N=K=256). Check NaN, check rel RMS.
4. Benchmark at production size (M=N=K=4096) with 20 warmup + 1000 timed iters, 5 s pre-sleep.
5. Profile with VTune or --collect gpu-hotspots to measure XVE ALU2 %. Target < 5% XVE for DPAS-bound kernel.
6. Iterate: reduce XVE ops via payload CSE, index increment, loop splitting.

Sample code (ready to compile)

Copy assets to a working directory, then compile and run:

icpx <file>.cpp -o <file>.exe -fsycl -fsycl-targets=spir64_gen -Xs "-device bmg -options -doubleGRF"
powershell.exe -Command "& './<file>.exe'"

To reproduce all results in order, see references/code-index.md.

Reference files