Skip to content

[BACKEND] Restrict broadcast-mul-reduce combine to f32 - #11164

Open
Chennesxu wants to merge 1 commit into
triton-lang:mainfrom
Chennesxu:fix-combine-reduce-match
Open

[BACKEND] Restrict broadcast-mul-reduce combine to f32#11164
Chennesxu wants to merge 1 commit into
triton-lang:mainfrom
Chennesxu:fix-combine-reduce-match

Conversation

@Chennesxu

Copy link
Copy Markdown
Contributor

CombineBroadcastMulReducePattern rewrites:

tl.sum(x[:, :, None] * y[None, :, :], axis=1)

into:

tt.dot(x, y)

Its isAddF32 predicate accepted any floating-point arith.addf of 32 bits or fewer, but the rewrite hardcodes an f32 zero scalar while typing the accumulator tensor from the source element type. For f16, this produces an ill-typed splat that fails verification:

%zero = arith.constant 0.0 : f32
%acc = tt.splat %zero : f32 -> tensor<32x32xf16>

An ordinary half-precision kernel therefore fails to compile:

@triton.jit
def kernel(...):
    x = tl.load(px + ox)  # (32, 16) fp16
    y = tl.load(py + oy)  # (16, 32) fp16
    r = tl.sum(x[:, :, None] * y[None, :, :], axis=1)
    tl.store(po + oo, r)

The predicate and rewrite have disagreed since the pattern was added in #1889, so no narrow-precision reduction could produce a valid tt.dot. This change therefore removes no working path: it matches only f32, the one type the current rewrite handles validly, while narrower reductions remain as tt.reduce instead of producing invalid IR.

Adds an f16 negative lit test to test/Triton/combine.mlir.

CombineBroadcastMulReducePattern matched any float addition of 32 bits or
fewer, but its rewrite hardcodes an f32 zero scalar while typing the
accumulator from the source element type. A half-precision reduction
therefore built `tt.splat : (f32) -> tensor<MxNxf16>`, which fails the
verifier, so a valid kernel using
`tl.sum(x[:, :, None] * y[None, :, :], axis=1)` in f16 failed to compile.

The predicate and rewrite have disagreed since the pattern was added in
triton-lang#1889, so no narrow-precision reduction could produce a valid tt.dot.
Match only f32, the only type the current rewrite handles validly, and
leave narrower reductions as tt.reduce instead of producing invalid IR.
@Chennesxu
Chennesxu requested a review from ptillet as a code owner August 4, 2026 08:11
// The rewrite below creates an f32 zero scalar, so it can only handle f32
// reductions without producing an ill-typed tt.splat.
if (auto addf = dyn_cast_or_null<arith::AddFOp>(op))
return addf.getType().getIntOrFloatBitWidth() <= 32;

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I remembered that this is a legacy pattern that we may consider to deprecate. cc @ThomasRaoux

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants