Merge OpenAI Triton commit ` c393f1b` by wdziurdz · Pull Request #7572 · intel/intel-xpu-backend-for-triton

wdziurdz · 2026-07-24T11:20:59Z

This PR changes the Triton base from 215af05 to c393f1b (Jul 23) .

Pass rate: 99.32%->

Restores the correctness portion of #10922 for ClusterBarrierInsertion, a tiny fix in a lowering, and proper typing of the barriers in CLC.

CPU Chrome events currently carry Proton's internal trace event ID for CPU-to-GPU flow matching, but that ID is not exposed by Python proton.scope. Preserve and export the original scope ID so an external tracer can deterministically link its scope-entry flow to the corresponding CPU event. This allows us to do something like lightseekorg/tokenspeed#740 (comment) --------- Signed-off-by: Lei Zhang <antiagainst@gmail.com>

Fixes several gfx1250 gluon test issues around async waits and input ranges and cache mods.

``` commit 909ba8d3508cf7177da65a4d1399f7bc5a3a4b98 Author: Changpeng Fang <changpeng.fang@amd.com> Date: Thu Jul 16 12:30:39 2026 -0700 [AMDGPU] Select upper 16-bits for inline constants in packed BF16 (#209861) From gfx1250 software programming guide, v_pk_*_bf16 instructions using inline constants must use OPSEL to select the upper 16-bits. Fixes: LCOMPILER-2445 ```

PR description written by Codex Cluster barrier `arrive` and `wait` don't work under warp specialization, aren't modeled by consan, gsan, etc. and generally have ill-defined semantics. Presently, they are only used by low-level lowering (TMEM allocation/deallocation). The mbarrier init fence is also inserted automatically by the compiler. Remove them from the public Gluon API to prevent misuse.

PR description written by Codex Atomics get cluster barriers, and also broadcast their result across CTAs in multi-CTA kernels. When those atomic ops are converted to GSAN instrumented ops, the corresponding GSAN ops aren't lowered the same way. Introduce an `AtomicOpInterface` and unify the atomic lowering logic. Without this, multi-CTA kernels with atomics miscompile under GSAN.

Before this PR scaled_upcast required unpacked scales which did result in unnecessary reshape + convert_layout, which did cause a roundtrip to LDS in some cases. With this PR we can do a `scaled_upcast` which directly feeds a dot without any expensive `convert_layouts` (they will be nops). The lowering is generalized to build a map from 8-packed elements to their scales value instead of hardcoding the distribution of scales to registers. This allows us to handle arbitrary "BLOCK_SCALE" values along the axis dim, the only real requirement is that 8 contiguous fp4 values use the same scale (intrinsic requirement). The verifier is extended to infer the scale layout from the output encoding which will also be used to verify the passed layout from Gluon. Note that we remove broadcasting dims for the comparison since the lowering will do the coordinate based lookup so broadcasting dims will be ignored anyways. The intrinsic also supports `opsel` to densely pack/select the fp8 scales into/from registers, this will be added in a follow up PR and will be done purely in the lowering of the op.

PR description written by Codex Refactors recurring LinearLayout operations into a small set of reusable utilities: - input and output basis masks; - checked optional `lstsq` factorization; - maximal identity-prefix factoring; and - order-preserving dimension renaming. The existing specialized `f2reduce` implementations stay in place. This PR does not add a general matrix subsystem or change memdesc semantics. ## Stack Merge bottom-up: - [ ] #10946 (this PR) - [ ] #10947

