PR Digest Iteration 57

PR Summary

Period: 2026-05-10 to 2026-05-23 (Iteration 57) | Total PRs: 59 (43 from Xe2/Xe3/Xe3P, 16 from Xe4) | Lines changed: +80,602 / -19,528

---

Triton XPU BE (Xe2/Xe3/Xe3P) (43 PRs, +22,897 / -6,571)

The team landed two new compiler analyses (TemporalReuseAnalysis and a ReuseAnalysis facade) and rebuilt AnnotateCacheControl on top of them, providing a foundation for upcoming cache-locality optimizations. The eviction-policy hint on tt.load and tt.descriptor_load was wired through to actual LSC cache modes for the first time. Several correctness fixes were made across stride, integer range, signed-division, and float32 reduction analysis to prevent silently wrong results. The team also hardened the Intel block-I/O lowering path with multiple bug fixes and a major refactor, and completed four upstream synchronizations with a 99.04%–99.06% test pass rate.

Key accomplishments:

• Introduced new compiler analyses underpinning cache-locality and reuse-driven optimizations, migrated AnnotateCacheControl to use them, and enabled the pass by default.
• Wired evict_first / evict_last hints through to LSC cache-control modes so load reuse decisions now take effect in hardware.
• Fixed several silent-correctness bugs in stride, integer range, signed division, and float32 reduction analysis.
• Strengthened the 2D block-load lowering path with bug fixes for transposed loads, rank-3 batch loads, tensor-of-pointers surface dimensions, and a major refactor that unifies dispatch.
• Made the backend more resilient on Linux LTS2 and Windows by recovering gracefully when zebin metadata is missing or incomplete.

Memory Access & Lowering

Eviction-policy hint wiring, refactors and bug fixes to the Intel 2D block-load path, plus sub-byte element-type fixes uncovered by tightened SPIR-V validation.

• #6775 Honor eviction_policy hints via LSC cache modes — Routes evict_first/evict_last hints on tt.load and tt.descriptor_load to the nearest LSC load-cache-control mode when no explicit cache_modifier is set. (+152/-3, @etiotto)
• #6962 Fix surface dimensions for tensor-of-pointers 2D block loads — Corrects 2D block surface width/height when lowering loads from tensor-of-pointers operands. (+19/-4, @whitneywhtsang)
• #6957 Widen sub-byte bit ops forbidden by SPV_INTEL_int4 — Cherry-picks #6927 to widen sub-byte bit operations the SPIR-V int4 extension forbids. (+231/-10, @anmyachev)
• #6947 Unify 2D block load lowering via lowerBlockLoad2D() callback pattern — Consolidates the multiple block-load lowering paths behind a single callback-driven dispatcher. (+162/-174, @whitneywhtsang)
• #6945 Extract Block2DLoadConfig and buildBlock2DLoadConfig() helper — Pulls the 2D block load configuration out into a reusable helper struct and builder. (+174/-159, @whitneywhtsang)
• #6944 Remove getBlockIOTileSize forwarding wrapper and add store assertion — Drops a redundant wrapper and asserts on unsupported store tile sizes. (+6/-14, @whitneywhtsang)
• #6932 Fix sub-byte element type allocation size mismatch in local_alloc — Aligns sub-byte allocation sizes between local_alloc and downstream consumers. (+82/-6, @app/copilot-swe-agent)
• #6927 Widen sub-byte bit ops forbidden by SPV_INTEL_int4 — Widens sub-byte and/or/xor operations the SPIR-V int4 extension forbids on int4 types. (+231/-10, @wdziurdz)
• #6909 Fix pitch computation for transposed pointer loads — Corrects the pitch field used for transposed 2D block loads from tensors of pointers. (+66/-15, @whitneywhtsang)
• #6908 Allow rank-3 pointer loads with batch dimensions — Extends rank-3 batched pointer loads to use the 2D block path. (+29/-9, @whitneywhtsang)
• #6898 Pass oneMatrixPerLoadForBT to tile validation — Threads the one-matrix-per-load flag into tile validation so BT layouts are correctly accepted/rejected. (+119/-47, @whitneywhtsang)
• #6887 Remove LoadOpToBlockIOConversion and DescriptorLoadOpToBlockIOConversion — Deletes the legacy conversion patterns now subsumed by the unified LowerTo2DBlockLoad pass. (+7/-701, @whitneywhtsang)

Analysis & Correctness

New reuse analyses, the AnnotateCacheControl migration, a DecomposeScaledBlocked fix, and several correctness fixes in stride, range, reduction, and membar analysis.

