[BUG] Unsigned condition analysis constructs a negative unsigned coefficient

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.14, source build at 3a1aeb08

System information

Reproduced on Windows (Python 3.12.13, PyTorch 2.11.0+cu130) and Linux (Python 3.13, PyTorch 2.7.1+cu128), using matching TileLang source builds. The reproducer only lowers IR; device compilation and execution are disabled.

Problem description

A valid uint32 zero comparison fails during block read/write-region analysis with cannot make uint from negative value -1. There is no negative constant in the user program. The negative value is introduced by the constraint solver.

Changing the input dtype to int32 succeeds; uint64 also fails.

Reproducible example code

import tilelang
import tilelang.language as T

@T.prim_func
def kernel(A: T.Tensor((32,), 'uint32'), B: T.Tensor((2, 32), 'int32')):
    with T.Kernel(1, threads=32):
        tx = T.get_thread_binding()
        for i in T.serial(2):
            scratch = T.alloc_var('int32', init=0)
            mask = A[tx]
            B[i, tx] = T.if_then_else(mask != 0, scratch + 1, 0)

tilelang.lower(
    kernel,
    target={'kind': 'cuda', 'arch': 'sm_120'},
    enable_host_codegen=False,
    enable_device_compile=False,
)

Traceback

PlanAndUpdateBufferAllocationLocation
  -> GetSBlockReadWriteRegion
  -> ConditionalBoundsContext::TrySolveCondition
  -> SolveInequalitiesToRange
  -> SolveLinearInequalities
  -> make_const / MakeConstScalar
tvm.error.InternalError: cannot make uint from negative value -1

Expected behavior

Successful lowering: B[i, tx] = int(A[tx] != 0). All accesses are in bounds and the local variable is initialized. Unsupported constraint refinement should conservatively retain valid access regions rather than reject this program.

Additional context

The loop-local allocation triggers region analysis. The analyzer visits both branches, including the implicit else condition mask == 0. Solving this equality introduces negative coefficients, but SolveLinearInequalities creates them using the unsigned variable's dtype. This conflates the program's unsigned value domain with the solver's internal coefficient representation.

The same failure reproduces by directly calling s_tir.analysis.GetSBlockReadWriteRegion on a small equivalent SBlock, without the TileLang pass pipeline. Merely wrapping -1 as an unsigned constant would not preserve ordered-integer inequality semantics.

Related search result: #1299 has the same error text but concerns the explicit literal 0xFFFFFFFF during expression construction; this reproducer has no such literal and fails in constraint analysis instead.