[BUG][Fuzzer][accepts-invalid] `T.any_sync` & the warp-vote/block-sync family silently truncate a float predicate to int (no error)

[Hughshine]

Required prerequisites

• I have read the documentation https://tilelang.com.
• I have searched the Issue Tracker that this hasn't already been reported. (comment there if it has.)

What version of TileLang are you using?

0.1.11 (built from source at cd37ed5).

System information

NVIDIA L40S (sm_89), Python 3.13, torch 2.11.0+cu130 (CUDA 13.0), tilelang 0.1.11 (tag v0.1.11 = cd37ed5fc35ae7a60a1277c8eb49028174ac51e6). The defect is in the Python builtin wrappers plus the int-predicate CUDA/HIP builtins they lower to; it is architecture-independent (the wrappers apply no dtype check on any target) — reproduced on sm_89 this session.

Problem description

T.any_sync over 32 lanes each holding the float32 value 0.5 returns 0 — even though 0.5 != 0, so the documented contract ("Non-zero if any thread in the mask has a non-zero predicate") says it must be non-zero. The kernel compiles and runs to completion with no error; the wrong value is deterministic and value-dependent (0.4 → false, 3.7 → true).

The cause is a missing dtype guard. The whole warp-vote / block-sync-with-predicate family — T.any_sync, T.all_sync, T.ballot_sync, T.ballot, T.syncthreads_count, T.syncthreads_and, T.syncthreads_or (7 wrappers, one root) — accepts a float predicate with no dtype check and emits it verbatim into the CUDA builtin whose predicate parameter is int, so C++ truncates the float toward zero and 0.5 becomes int 0 = "false":

Out[threadIdx.x] = __any_sync((uint)4294967295, A[threadIdx.x]);   // A is float32

A control in the same script isolates the defect to the un-guarded float predicate: T.any_sync(A[tx] != 0.0) (a bool predicate) over the same 0.5 lanes returns 1.