• #6984 Extend TemporalReuseAnalysis with provenTemporalReuse — Adds provenTemporalReuse(), a positive-proof query requiring every enclosing loop to be classified Held or Invariant; complements the existing suppress-side hasTemporalReuse. (+408/-50, @etiotto)
• #6983 Extend SpatialReuseAnalysis with std::optional known-* APIs — Adds knownCrossSubgroupReuse() and knownWarpBroadcastFactor() which return positive proof only (no conservative fallback), so callers that force actions like setting EVICT_LAST get a safe signal. (+417/-14, @etiotto)
• #6979 Enable AnnotateCacheControl pass by default — Flips TRITON_INTEL_DISABLE_ANNOTATE_CACHE_CONTROL default from disabled to enabled so the cache-control pass runs on every compilation. (+276/-13, @etiotto)
• #6977 DecomposeScaledBlocked — use deduceScaleFactor() instead of hardcoded 32 — Replaces the hardcoded scale factor of 32 with deduceScaleFactor() so the Intel fallback correctly handles NVFP4's factor=16. (+44/-3, @etiotto)
• #6969 Fix bug in Intel Membar analyze — Corrects an analysis bug in the Intel membar pass that could produce incorrect barrier placement. (+25/-2, @anmyachev)
• #6948 Rewrite AnnotateCacheControl on top of ReuseAnalysis + AliasAnalysis — Replaces ad-hoc cache-control heuristics with a ReuseAnalysis-driven pass that uses alias info to assign cache hints. (+505/-237, @etiotto)
• #6931 RemoveLayoutConversions — Guard hoist bitwidth check on int/float element type — Restricts the bitwidth-comparison guard in convert-hoisting to integer/float element types only. (+111/-12, @etiotto)
• #6925 IntegerRangeAnalysis — Fix std::optional misuse in getTripCount — Cherry-picks #6922 to fix optional misuse that produced wrong loop-trip-count results and broken cmpi folds. (+38/-6, @whitneywhtsang)
• #6922 IntegerRangeAnalysis — Fix std::optional misuse in getTripCount — Replaces optional tri-state misuse with explicit return paths in the trip-count analysis. (+38/-6, @whitneywhtsang)
• #6919 Fix float32 reduction accuracy regression in ReduceOpToLLVM — Cherry-picks #6914 to restore float32 reduction accuracy after a recent regression. (+68/-14, @whitneywhtsang)
• #6914 Fix float32 reduction accuracy regression in ReduceOpToLLVM — Restores the previous reduction order in ReduceOpToLLVM lowering to recover float32 numerical accuracy. (+68/-14, @januszjah)
• #6896 AxisInfo — Fix signed division AxisInfo bug and revert PyTorch workaround — Fixes signed-division stride inference and removes the previous PyTorch-side workaround. (+79/-6, @whitneywhtsang)
• #6875 Analysis — Add ReuseAnalysis facade — Adds a unified facade over the underlying reuse analyses for use by downstream passes. (+275/-0, @etiotto)
• #6860 Analysis — Add TemporalReuseAnalysis — Introduces a new analysis that detects loads whose values are reused across loop iterations. (+1492/-0, @etiotto)

Backend Robustness

Fixes to make the backend recover gracefully when LTS2 IGC produces degenerate zebin output, plus Windows static-launcher fixes.

• #6967 Fix static launcher for msvc — Cherry-picks #6965 to fix the Windows static launcher build under MSVC. (+34/-52, @anmyachev)
• #6966 Trigger 256-GRF retry when zebin is degenerate (LTS2 IGC silent failure) — Cherry-picks #6958 so a degenerate zebin from LTS2 IGC triggers an automatic 256-GRF recompile. (+31/-3, @anmyachev)
• #6965 Fix static launcher for msvc — Repairs the Windows static launcher under MSVC after recent build changes. (+34/-52, @anmyachev)
• #6958 Trigger 256-GRF retry when zebin is degenerate (LTS2 IGC silent failure) — Detects degenerate zebin output from LTS2 IGC and retries compilation with 256-GRF mode. (+31/-3, @dev-tomek)
• #6940 Don't use removed sycl::ext::intel::device_ptr in AOT tests — Updates AOT tests to stop using a sycl::ext::intel::device_ptr API removed in DLE2026. (+3/-6, @anmyachev)
• #6930 Don't raise an error when .ze_info section is missing from zebin — Cherry-picks #6921 so a missing .ze_info section is treated as a recoverable case. (+110/-2, @anmyachev)
• #6982 Build oneDNN library in case it is downloaded from the web — Ensures the downloaded oneDNN library is built before linking so the benchmark target doesn't fail to find -ldnnl. (+10/-2, @gshimansky)
• #6921 Don't raise an error when .ze_info section is missing from zebin — Replaces the hard error on missing .ze_info with graceful handling so kernels still load. (+110/-2, @dev-tomek)

Upstream Alignment

4 upstream merges from OpenAI Triton (commits a4b65bf, 44c91e5, 0bdab22, c7391a2) — pass rate 99.04%–99.06%.

Tooling, Benchmarks & CI

Benchmark and CI infrastructure improvements, vLLM/Proton integration work, and NVL architecture support.

• #6994 Add NVL archs — Adds NVL architecture entries for use by downstream test and deployment paths. (+11/-0, @anmyachev)
• #6959 Allow benchmarks to run as pytest tests — Wires benchmark scripts so they can be invoked from pytest, simplifying nightly runs. (+241/-2, @gshimansky)
• #6935 Add autotune for fused moe kernel — Enables autotuning of the fused MoE kernel used by vLLM. (+39/-6, @chengjunlu)
• #6915 Show COMPILER_VERSION in Report environment details step — Surfaces the IGC compiler version in CI environment reports for easier debugging. (+88/-85, @anmyachev)
• #6900 Suppress false 'No junit xml files' warnings for tutorial benchmark subdirs — Filters out spurious junit-xml warnings emitted when tutorial benchmark dirs are empty. (+47/-0, @januszjah)
• #6881 Remove XE-TLA from gemm_benchmark — Drops the XE-TLA reference path from the gemm benchmark. (+3/-34, @slawblauciak)
• #6766 Migrate benchmark workflows to triton-benchmarks CLI — Migrates benchmark workflows to the new triton-benchmarks CLI driver. (+287/-166, @januszjah)
• #6758 vLLM — update pin past unified attention 2d and 3d unification refactor — Bumps the vLLM pin past the unified attention 2D/3D refactor. (+255/-248, @quinnlp)
• #6709 Proton — Enable tensor metrics API for XPU — Enables Proton's tensor metrics API on XPU devices. (+184/-85, @wdziurdz)