[BUG][Fuzzer][ice-on-valid-code] a non-bool buffer element as a bare if/assert condition aborts with an internal is_bool() assert instead of testing != 0 #3009

Required prerequisites

• I have searched the Issue Tracker and no similar report exists.

What version of TileLang are you using?

0.1.13

System information

TileLang 0.1.13 / L40S (sm_89), CUDA 12.8. The abort happens during frontend compilation (TIR analysis), before any codegen, so it is GPU-independent.

Problem description

Guarding work with if <mask>[i]: — the block-sparse idiom the shipped examples use (examples/blocksparse_gemm, examples/blocksparse_attention) — aborts compilation as soon as the mask buffer is a non-bool type. if A[i]: on an int buffer is the ordinary Python spelling of "if this element is non-zero", but it aborts with an internal TVM assert:

tvm.error.InternalError: Check failed: (a.dtype().is_bool()) is false:

The identical kernel written with the explicit comparison if A[i] != 0: compiles and runs (control below). So the integer condition is silently accepted at parse time and only detonates later, in a downstream pass that composes the raw int condition with a boolean operator, with a message that never mentions the user's if.

The abort is the same on the @T.prim_func path and the @tilelang.jit eager path (both go through the same if visitor). Where it detonates depends on the branch shape: with an else branch the LegalizeNegativeIndex pass builds !condition and trips the bool-only Not constructor (bare Check failed: (a.dtype().is_bool())); with only a then branch the buffer-index flattener composes the condition with && and trips the logical_and RHS check (Check failed: (rhs.dtype().is_bool()) ... Expected boolean argument as RHS of && operator (logical AND), but received A[i] of type int32). Same root — the raw int32 condition is never coerced to bool — two detonation sites.

Not a regression — see Provenance.

Reproducible example code

import tilelang, tilelang.language as T

# FAIL: integer truthiness in an `if`. `if A[i]:` on an int buffer is the normal
# Python spelling of `A[i] != 0`.
@T.prim_func
def k(A: T.Tensor((8,), "int32"), B: T.Tensor((8,), "int32")):
    with T.Kernel(1, threads=8) as bx:
        i = T.get_thread_binding()
        if A[i]:            # <-- int32 condition
            B[i] = 1
        else:
            B[i] = 0
tilelang.compile(k, out_idx=[1]).get_kernel_source()   # InternalError: is_bool() ...

# PASS control: explicit `!= 0` compiles and runs — isolates the defect to the
# bare integer condition, not the surrounding kernel.
@T.prim_func
def k_ok(A: T.Tensor((8,), "int32"), B: T.Tensor((8,), "int32")):
    with T.Kernel(1, threads=8) as bx:
        i = T.get_thread_binding()
        if A[i] != 0:
            B[i] = 1
        else:
            B[i] = 0
tilelang.compile(k_ok, out_idx=[1]).get_kernel_source()  # -> compiles OK

Traceback

# with an `else` branch (LegalizeNegativeIndex builds `!condition`):
  in tvm::tl::LegalizeNegativeIndex(tvm::tirx::PrimFunc)
  in tvm::arith::IRVisitorWithAnalyzer::VisitStmt_(tvm::tirx::IfThenElseNode const*)
  File ".../3rdparty/tvm/src/tirx/ir/expr.cc", line 471, in tvm::tirx::Not::Not(tvm::PrimExpr, tvm::Span)
tvm.error.InternalError: Check failed: (a.dtype().is_bool()) is false:

# with a `then`-only branch (buffer-index flattener composes `condition && ...`):
  in tvm::tl::BufferFlattener::GetSimplifiedElemOffset(...)
  in tvm::operator&&(tvm::PrimExpr, tvm::PrimExpr)
  in tvm::logical_and(tvm::PrimExpr, tvm::PrimExpr, tvm::Span)
tvm.error.InternalError: Check failed: (rhs.dtype().is_bool()) is false:
Expected boolean argument as RHS of && operator (logical AND), but received A[i] of type int32

Expected behavior

if A[i]: on a non-boolean condition should be treated as if A[i] != 0: (Python truthiness), which the compiler already handles correctly — the != 0 control produces a working kernel. Failing that, a parse-time diagnostic naming the if and the non-boolean condition would at least point the user at the real line, instead of an internal assert in a region-analysis pass.

Additional context

Root cause. The if visitor accepts a non-boolean condition without coercing it to bool, so an integer PrimExpr becomes the IfThenElse condition; a later pass applies !/&& to that condition and asserts it must be boolean. In visit_if, tilelang/language/parser/parser.py:475-504 the evaluated predicate is passed straight into T.If(...) whenever it is a PrimExpr/ExprOp — the else clause only rejects a predicate that is neither PrimExpr nor bool, so an int32 element slips through unchecked and un-cast.