…h-size recompiles) (#10875) Fixes #10872. `topk_forward` allocates the bitmatrix with `stride(1) = cdiv(n_rows, 32) * 32`, but `_topk_forward` and `_sum_bitmatrix_rows` declare their bitmatrix strides `tl.constexpr`, so every new 32-row bucket of `n_rows` triggers a full recompile at run time (~700ms each). In LLM serving `n_rows` is the per-batch token count, so this fires continually on the request path — originally observed as recurring ~1s TTFT stalls on a production gpt-oss-120b deployment (4×H100, TP4; details and standalone repro in #10872). The strides only feed address arithmetic: `stride_rm == 1` still specializes to a constexpr via the `==1` rule, and `stride_rn` is always a multiple of 32 so the divisibility-16 specialization keeps the alignment facts. Generated PTX is unchanged for both kernels (identical global memory-op histograms; neither kernel has any vectorised access involving these strides — the only wide loads are the logits reads on a literal stride-1 axis), kernel timing is flat, and outputs are bit-identical. Controlled serving A/B (sglang `triton_kernel` MoE backend, gpt-oss-20b on 1×H100 — same kernels/quantization as 120b, sized for paired experiments; cold engine, identical 200-request varied-length sequences at 1 QPS): - steady state: stalls from these two kernels drop to zero — stock still hits fresh recompiles indefinitely as new lengths arrive (p99 TTFT 682ms, stalls ongoing after 200 warmup requests); fixed: p99 316ms = p50, no stalls. - whole cold run including startup: 66/200 → 14/200 requests >500ms. The residual 14 are first-encounter compiles of unrelated kernels' bounded config space (matmul block_m/split_k variants): 9 in the initial burst plus 5 isolated later first-touches — at identical request positions with near-identical durations whether or not this PR's patches are applied (and present in the stock run at the same positions), confirming they are independent of these two kernels. Zero events in the subsequent 100-request phase. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

Fixes #10981

On pre-Rubin we just emit one arrive instruction that's executed from the first `num_ctas` threads. On Rubin, we use the new `multitcast` arrive.

We found this interesting problem with in multicta consan where you have a pattern that looks as follows in the CLC partition: ``` python @gluon.jit def consumer(slot, sink, consumed, layout: ttgl.constexpr): value = slot.load(layout) sink.store(value) mbarrier.arrive(consumed, count=1) @gluon.jit def clc_partition(slot, result, clc_bar, consumed): mbarrier.wait(consumed, 0, deps=[slot]) mbarrier.expect(clc_bar, 16) clc.try_cancel(result, clc_bar) mbarrier.wait(clc_bar, 0) clc.load_result(result) slot.store(gl.full([1], 1, gl.int64, layout=layout)) ``` Assume the kernel has 2 CTAs. Then, the way that CLC works on multicast is that CTA0 tries to cancel a block, and if it manages to do so, multicasts the result to all CTAs (into `result` in this example). What's rather surprising is that waiting on CTA1 on the the `clc_bar` where this arrive is broadcasted DOES NOT transfer the visibility of the `slot.load` from CTA0 to CTA1 so it sees `slot.store` as a race. This is quite crazy as clearly there isn't a way for `slot.store` to execute before the `slot.load` in the `consumer` partition has finished! According to codex: ``` In this case there is an operational chain: load finishes → consumed arrival → consumed wait finishes → CLC is issued → CLC completes → CLC wait finishes → store executes so it is entirely reasonable to conclude that the load/store cannot physically overlap. What PTX declines to infer is a transitive memory-synchronization edge through CLC. CLC guarantees completion and visibility of its own response writes; it does not say that completion publishes arbitrary operations that happened before the request was issued. This is intentional conservatism in the model, not an internal contradiction: completion/execution order and cross-thread memory causality are separate concepts. In fact, the official CLC example explicitly inserts synchronization between iterations before reusing the response storage; it does not rely on CLC completion to order prior accesses. PTX CLC example ``` This PR adds a `0 <= mask < num_ctas` whose on-bits represent the bits of the CTA that we shall use as the origin of the arrive / expect. In other words, we predicate the instruction by `cta_id & mask`. This allows us in this case to send the arrive / expect from CTA0 to all CTAs effectively sending the frontier visibility from CTA0 to all other CTAs when they wait on the barrier. e.g. for 8 CTAS and `mask = 0x011`: ``` CTA0 -> CTA0, CTA4 CTA1 -> CTA1, CTA5 CTA2 -> CTA2, CTA6 CTA3 -> CTA3, CTA7 CTA4-7 -> predicated ``` and 8CTAs for `mask = 0b110`: ``` CTA0 -> CTA0, CTA1 CTA1 -> predicated CTA2 -> CTA2, CTA3 CTA3 -> predicated CTA4 -> CTA4, CTA5 CTA5 -> predicated CTA6 -> CTA6, CTA7 CTA7 -> predicated ... I also renamed and reordered the masks that `arrive` may receive.

…ng the same accumulator" (#10988) reland triton-lang/triton#10858. It was revert due to a wrong bisect.

## Summary Add an opt-in `TRITON_FPSAN_HOMOMORPHIC_CASTS` compilation flag that restores truncating FPSan payload downcasts. The default FPSan behavior remains unchanged. When enabled, narrowing drops the high payload bits without folding them into the retained bits, allowing independently computed BF16 matmul partials to be combined bitwise: ```python x @ w == x[:, :half] @ w[:half, :] + x[:, half:] @ w[half:, :] ``` This permits distributed BF16 reductions to retain the normal FP32 matmul accumulator while preserving FPSan parity across sharding layouts. The flag can be configured programmatically: ```python triton.knobs.compilation.fpsan_homomorphic_casts = True ``` or through the environment: ```bash TRITON_FPSAN_HOMOMORPHIC_CASTS=1 ``` and is included in the compiler cache key. ## Test plan Added a deterministic BF16-dot sharding test using a payload carry witness: ```text 0x00ff * 0x0101 + 1 = 0x00010000 ``` The test is strict-XFAIL with the flag disabled and passes with it enabled. Validated on GB300: ```text 13 passed, 1 xfailed FPSan lit checks: PASS ```

PR description written by Codex Allow non-duplicating cast/broadcast layout-conversion hoists when rematerialization splitting is disabled, while retaining the existing multi-use guard. This fixes four ConSan matmul compilations whose shared-memory requirement was 236,764 bytes; SM100 compilation now requires 228,572 bytes with the same 12-warp, 2-CTA configuration. The focused removal, ConSan, and FPSan lit tests pass. This draft intentionally keeps the change scoped: forward layout-conflict resolution is still not shared-memory-aware and can increase conversion scratch for unrelated multi-layout cases. ## Stack Merge bottom-up: - [ ] #11003 (this PR)

…#10850) In hex mode, bitcast floating-point operands to same-width integers before printf promotion so their raw IEEE-754 bits live on NVIDIA and AMD backends. Format interpreter hex output using the operand dtype width and view floating-point values as unsigned integers. Extend device_print hex coverage to float16, float32, and float64, verify exact interpreter bit patterns, and check hard-coded fp32 canonical values for 1.0, negative zero, and positive and negative infinity. Tested with make and the full test_print matrix: 35 GPU tests and 35 interpreter tests pass.  # New contributor declaration - [x] I am not making a trivial change, such as fixing a typo in a comment. - [x] I have written a PR description following these [rules](https://cbea.ms/git-commit/#why-not-how). - [x] I have run `pre-commit run --from-ref origin/main --to-ref HEAD`. - Select one of the following. - [x] I have added tests. - `/test` for `lit` tests - `/unittest` for C++ tests - `/python/test` for end-to-end tests - [ ] This PR does not need a test because `FILL THIS IN`. - Select one of the following. - [x] I have not added any `lit` tests. - [ ] The `lit` tests I have added follow these [best practices](https://mlir.llvm.org/getting_started/TestingGuide/#filecheck-best-practices), including the "tests should be minimal" section. (Usually running Python code and using the instructions it generates is not minimal.)

… contiguosu subviews (#10997) PR description written by Codex As per title ## Stack Merge bottom-up: - [ ] #11006 - [ ] #10997 (this PR)

PR description written by Codex Fix multibuffer shared-memory subslices when the leading pipelining dimension is outside the layout rank, including physical offsets and buffer-region analysis. Allow tensor-memory scale subviews, centralize pipelining-dimension handling, and preserve the existing reinterpret and subview restrictions. Validated with the focused lit, C++ layout, Gluon GPU, and ConSan suites on GB300. ## Stack Merge bottom-up: - [ ] #10856 (this PR) --------- Co-authored-by: lezcano <lezcano-93@hotmail.com>

# New contributor declaration - [x] I am not making a trivial change, such as fixing a typo in a comment. - [x] I have written a PR description following these [rules](https://cbea.ms/git-commit/#why-not-how). - [x] I have run `pre-commit run --from-ref origin/main --to-ref HEAD`. - Select one of the following. - [x] I have added tests. - `/test` for `lit` tests - `/unittest` for C++ tests - `/python/test` for end-to-end tests - Select one of the following. - [x] I have not added any `lit` tests. --- A fractional autotune `top_k` describes the fraction of candidate configurations to benchmark. When `early_config_prune` was also configured, the autotuner calculated that fraction from the original configuration list instead of the surviving list. ```text 10 configured | v 6 survive early pruning | +-- old: top_k = 50% of 10 -> benchmark 5 | +-- new: top_k = 50% of 6 -> benchmark 3 ``` This change calculates fractional `top_k` from `pruned_configs`, the set produced by `early_config_prune`. Integer `top_k`, performance-model ordering, and the existing guarantee that a fractional value retains at least one configuration are unchanged. The new `test_prune_configs_fractional_top_k_after_early_prune` exercises `Autotuner.run()` and verifies that early pruning retains 6 of 10 configurations, the performance model evaluates all 6 survivors, and `top_k=0.5` benchmarks exactly the fastest 3. Before the fix: ```text assert [1, 2, 3, 4, 5] == [1, 2, 3] 1 failed ``` After the fix: ```text 1 passed ``` ## Test plan - `test_prune_configs_fractional_top_k_after_early_prune` (`1 passed`) - `pre-commit run --from-ref origin/main --to-ref HEAD` - `git diff --check` Fixes #10964.

…p code (#11008)

PR description written by Codex Synchronize participating warps after sampling the virtual cluster-barrier phase so they cannot wait for different rendezvous epochs. Cover implicit kernel-exit and explicit cluster barriers in two-CTA and four-CTA warp-specialized kernels. ## Stack Merge bottom-up: - [ ] #11010 (this PR)

Allow `@aggregate` types to participate in operator resolution with clear rules on resolution precedence for forward and reversed operators. Actual functional diff is +92, -57. Big test case adds an example on how to automatically decompose non-power-of-2 tensors using metaprogramming.

…32 (#10883) ## Summary `computation_type_impl` already upcast `float16`/`bfloat16` to `float32` for `/` and `%`, because those operations have no native form below fp32. But the fp8 rule and the scalar branch didn't apply the same logic, so `fp8 / fp8`, `fp8 % fp8`, and `scalar / fp8_tensor` kept fp8/fp16 types for division and modulo. Per lezcano's review, this PR now unifies the rule instead of patching per-dtype: a single rule between rules 1 and 2 — ```python # 2) if we have a div or a mod we upcast to `fp32` as div and mod are # not supported for floats with bitwidth < 32 if div_or_mod and (a_ty.is_floating() or b_ty.is_floating()): return tl.float32 ``` — and removes the per-dtype `div_or_mod` branches it subsumes. The scalar branch now checks `tensor_ty.is_floating()` so scalar-vs-fp8 division is covered by the same rule. `triton-semantics.rst` documents the fp32 upcast for div/mod. ## Test plan - New `test_fp8_div_mod_promotion` covers the matrix: fp8 same/mixed, fp16, bf16, int (keeps integer promotion for `//` and `%`), and scalar-vs-tensor operands, for both `/`/`%` and non-division ops. - Verified on Linux with no GPU: the full 18-case type matrix through `computation_type_impl` passes, and the parser-level `static_assert` kernel passes.

This changes how GSan models stream synchronization. Each stream maintains a ring buffer of three vector clocks, and each kernel receives a monotonically increasing `kernel_id`. - At kernel completion, each CTA publishes its vector clock using an atomic max on `stream_clock[kernel_id % 3]`. - For a normal launch, the kernel initializes its clock from `stream_clock[(kernel_id - 1) % 3]`, acquiring the previous kernel’s aggregated clock. - For a PDL launch, the kernel initializes from `stream_clock[(kernel_id - 2) % 3]`, acquiring the kernel from two launches ago. - When a PDL kernel executes `gdc_wait`, it acquires `stream_clock[(kernel_id - 1) % 3]`, synchronizing with its immediate predecessor. This also completely removes the need for the separate `synchronize_launch_stream` function.

Currently containsLocalBarrier also assumes that the barrier happens before any memory effects, so an atomic with acquire effects which has the barrier after the memory effect is not handled correctly. This splits it into 3 stages: 1. beforeMemoryEffects -> e.g. atomic with release semantics 2. betweenMemoryEffects -> e.g. a cross-warp convert layout 3. afterMemoryEffects -> e.g. atomic_poll without a timeout Doing this allows us to model the outgoing memory effects more precisely, but not unconditionally skip inserting barriers before the op.

This PR is best reviewed commit by commit! We allow reshapes of subslices. The rules are: - The dst type may not be a subslice itself - The range covered by the dst type should be a contiguous chunk of memory owned by src - An memdesc without a subview is considered contiguous in all its image. If it does not cover all its image (because it has broadcasting) that's fine, because the allocator will always allocate a full contiguous chunk of memory for a given tensor. We also tighten a number of verifiers here and there. --------- Co-authored-by: lezcano <lezcano-93@hotmail.com>

This spiraled a little bit from noting that `getAllocationShapePerCTA` was incorrect for `SharedLinearLayout`s with broadcasting. I also noted that we never really need the "allocation shape" as it's not even a well defined concept in the presence of broadcasting. Instead, we compute the number of elements that we need to allocate, which is just the number of `kOffset` bases of the layout (+ pipelining) We also kill `getMemDescSize` which had similar issues. There is probably a bigger clean-up around computing shapes and things around shapes as we have a million helpers that compute similar but subtly different quantities without using the associated LinearLayouts, but I'm leaving that one for another day.

…++ arg-eval order GCC and Clang emit the xor operands in different order, so pin the two swappable ops with CHECK-DAG. Signed-off-by: Witold Dziurdz <witold.dziurdz@intel.com>

lezcano and others added 29 commits July 20, 2026 12:12

[BACKEND] Several cross-CTA synchronisation fixes (#10975)

63f3fd7

Restores the correctness portion of #10922 for ClusterBarrierInsertion, a tiny fix in a lowering, and proper typing of the barriers in CLC.

[AMD] Fix serval gfx1250 tests (#10978)

12ebfbb

Fixes several gfx1250 gluon test issues around async waits and input ranges and cache mods.

[Frontend] Pass simple flag when constructing AnnAssign nodes (#10986)

6bf75c0

Fixes #10981

[BACKEND] Improve cluster_barrier codegen within WS (#10992)

d66def3

On pre-Rubin we just emit one arrive instruction that's executed from the first `num_ctas` threads. On Rubin, we use the new `multitcast` arrive.

reland "[Blackwell] Fix incorrectly pipelined back-to-back MMAs shari…

18f97b6

…ng the same accumulator" (#10988) reland triton-lang/triton#10858. It was revert due to a wrong bisect.

[AMD][gfx1250] Support proton clock instructions (#10813)

5b90e63

[BACKEND] Support subslices at nonaligned offsets and reinterprets of…

b5c76c8

… contiguosu subviews (#10997) PR description written by Codex As per title ## Stack Merge bottom-up: - [ ] #11006 - [ ] #10997 (this PR)

[PROTON] Fix a critical problem in runtime exit callbacks and clean u…

5fcebcb

…p code (#11008)

wdziurdz requested a review from anmyachev July 24, 2026 11:21

wdziurdz requested a review from whitneywhtsang July 24, 2026 11:21

wdziurdz force-pushed the dev/wdziurdz/merge-53 branch from 7df3c97 to a36a4cd Compare July 24, 2026 11:30

Merge commit 'c393f1ba4374019b3f565175c99b1566607d6c49'

a6b01bf

wdziurdz force-pushed the dev/wdziurdz/merge-53 branch from a36a4cd to a6b01bf Compare July 24, 2026 11:40

[TEST] Use RUBIN-DAG for cluster barrier mask checks independent of C…

40469fb

…++ arg-eval order GCC and Clang emit the xor operands in different order, so pin the two swappable ops with CHECK-DAG. Signed-off-by: Witold Dziurdz <witold.dziurdz@intel.com>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Merge OpenAI Triton commit `c393f1b`#7572

Merge OpenAI Triton commit `c393f1b`#7572
wdziurdz wants to merge 31 commits into
mainfrom
dev/wdziurdz/merge-53

wdziurdz commented Jul 24, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

15 participants

Uh oh!

Conversation

wdziurdz commented Jul 24, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

15 participants