[BUG][Fuzzer][ice-on-valid-code] `T.q_multiply_shift` crashes on a runtime (non-constant) multiplier instead of computing the result

Required prerequisites

• I have searched the issue tracker and did not find an existing report for this defect.
• I have tested with the latest release.

What version of TileLang are you using?

0.1.13 (submodule 3rdparty/tvm at 8df8ebd6)

System information

TileLang 0.1.13 / CUDA 12.8 / PyTorch 2.8.0, single NVIDIA A10G (Ampere, sm_86). The crash is in a target-independent TIR legalization pass, so it is not arch-specific.

Problem description

T.q_multiply_shift(x, y, q, s) crashes the compiler whenever the multiplier y is a runtime value (e.g. a BufferLoad) rather than a compile-time integer constant. The op is a documented public primitive computing the rounded fixed-point product round(x*y*2^-s) on Q-format integers, and its signature declares both operands as PrimExpr with no constant requirement. Every invocation with a per-element y — thread-indexed, in T.Parallel, or in T.serial — aborts during legalization with:

InternalError: Check failed: (broadcast_node != nullptr) is false:

The multiplier is the only load-bearing ingredient: a runtime x with a constant y compiles fine, while a constant x with a runtime y crashes.

Not a regression — see Provenance.

Reproducible example code

import tilelang, tilelang.language as T, torch

N = 16

@T.prim_func
def main(X: T.Tensor((N,), "int32"), Y: T.Tensor((N,), "int32"), Out: T.Tensor((N,), "int32")):
    with T.Kernel(1, threads=N):
        tx = T.get_thread_binding(0)
        # Multiplier Y[tx] is a runtime value -> crashes in legalization.
        Out[tx] = T.q_multiply_shift(X[tx], Y[tx], 31, 1)

k = tilelang.compile(main, out_idx=[2])   # InternalError: Check failed: (broadcast_node != nullptr)

# --- Working control: the SAME rounded fixed-point multiply, done by hand in int64.
#     Compiles and returns the correct value, proving the crash is the op, not the harness.
@T.prim_func
def control(X: T.Tensor((N,), "int32"), Y: T.Tensor((N,), "int32"), Out: T.Tensor((N,), "int32")):
    with T.Kernel(1, threads=N):
        tx = T.get_thread_binding(0)
        s = 32                                   # total right shift = q(31) + s(1)
        prod = T.cast(X[tx], "int64") * T.cast(Y[tx], "int64")
        Out[tx] = T.cast((prod + (T.cast(1, "int64") << (s - 1))) >> s, "int32")

kc = tilelang.compile(control, out_idx=[2])
x = torch.full((N,), 2**30, dtype=torch.int32).cuda()
y = torch.full((N,), 2**30, dtype=torch.int32).cuda()
print(kc(x, y)[0].item())   # -> 268435456  == 2**28 == round(2**30 * 2**30 * 2**-32)  (correct)

Traceback

[TileLang] TileLang begins to compile kernel `main`
Fatal: Check failed: (broadcast_node != nullptr) is false:
  ...
  tvm::tl::IntrinInjecter::VisitExpr_(tvm::tir::CallNode const*)
  tvm::runtime::detail::LogFatal::~LogFatal()
tvm.error.InternalError: Check failed: (broadcast_node != nullptr) is false:

Expected behavior

Compilation should succeed and the kernel should compute round(x*y*2^-s) for a runtime multiplier, the same value the manual int64 control returns (2**28 for the inputs above). A runtime y is computable: the op's own general lowering path treats y as an arbitrary PrimExpr (it casts it to int64 and multiplies), so nothing about a non-constant multiplier is unsupported — only the compile-time power-of-2 shortcut needs a constant, and when the value is unknown the shortcut should simply be skipped. If a constant multiplier were genuinely required, the op should reject a non-constant y with a clear frontend error rather than aborting deep in legalization; the evidence suggests computing is the natural outcome, but which is appropriate is a call for the maintainers.

Additional context

Root cause. The tir.q_multiply_shift FLegalize lowering assumes the multiplier y is always a compile-time constant, so it fails on any runtime y. Before choosing between the power-of-2 fast path and the general path, the lowering unconditionally calls a helper to read y's integer value, and that helper hard-asserts the node is an IntImm or a Broadcast(IntImm) — a BufferLoad is neither, so the assert fires.