Suggested fix. In visit_if (and the sibling visit_assert), when the predicate is a PrimExpr that is not already boolean, coerce it to bool (cond != 0) before entering the T.If/T.Assert frame — matching Python truthiness, what the user's explicit != 0 already does, and what visit_while effectively already does (its lowering emits if (!(cond)) break;, so an int while compiles). Alternatively, reject a non-boolean condition at parse time with a diagnostic naming the statement — as T.if_then_else already does (if_then_else only accept the condition to be boolean type), which is a far better message than the if path's internal is_bool() assert.

Provenance. The unchecked-condition handling is in TileLang's own if visitor (tilelang/language/parser/parser.py), and the bool-only asserts it flows into are in the bundled tile-ai/tvm op/expr layer. It aborts on 0.1.13 and is not a regression introduced by a recent change. I did not pin an introducing PR.

Root, two levels. (a) SOURCE-level: visit_if, parser.py:474-504 passes the evaluated predicate straight into T.If(...) whenever it is a PrimExpr/ExprOp, with no coercion to bool; the sibling condition-takers visit_assert, parser.py:508-524 and visit_while, parser.py:201-216 share the same uncoerced-condition shape. (b) OPERATOR-level: the raw non-bool PrimExpr reaches a bundled-tvm op/stmt that asserts is_bool()/is_predicate_dtype()/==Bool() — Not (expr.cc:471), logical_and (op.cc:705-708), AssertStmt (stmt.cc:107), if_then_else (op.cc:634). Any bool-only op the uncoerced condition later flows into is a detonation site; which one fires depends on the construct and branch shape.

Generalization (0.1.13, L40S). All cells run this session, one kernel per fresh process — observed, not asserted.

Class = a non-bool buffer element (any width, signed/unsigned, or float) used as a bare condition in if / assert / not / if_then_else; all abort with a bool-only op assert. Boundary: an explicit != 0, a bool condition, and a while loop all compile (the while path does the coercion the if path omits — this is the fix precedent). if_then_else and assert reach the abort through their own tvm ops with clearer messages; they are the same defect class, worth the same one-line coercion fix, not separate bugs.

Dedup. I searched the open and closed tracker and found no existing report of an integer if condition aborting on the is_bool() assert. (Sibling-in-spirit but distinct root from the float-%/// abort: that is the int-only op-assert on floormod/floordiv; this is the bool-only op-assert on Not/&&. Different operators, asserts, and fix sites, so filed separately.)

Impact. The trigger is a bare non-bool buffer element used as a condition (if / assert / not / if_then_else), which the shipped idiom currently sidesteps by declaring every mask bool. When it fires it is a compile-time abort — loud and immediate, caught before any code is generated, so it corrupts no output and cannot reach a workload silently; the cost is that a valid if M[i]: on an integer mask won't build, and (for the if/not paths) the internal is_bool() assert names a downstream pass (LegalizeNegativeIndex or the index flattener) rather than the user's line. Fixing it closes the whole class — non-bool truthiness compiles like the explicit != 0 — and replaces the internal assert with either a working kernel or a diagnostic that points at the real line. The while path already does this coercion correctly and is the fix precedent.

Reach. Shipped examples located and run this session (0.1.13). The if <mask>[i]: idiom is shipped and I verified each cited mask is bool in the v0.1.13 tree: examples/blocksparse_gemm/example_blocksparse_gemm.py:82 guards on BlockMask, declared T.Tensor(..., "bool") at line 69; examples/blocksparse_attention/.../sparse_gqa_decode_varlen_mask.py:78 guards on block_mask, declared T.Tensor(shape_mask, T.bool) at line 42; examples/deepseek_v32/sparse_mla_fwd_pipelined.py:258 guards on is_kv_valid, allocated T.alloc_shared([BI], "bool", ...) at line 92 (the seesaw variant lines 219/251 the same). I RAN the shipped blocksparse kernel (the test_pipeline_order_stage body from testing/python/language/test_tilelang_language_pipeline.py, BlockMask: T.Tensor(..., "bool")) verbatim: it compiles (get_kernel_source() → 8700-char kernel). So every shipped site DODGES — solely because the mask is bool. To confirm the crash is one dtype away in the real kernel, I re-ran that exact shipped kernel with BlockMask changed to int8 and the if directly on the buffer element: it aborts logical_and RHS op.cc:707 — Expected boolean argument as RHS of && ... but received BlockMask[(by*2+bx)*8] of type int8. Storing a 0/1 mask as int8/uint8 (a common flag encoding) trips the bug in an otherwise-shipped kernel. CI stays green because every shipped mask is bool.