The PTSS Cliff: A Kernel That Won't Build
Why a large-K FMA matmul on a non-DPAS Intel iGPU fails before it ever runs
Where Bytes Can Live
Every GPU optimization argument is, at heart, an argument about where a value sits when you need it. Get the location wrong and either the memory bus or the register file bottlenecks you. Three locations matter for this story.
GRF — Register File
128 registers × 32 bytes = 4 KB per thread, 1-cycle access. The entire universe of things a thread can name at once. Overflow it and the compiler spills.
SLM — Shared Local Memory
A software-managed scratchpad shared across a workgroup. 64 KB on Xe-LPG iGPUs, one or two cycles to access. CUDA calls it "shared memory."
PTSS — Per-Thread Scratch
A DRAM region (backed by L3) where spilled registers land. Capped at 256 KB per thread. Spill past that and the kernel does not even build.
Smaller, faster, more private as you go up: DRAM/PTSS → L3 → SLM → GRF. This pass moves bytes between these levels to dodge a specific disaster.
Two Paths From a tt.dot
When you write tl.dot(a, b, acc) in a Triton kernel, the compiler picks one of two lowering paths depending on the hardware.
DPAS Path
On hardware with a systolic matrix unit (PVC, BMG, ARL-H Xe2), the dot lowers to a few dpas instructions per K step. Register pressure stays bounded. This pass is a no-op here.
FMA Path
On non-DPAS iGPUs (ARL-S, MTL, A770), there is no matrix unit. The dot falls back to a hand-coded FMA loop nest, fully unrolled along K. This is the path that breaks.
The Arithmetic of "Fully Unrolled K"
Take a concrete shape: a 64×128×128 f16 GEMM where one thread covers an 8×8 micro-tile of the output. For its 8×8 chunk of C, the thread needs:
# One thread owns an 8x8 block of C
C[i][j] = sum(A[i][k] * B[k][j]
for k in range(128))
# 64 output cells x 128 K-steps
# = 8192 FMAs per thread, unrolled flatEach thread is responsible for 64 output cells.
Each cell needs a 128-step dot product along K.
The Triton lowering plus IGC's backend fully unroll that K loop — all 8192 FMAs live flat, with no loop structure left to bound register pressure.
There is no TTGIR-level K-tile knob in upstream's pipeliner, and IGC's unroll heuristic does not fire at this pressure regime.
The Cliff: 345 KB Spill, 256 KB Cap
A SIMD16 thread on Xe-LPG has a 4 KB GRF in the default small register mode. The fully-unrolled schedule wants far more than that, so IGC spills the surplus into PTSS. But ARL-S, A770, and MTL iGPUs cap PTSS at 256 KB per thread, enforced by Intel's compute-runtime (NEO) driver.
On the 64×128×128 f16 GEMM, the spilled schedule wants 345 KB per thread — an 89 KB overshoot past the ceiling.
The backend cannot fit the schedule into the 256 KB scratch budget.
ZE_RESULT_ERROR_INVALID_NATIVE_BINARYThe kernel is rejected at zeKernelCreate time. No SPIR-V is emitted.
The kernel does not run slowly — it does not run at all. And there is no compiler flag to flip: the -ftarget-register-alloc-mode large-GRF switch is restricted to PVC, so it is not a portable escape hatch.
Check Your Understanding
Which memory region's 256 KB cap causes the build to fail outright?
Which lowering path hits the PTSS cliff?
The Idea: Stage in SLM, Tile K
Trade a register-bound flat schedule for a small, bounded loop fed from the scratchpad
Two Moves, One Goal
The fix is a new TTGIR pass — tritonintelgpu-stage-large-fma-dots-via-slm — that rewrites a doomed tt.dot into a shape IGC can actually schedule. It makes exactly two moves.
Move 1 — Stage operands in SLM
Instead of keeping the full A and B tiles live in registers, write them once to SLM. The scratchpad becomes the home for the operands; the GRF only holds the slice in flight.
Move 2 — Tile the K dimension
Split the 128-deep dot into chunks of K_TILE. Loop over the chunks, accumulating into C. Only K_TILE columns of each operand are live in registers at any instant.
We swap a flat, fully-unrolled schedule that spills 345 KB to PTSS for a small, looped schedule whose live set is bounded by K_TILE. SLM bandwidth replaces register pressure — and SLM is the one budget we have headroom in.
Watch the Bytes Move
Step through how a single K-tiled, SLM-staged dot processes its operands. Notice that the register file only ever holds one K_TILE-wide slice at a time.
Why SLM Is the Right Place to Spill
We are not eliminating the data movement — the operands are big, they have to live somewhere. The insight is which budget to spend.
Hard 256 KB per-thread cap — overflow means no binary. Backed by DRAM via L3 with high miss latency, private per thread so nothing is reused, and there is no TTGIR-level knob to control it.
A 64 KB workgroup budget with real headroom at these tile sizes, one- to two-cycle access instead of DRAM latency, shared across the workgroup so operands can be reused, and explicitly sized and placed by the pass itself.
The Algorithm & the IR Rewrite
Five steps that turn one oversized tt.dot into a tiled, SLM-staged loop
The Pass, Step by Step
The pass walks every tt.dot in the module. For each one that lands on the FMA path and trips the size gate (Module 4), it performs this rewrite:
Reserve two ttg.local_alloc buffers — one for the A operand tile, one for B — sized to the operand tiles in their native dtype.
Emit ttg.local_store for A and B before the dot. The full operands now live in the scratchpad, not in registers.
Wrap the accumulation in an scf.for over K / K_TILE iterations, carrying the C accumulator as a loop-carried value.
Inside the loop, ttg.local_load one K_TILE-wide slice of A and B, run a small tt.dot on it, and add into the accumulator.
After the loop, the carried accumulator is the original dot's result. Replace all uses of the old tt.dot with it and erase the original.
Before → After, in TTGIR
Here is the shape of the rewrite. On the left, the original single oversized dot; on the right, what the pass leaves behind.
// K = 128, fully unrolled by the
// FMA lowering -> 345 KB spill
%c = tt.dot %a, %b, %c0
: tensor<64x128xf16>
* tensor<128x128xf16>
-> tensor<64x128xf32>%sA = ttg.local_alloc / %sB = ttg.local_alloc — SLM scratch for both operands.
ttg.local_store %a, %sA and %b, %sB — operands written to SLM once.
%c = scf.for %k = 0 to 128 step K_TILE iter_args(%acc = %c0) — the K-tile loop.
Inside: local_load a K_TILE slice of each, a small tt.dot, then scf.yield %acc + %partial.
All uses of the old result are RAUW'd to the loop's result; the original dot is erased.
Every op the pass emits — local_alloc, local_store, local_load, scf.for, tt.dot — already has a lowering. The pass composes existing primitives; it introduces no new codegen.
Check Your Understanding
Which ops does the pass use to move operands into and out of SLM?
Picking K_TILE & the Four Gates
A byte budget decides the tile size; four predicates decide whether to fire at all
The SLM Footprint
Before tiling, the pass checks that staging both operands in SLM actually fits the workgroup budget. For the running 64×128×128 f16 example:
# A tile: 64 x 128 f16
A = 64 * 128 * 2 = 16384 B
# B tile: 128 x 128 f16
B = 128 * 128 * 2 = 32768 B
total = 48 KB < 64 KB SLM ✓Staging both operands costs 48 KB of the 64 KB workgroup SLM — it fits, with 16 KB to spare.
If the operands did not fit, the pass would bail rather than blow the SLM budget and serialize the workgroup.
K_TILE then controls only how much of each operand is live in registers per loop iteration — the SLM copy stays whole.
Choosing K_TILE
The tile size is the one tuning knob. Pick it too large and the per-iteration live set spills again; too small and loop overhead and SLM traffic dominate.
Too large
A wide slice keeps too many columns live per iteration — register pressure climbs back toward the PTSS cliff we were escaping.
Too small
Many tiny iterations mean more loop bookkeeping and more local_load round-trips — SLM bandwidth and loop overhead start to dominate.
Just right
The largest K_TILE whose live set still fits the GRF budget — maximal work per iteration with no spill. The pass solves for this directly.
The Four Gates
The pass is conservative by design. A dot is rewritten only if it passes all four gates — otherwise it is left exactly as-is, guaranteeing the pass can never make a healthy kernel worse.
The target must lack a DPAS unit. On DPAS hardware the dot is already cheap, so the pass refuses to touch it.
The dot's estimated spill must actually exceed the budget. Small dots that schedule fine are left alone.
K must be large enough, and cleanly divisible, to split into K_TILE chunks without a ragged remainder.
The staged operands must fit the workgroup SLM budget. If staging would overflow SLM, the pass declines.
Because all four gates must pass and the fallback is "do nothing," the worst case for any kernel the pass doesn't help is a no-op. That property is what makes it safe to enable by default.
Check Your Understanding
How many of the four gates must a dot pass to be rewritten?
Pipeline Placement & Design Choices
Where the pass sits in the compiler, and how the shipped design diverged from the first sketch
Where It Runs
Placement matters as much as the rewrite. The pass runs in the TTGIR phase — after layouts are assigned (so it knows the operand shapes and the FMA-vs-DPAS decision) but before the software pipeliner and SLM allocation finalize.
The pass needs the resolved operand layouts and the FMA/DPAS verdict to evaluate its gates. Running earlier would mean guessing.
The new scf.for K-loop and SLM ops are handed to the existing pipeliner and allocator, which schedule and place them like any other loop.
It lives in the intel third-party tree as a target-specific TTGIR pass, so it only ever runs for Intel GPU compilation.
Because the pass emits ordinary scf.for + SLM ops, the downstream pipeliner, allocator, and lowerings treat its output as a normal loop. No special-casing anywhere else in the stack.
Five Ways the Final Design Diverged
The shipped pass is not the napkin sketch from the original proposal (PR #7276). Five deliberate changes hardened it for production.
The sketch patched late, near codegen. The shipped pass works in TTGIR so it can reuse local_alloc/local_load and ride the existing pipeliner.
The sketch assumed the caller only invoked it on FMA dots. The shipped pass checks the path itself, so it is safe to run unconditionally.
Staging only the larger operand left the smaller one spilling. The final design stages A and B for a clean, bounded live set.
K_TILE, not hard-codedA fixed tile size regressed some shapes. The pass now derives K_TILE from the register and SLM budgets per dot.
The sketch always rewrote. The shipped pass added the size, tileability, and SLM-fit gates so the fallback is always a guaranteed no-op.
Validation, Prior Art & Risk
What the pass actually recovers on ARL-S, where the idea comes from, and why it is safe
Measured on ARL-S
The pass was validated on an Arrow Lake-S iGPU across three groups of kernels: ones that used to fail, ones that already worked but were slow, and ones that should be untouched.
Every one of the nine kernels that hit the INVALID_NATIVE_BINARY cliff now builds and runs correctly. The PTSS overflow is gone.
Kernels that compiled but spilled heavily got faster too — a 1.86× geometric-mean speedup from trading PTSS round-trips for bounded SLM traffic.
DPAS-path and small-dot kernels produced byte-identical IR with the pass on or off — confirming the four gates hold.
Not a New Idea — A Ported One
Staging operands in scratchpad and tiling the contraction dimension is the oldest trick in the high-performance GEMM book. This pass brings it to a path that was missing it.
Every mature GEMM stack tiles K through scratchpad — but only on its fast path. The non-DPAS FMA fallback on Intel iGPUs never got that treatment. This pass closes that specific gap.
Risk & Composability
The pass ships behind the four gates and emits only standard TTGIR. That bounds its blast radius tightly.
Bounded downside
If any gate fails, the dot is untouched. The worst case for an unhelped kernel is a no-op — never a regression.
Composes cleanly
Output is ordinary scf.for + SLM ops, so the pipeliner, allocator, and lowerings need no special-casing.
Watch SLM pressure
Staging consumes SLM that other parts of a kernel might want. Gate 4 guards the budget, but occupancy on SLM-heavy kernels is worth monitoring.
Final Check
On ARL-S, what happened to the nine kernels in Group A that previously failed to build?
What is the worst-case outcome for a kernel that doesn't clear the gates?
This pass landed as intel/intel-xpu-backend-for-triton#7291, evolving the original sketch from PR #7276.