Suggested fix. Make the power-of-2 detection tolerate a non-constant y: attempt the constant extraction and, when y is not a compile-time integer, skip the fast path and fall through to the general QMultiplyShift lowering (which already handles a symbolic y). Concretely, get_int_value could return an optional / sentinel instead of asserting, and the == (1 << 30) test would then be false for a non-constant y.

Provenance. The offending guard and helper are present in the TileLang TVM fork at the pinned submodule commit 8df8ebd6 shipped with 0.1.13. The pattern derives from upstream Apache TVM's q_multiply_shift legalization and is long-standing; the exact introducing commit is not pinned here. Not a recent regression.

Dedup. No existing report covers a q_multiply_shift legalization crash; the mechanism (a constant-only guard on the multiplier) is distinct from the other intrinsic-lowering issues on record (math-intrinsic dtype coverage, T.Broadcast codegen arity, LayoutInference divide-by-zero).

This is one of three distinct defects in the same tirx.q_multiply_shift legalization (intrin_rule.cc): (1) this constant-asserting get_int_value helper (runtime y → broadcast_node != nullptr); (2) the power-of-2 fast path off-by-one at s == 1 (1 << -1 shift-range ICHECK); (3) the per-axis op forwarding an integer is_lshift_required into a boolean Select (is_bool() assert). Same file and op family, three separate root causes on different lines — this report covers only (1).

Reach. Triggering requires only that the multiplier y be a runtime value — the natural way to use the op for per-element quantized rescaling (a rescale factor loaded from a tensor). A compile-time-constant y avoids the crash, so a program that hard-codes the multiplier is unaffected. q_multiply_shift is a public exported op (tilelang/language/tir/ir.py) but is used in no shipped examples/ and exercised by no test in the repo, which is why CI does not catch it.

No-example verification (ran this session on 0.1.13): git grep q_multiply_shift over the v0.1.13 tree (excluding the bundled 3rdparty/tvm) returns only the op definition (op.py:3215), the .pyi stub, and the _op_wrapper re-exports in ast/ir.py + tir/ir.py — zero call sites under examples/, testing/, or docs/. There is therefore no shipped example or test to run for this op; the "no example exercises it" claim is confirmed by direct grep, not inherited. The op's docstring (op.py) documents it as out = round(x*y*2^-s) with both operands typed PrimExpr and no constant requirement, matching the Problem statement.

Generalization (§17) — root + 4-axis sweep, all cells re-run on 0.1.13.

Two-level root.

• SOURCE-level: the constant-asserting get_int_value lambda in the tirx.q_multiply_shift FLegalize (intrin_rule.cc L270-279): it reads an integer value to steer the power-of-2 fast path, but TVM_FFI_ICHECKs that the node is an IntImm or Broadcast(IntImm) instead of returning an optional. Any other node (a BufferLoad) aborts. It is called at two sites: on y unconditionally (L283) and on s inside the pow2 branch (L285).
• OPERATOR-level: the fixed-point-rescale op family q_multiply_shift / q_multiply_shift_per_axis, both legalized in this same file and both ultimately delegating to QMultiplyShift (L224-255).

4-axis findings (input → observed → same-root?):

Boundary / reframe decision. The root is narrow and specific: it is exactly the get_int_value helper in the scalar q_multiply_shift legalizer. The sibling op q_multiply_shift_per_axis does not share it (compiles with a runtime multiplier), and no other legalizer in the file uses the idiom. So the report is kept specific to q_multiply_shift, not reframed to a class. Within that op the root spans two call sites (y at L283, s at L285) — both verified, both closed by one sentinel fix — which the report already covers as a single root.

Distinct adjacent bug found while sweeping (not filed). With a constant non-pow2 multiplier the general path is reached, and QMultiplyShift hard-asserts y.dtype().bits() == 32 (L227); an int16 program crashes with Check failed: (y.dtype().code()==kDLInt && y.dtype().bits()==32). This is the documented "int32-only" limitation surfacing as a deep legalization ICHECK rather than a frontend dtype error — a distinct root (dtype restriction, not the constant guard). Recorded here, not filed.

Example run (PART 1). No shipped example/test exercises q_multiply_shift at v0.1.13 (grep result above), so there was no example to run; the crash is demonstrated only by the minimal repro in this report, which reproduces on 0.1.13.

Start in 3rdparty/tvm/src/target/intrin_rule.cc at get_int_value and the q_multiply_shift FLegalize calls around lines 270-285; read QMultiplyShift at lines 224-255 to understand the general path. Reproduce the issue with the provided runtime-Y example, since the repository has no shipped test or example for this op. Done means runtime multipliers compile and produce the expected rounded result without the legalization assertion.