This is present since the family was added (#1858, 2026-04-13); not a regression.

Family confirmation — all 7 ops, all lanes = 0.5 vs lane 0 = 3.7

Each op run over 32 lanes. all-0.5 = every lane holds 0.5; lane0=3.7 = lane 0 holds 3.7 (→ int 3), the rest 0.5 (→ int 0).

op	all lanes 0.5	lane 0 = 3.7	contract
any_sync	0	1	non-zero if ANY lane non-zero → both should be non-zero
all_sync	0	0	non-zero only if ALL lanes non-zero → all-0.5 should be non-zero
ballot_sync	0	1 (bit 0 set)	bitmask of lanes with non-zero predicate → all-0.5 should be 0xFFFFFFFF
ballot	0	1	full-warp wrapper of ballot_sync
syncthreads_count	0	1	count of threads with non-zero predicate → all-0.5 should be 32
syncthreads_and	0	0	non-zero only if ALL non-zero → all-0.5 should be non-zero
syncthreads_or	0	1	non-zero if ANY non-zero → both should be non-zero

The lane0=3.7 column is a non-vacuity check: the any-style ops (any_sync/ballot*/syncthreads_count/syncthreads_or) flip to non-zero when a lane truncates to a non-zero int (3.7 → 3), so the op is live and value-dependent, not stuck at 0. The all-style ops (all_sync/syncthreads_and) stay 0 because the other 31 lanes still truncate 0.5 → 0. Every cell that should be non-zero for all-0.5 reports 0 — the fractional truncation, not the op.

Reproducible example code

import tilelang
import tilelang.language as T
import torch

def build(guard):
    @tilelang.jit
    def k(A: T.Tensor((32,), "float32"), Out: T.Tensor((32,), "int32")):
        with T.Kernel(1, threads=32):
            tx = T.get_thread_binding(0)
            if guard:
                Out[tx] = T.any_sync(A[tx] != 0.0)   # bool predicate (control)
            else:
                Out[tx] = T.any_sync(A[tx])          # raw float predicate
    return k

A = torch.full((32,), 0.5, device="cuda", dtype=torch.float32)   # all lanes 0.5 (non-zero!)
Out = torch.zeros(32, dtype=torch.int32, device="cuda")
build(False)(A, Out); torch.cuda.synchronize()
print("any_sync(float 0.5)  ->", Out[0].item())   # -> 0  (WRONG; 0.5 != 0, so should be non-zero)

Out2 = torch.zeros(32, dtype=torch.int32, device="cuda")
build(True)(A, Out2); torch.cuda.synchronize()
print("any_sync(A != 0.0)   ->", Out2[0].item())  # -> 1  (control, correct)

Traceback

No traceback — the kernel compiles and runs to completion; the result is silently wrong and deterministic (reproduced twice, identical).

Expected behavior

A float predicate is not a declared input here: each wrapper's signature is predicate: int | PrimExpr and its docstring says "Integer condition to test", and the underlying __any_sync(mask, predicate) / __syncthreads_* builtins take an int predicate. So T.any_sync and its family should reject a float predicate at the frontend with a clear message, rather than silently truncating it — the project's own T.if_then_else already rejects a float condition ("only accept the condition to be boolean type"), and legalize_safe_memory_access likewise errors on a non-boolean condition, so the reject pattern already exists in-tree. The predicate's dtype is available at the wrapper (DataType(predicate.dtype).is_float()), so the check is a one-liner; only is_float need be rejected — any integer width (int8/16/32/64) and bool are legitimate and need no distinction, since the builtin tests != 0 and width doesn't change the truth value. (If float masks are ever intended to be supported, the wrapper would need to define the truth-value semantics explicitly — e.g. coerce predicate != 0 — rather than let C++ truncate; but that is a feature choice, whereas the current silent truncation of a non-declared input is the defect.)

Additional context

Root cause. Each of the 7 wrappers calls tirx.call_intrin(...) on the predicate with no dtype check, so a float PrimExpr flows unmodified into codegen, which emits the CUDA __*_sync / __syncthreads_* builtin whose predicate parameter is int. The wrappers are in tilelang/language/builtin.py: any_sync L1022-L1039, and the block-sync trio syncthreads_count/_and/_or at L1102-L1120, each a one-line return tirx.call_intrin(...) on the predicate. all_sync (L1042-L1058), ballot_sync (L1061-L1075), and ballot (L1078-L1085) have the identical shape. Notably each docstring already declares the argument "Integer condition to test", so the int contract is real but unenforced.

Suggested fix. Add a frontend dtype check in the wrappers before the call_intrin — if DataType(predicate.dtype).is_float(): raise ValueError(...) — rejecting a float predicate the way T.if_then_else already rejects a float condition. Integer widths and bool pass unchanged (the builtin tests != 0; width is irrelevant). One check in the shared path covers all 7 wrappers. (Proposed, not built.)

Provenance. The whole family (tl.any_sync/all_sync/ballot_sync/ballot/syncthreads_count/_and/_or) was introduced in #1858 "Add TIR builtins for warp-level vote and block-level predicate sync" (merged 2026-04-13), with no dtype guard from the start; this is not a regression. (git log -S "tl.any_sync" -- tilelang/language/builtin.py points to that single commit.)

Related. Same missing-frontend-dtype-guard family as other reject-hardening reports, but a distinct surface (the warp-vote / block-sync intrinsics). Its closest kin is the T.dp4a non-int8 case (a generic/unguarded op silently accepting a dtype it never declared and returning garbage rather than rejecting) — same "should-reject, silently passed" shape, different op. It differs from the loud reject-hardening siblings (which abort at compile) and from #2569 (T.shfl_* on fp16/bf16 fails to compile — a different op and a loud face, not a silent predicate truncation). I searched the open and closed tracker for any_sync, ballot, syncthreads_count, and "float predicate" and found no existing report of this defect.