Repository navigation
PR Digest Iteration 62
Period: 2026-07-19 to 2026-08-01 (Iteration 62) | Total PRs: 52 (41 from Xe2/Xe3/Xe3P, 11 from Xe4) | Lines changed: +40,936 / -11,610
The block-pointer API was fully removed from the frontend after its deprecation period. A new LTS DPAS path via GenISA builtins was added to work around SPIR-V DPAS rejections on older drivers. Performance regressions were fixed in masked store vectorization, HostTensorDescriptor specialization, and the eviction-policy lowering. The atomic_poll crash on 16-bit types was resolved by routing through a 32-bit load on architectures lacking 16-bit atomic support. A new SLM-staged K-tiling pass for non-DPAS hardware prevents PTSS overflow on large matrix kernels. Five upstream synchronizations were completed with a 99.32%–99.44% test pass rate.
Key accomplishments:
- Removed the deprecated block-pointer API from the Triton frontend, eliminating ongoing maintenance burden
- Added a GenISA fallback DPAS path for LTS IGC drivers that reject SPIR-V
__spirv_SubgroupMatrixMultiplyAccumulateINTELwith non-8-component operands - Fixed masked store vectorization to handle non-statically-uniform masks, recovering a 28% bandwidth improvement on PVC
- Fixed atomic_poll crash on Intel GPUs without 16-bit atomic support by routing through a 32-bit aligned load
- Rewrote HostTensorDescriptor specialization to encode full shape values in the key, fixing FuseReshape regressions and enabling better 2D block I/O
- Added StageLargeFMADotsViaSLM pass for non-DPAS hardware to prevent PTSS overflow on large blocked-encoding dots
- Fixed vLLM MoE WorkspaceManager initialization on XPU, unblocking MoE test coverage on B580
DPAS fallback for LTS drivers and 16-bit atomic fix.
-
#7583 Use GenISA for DPAS on LTS driver — Adds a
llvm.genx.GenISA.sub.group.dpasfallback path for MatrixDPASOp gated onttig.is_lts, matching the existing 2D block I/O fallback pattern. (+97/-2, @whitneywhtsang) - #7573 Fix tt.atomic_poll runtime crash for 16-bit types — On Intel GPUs without 16-bit atomic support, lowers atomic_poll via a 4-byte-aligned 32-bit load + half-word extraction instead of crashing. Fixes #7390. (+247/-56, @wdziurdz)
Vectorized masked stores, host-descriptor specialization, and 2D block load/store fixes.
- #7577 Vectorize masked StoreOp when mask isn't statically uniform — Enables SIMD vectorization for store ops with non-uniform masks by extracting the mask per-lane, recovering 28% bandwidth improvement on PVC issue #7571. (+142/-30, @anmyachev)
- #7533 Include shape values in HostTensorDescriptor specialization (cherry-pick) — Cherry-picks #7519 to encode all shape values in the specialization key, fixing FuseReshape regression and enabling compile-time divisibility proofs. (+42/-0, @whitneywhtsang)
- #7519 Include shape values in HostTensorDescriptor specialization — Replaces the "L" last-dim divisibility key with full shape encoding, fixing #7471 regression. (+38/-35, @whitneywhtsang)
- #7528 Revert eviction_policy LSC cache modes (cherry-pick) — Cherry-picks the revert of #6775 due to regressions. (+3/-152, @anmyachev)
- #7521 Revert eviction_policy LSC cache modes — Reverts the eviction_policy hint implementation that caused performance regressions. (+3/-152, @anmyachev)
- #7497 Fix crash on broadcast layouts in LoadStore/DPAS — Pads the unpacked value count back to the full register index expectation for broadcast-encoded loads and dots, fixing #7492. (+27/-22, @wdziurdz)
- #7490 Add TritonGen op for gathering load — Introduces a TritonGEN dialect op for predicated sub-group gather loads. (+242/-2, @chengjunlu)
- #7487 Support transpose B matrix load for sub-group-size=32 — Extends transpose 2D block load support for sub-group-size=32. (+259/-48, @chengjunlu)
- #7454 Add SPV_INTEL_int4 and update expandSubByteVectorBitcast — Adds the SPV_INTEL_int4 extension and updates the sub-byte vector bitcast expansion. (+9/-10, @whitneywhtsang)
- #7537 Update spirv-llvm-translator.conf — Updates the SPIR-V LLVM translator configuration. (+142/-1, @whitneywhtsang)
Block pointer removal, SLM-staged dot pass, and driver cleanup.
-
#7557 Remove block pointer support from the frontend — Removes the deprecated
tl.make_block_ptr/tl.advanceblock-pointer API after its deprecation period, reducing compiler maintenance burden. (+22/-672, @whitneywhtsang) -
#7291 Add StageLargeFMADotsViaSLM pass — Adds
tritonintelgpu-stage-large-fma-dots-via-slmthat stages blocked-encodingtt.dotoperands in SLM and unrolls K, preventing PTSS overflow on non-DPAS hardware whenperThreadBytes > 4096. (+997/-0, @sgurwinderr) - #7558 Partial sync DecomposeScaledBlocked from upstream — Syncs a non-functional portion of DecomposeScaledBlocked with upstream. (+14/-19, @whitneywhtsang)
- #7611 Clean up redundant code in driver.c by unifying the rebuild path — Refactors the Level Zero kernel/module compilation to centralize cleanup and unify the large-GRF retry logic. (+101/-93, @chengjunlu)
- #7551 Improve codegen retry for CRI — Updates codegen retry behavior to support larger GRF modes including 512-GRF for CRI in both the Level Zero driver and ocloc-based ZEBIN paths. (+80/-62, @chengjunlu)
- #7568 Remove deprecated block pointer from batched flash attention — Migrates the batched flash attention reference implementation off the deprecated block-pointer API. (+3/-29, @whitneywhtsang)
vLLM MoE fixes, skiplist work, CodeQL, and CI improvements.
-
#7462 Fix WorkspaceManager initialization for vLLM MoE tests on XPU — Initializes WorkspaceManager with
torch.device("xpu:0")in conftest.py, unblocking MoE tests on B580. (+141/-607, @sarah12121212) - #7556 Add test_triton_heuristics test based on reproducer from issue #7542 — Adds a regression test for the heuristics path exercised by #7542. (+89/-0, @slawblauciak)
- #7582 Create codeql.yml — Adds a CodeQL static analysis workflow to the repository. (+85/-0, @exolyr)
- #7599 Skip test_mamba_ssm_ssd tests that silently abort pytest — Creates #7600 and skips Mamba SSM tests that cause silent pytest worker aborts. (+18/-0, @quinnlp)
- #7598 Revert: Avoid creating sycl::event when submitting kernels — Reverts #6361 after benchmark regression discovered. (+24/-122, @quinnlp)
- #7594 Skip kda tests using block ptrs — Skips vLLM kda tests that depend on the now-removed block pointer API. (+70/-0, @quinnlp)
- #7547 Sign triton-shim wheels with keyless Sigstore — Adds cosign-based keyless wheel signing to the CI pipeline. (+47/-2, @kwasd)
- #7545 Build benchmarks from branch source when branch-specific config exists — Makes the benchmark build use the branch's source when a branch-specific config file is present. (+23/-6, @quinnlp)
- #7538 Verify nightly vllm-xpu-kernels wheel matches pin — Adds a check that the nightly vLLM XPU kernels wheel matches the pinned version. (+23/-7, @quinnlp)
- #7560 Move vLLM test schedules into per-arch wrapper workflows — Restructures vLLM CI test scheduling into per-arch wrapper jobs. (+34/-19, @quinnlp)
- #7522 Fix test_fast_sin_cos reference to avoid device-side fp64 — Updates the test reference to avoid triggering device-side fp64 dispatch. (+8/-4, @kwasd)
- #7527 Skip test_debug.py failures on B580/Xe2 — Adds skip guards for debug test failures specific to B580. (+42/-0, @kwasd)
- #7534 Bump actions/setup-python from 6 to 7 — Updates the GitHub Actions setup-python action to v7. (+18/-18, @app/dependabot)
- #7438 Extend vLLM MoE XPU patcher to allow more tests — Broadens the XPU patcher scope so more MoE tests can run without modification. (+31/-5, @sarah12121212)
- #7435 Capture hardware details with new agama versioning scheme (listed as #7434) — Updates the hardware-detail capture script for the new agama versioning scheme. (+20/-2, @kwasd)
- #7400 Increase GDN_ATTN test coverage for priority kernels — Adds GDN_ATTN test variants for vLLM priority kernels. (+17/-1, @quinnlp)
- #7348 Update spec_decode skiplist — Updates the speculative decode skiplist for current failing tests. (+65/-21, @sarah12121212)
- #7331 Increase TRITON_ATTN test coverage for priority kernels — Expands triton attention test coverage for priority vLLM kernels. (+385/-2, @quinnlp)
- #7317 Fix backward autotuning config for flash-attention benchmark — Fixes the autotune configuration for the flash-attention backward benchmark. (+11/-36, @sgurwinderr)
- #6361 Avoid creating sycl::event when submitting kernels — Reduces SYCL event creation overhead in the kernel submission path (later reverted in #7598). (+122/-24, @uditagarwal97)
- #6359 Remove GEMM benchmark with legacy block pointer — Removes the legacy block-pointer GEMM benchmark. (+2/-164, @anmyachev)
- #7535 Update README and tutorial 6 — Updates the README and tutorial 6 for current API. (+10/-6, @whitneywhtsang)
5 upstream merges from OpenAI Triton (commits e7eab7d, 3d6e050, 215af05, c393f1b, d16541f) — pass rate 99.32%–99.44%.