[BUG][Fuzzer][ice-on-valid-code] `T.popcount` aborts with `Unresolved call ir.Op(name="tirx.popcount")` on every integer dtype except `uint32`/`uint64` (all signed widths and `uint8`/`uint16`) instead of counting set bits

### Required prerequisites

- [x] I have read the documentation .
- [x] I have searched the [Issue Tracker](https://github.com/tile-ai/tilelang/issues) that this hasn't already been reported. (comment there if it has.)

### What version of TileLang are you using?

0.1.13 (latest release)

### System information

TileLang 0.1.13 / NVIDIA L40S (`sm_89`) / CUDA 13.0 / PyTorch 2.13 / Python 3.13. The failure is in CUDA source codegen (the dtype→intrinsic-name mapping), before any arch-specific lowering, so it is architecture-independent.

### Problem description

`T.popcount` fails to compile on a **signed** integer buffer with an `InternalError: Unresolved call ir.Op(... name="tirx.popcount")`, while the same kernel on the corresponding **unsigned** dtype compiles and returns the correct bit count. `popcount` counts set bits in a bit pattern — a well-defined operation on any integer regardless of C signedness — and the docstring ("Count the number of set bits in input `x`") places no signedness restriction on the input, so the signed case should compile to the same `__popc`/`__popcll` and return the same result.

The sibling bit-counting intrinsic `T.clz` compiles fine on `int32`, so it is `popcount`'s dtype dispatch specifically that rejects signed integers.

dtype matrix (run this session on sm_89)

| dtype | `T.popcount` |
|---|---|
| `uint32` | compiles, correct (`__popc`) — got `[0,1,2,3,8,31,32,24]`, matches `bin().count("1")` |
| `int32` | **ICE** — `Unresolved call ir.Op(..."tirx.popcount")` |
| `uint64` | compiles (`__popcll`) |
| `int64` | **ICE** |
| `int16` | **ICE** |

Control: `T.clz` on `int32` compiles.

Not a regression — see Provenance.

### Reproducible example code

```python
import tilelang
import tilelang.language as T

def build(dtype):
@T.prim_func
def f(A: T.Tensor((256,), dtype), Out: T.Tensor((256,), dtype)):
with T.Kernel(1, threads=256):
for i in T.Parallel(256):
Out[i] = T.popcount(A[i])
return tilelang.compile(f, out_idx=[1])

build("uint32") # OK -> lowers to __popc, returns the correct bit count
build("int32") # InternalError: Unresolved call ir.Op(name="tirx.popcount")
```

### Traceback

InternalError raised in CUDA codegen (int32)

```
File ".../tilelang/engine/lower.py", line 251, in device_codegen
return resolve_device_codegen(target).lower(device_mod, target, compile_device=True)
File ".../tilelang/backend/device_codegen.py", line 22, in build
return tvm.ffi.get_global_func(global_func_name)(mod, target)
...
in tvm::codegen::CodeGenTileLangCUDA::AddFunction(...)
in tvm::codegen::CodeGenC::VisitExpr_(tvm::tir::CallNode const*, std::ostream&)
tvm.error.InternalError: Unresolved call ir.Op(name="tirx.popcount")
```

### Expected behavior

`T.popcount` on a signed integer should compile and return the number of set bits in the value's bit pattern, identically to the unsigned dtype of the same width. On the same inputs, `uint32` popcount already lowers to `__popc` and returns the correct count on hardware (verified this session: `[0,1,3,7,255,0x7fffffff,0xffffffff,0xdeadbeef]` → `[0,1,2,3,8,31,32,24]`), so the operation is computable at this width; the signed case takes no other path, it is simply dropped by the dtype dispatch. The sibling `T.clz` accepts signed integers, so `popcount` accepting them would also be internally consistent.

### Additional context

**Root cause.** The CUDA popcount lowering does not handle signed integer dtypes — it maps only `uint32`/`uint64` to `__popc`/`__popcll` and returns an empty intrinsic name for everything else, so codegen has no symbol to emit and aborts. `CUDAPopcount::operator()` guards the whole mapping behind `if (t.is_uint())` and otherwise `return ""` — in the bundled TVM fork this is `3rdparty/tvm/src/target/cuda/intrin_rule_cuda.cc` (the `is_uint()` guard at line 117, empty-name fallthrough at 124/127), registered for the CUDA `tirx.popcount` intrinsic via `DispatchPureExtern` at lines 228-229 of that file. An empty name is what `DispatchPureExtern` yields when no intrinsic is found, which surfaces as `Unresolved call ir.Op(name="tirx.popcount")`. TileLang also carries an identical (but now unregistered / dead) copy of this struct in its own tree at [`src/cuda/codegen/intrin_rule_cuda.cc#L92-L105`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/intrin_rule_cuda.cc#L92-L105) (`is_uint()` guard at L94); the live lowering is the fork copy.

**Suggested fix.** One direction is to widen the guard so a signed integer of the same width takes the same branch — `__popc` reinterprets its argument's bits, so `__popc((unsigned)x)` gives the correct count for a 32-bit signed value (likewise `__popcll` for 64-bit). Whether to also lower the sub-word widths (`int8`/`int16`/`uint8`/`uint16`, which take the `default: return ""` path as well) is a separate question depending on whether TileLang means to support popcount below 32 bits.

**Provenance.** Not a TileLang regression: the `is_uint()`-only mapping is inherited from the upstream TVM CUDA intrinsic rules and predates the current file. The sibling `tir.clz` was later routed through the signed-aware `CUDAMath` dispatcher (upstream apache/tvm #16952, "Enhance CLZ intrinsic support"), which is why `clz` accepts signed integers today while `popcount` was left on the unsigned-only mapping.

**Generalization (tested this session, 0.1.13 / L40S).** This is a *missing-lowering* class keyed on the `if (t.is_uint())` guard plus the `>= 32-bit` mapping, NOT "signed vs unsigned". Every integer dtype run in its own fresh process (`popcount(7)` → expect `3`):

| dtype | result | | dtype | result |
|---|---|---|---|---|
| `int8` | ICE | | `uint8` | ICE |
| `int16` | ICE | | `uint16` | ICE |
| `int32` | ICE | | `uint32` | **OK (3)** |
| `int64` | ICE | | `uint64` | **OK (3)** |

So the boundary is exactly "`uint32`/`uint64` lower; everything else aborts" — all four signed widths AND `uint8`/`uint16` fail. The common suspicion "signed fails, unsigned works" is **refuted**: `uint8`/`uint16` fail too. The single root is the CUDA popcount mapping only handling `uint32`→`__popc` / `uint64`→`__popcll` and returning `""` (→ `Unresolved`) for all other widths/signedness. Title and Root cause are framed to that real boundary.

Also turned the *operator* knob (sibling bit-intrinsics): `T.clz(int32)` **compiles** (returns 29 for 7), while `T.popcount(int32)` aborts — a direct in-tree confirmation of the Provenance note that `clz` was routed through the signedness-aware `CUDAMath` dispatcher while `popcount` was left on the unsigned-only mapping. (`T.ctz` is not exposed.) So the defect is specific to `popcount`'s mapping, not to bit-counting intrinsics in general.

**Dedup.** I searched the open and closed tracker (`popcount`, `popcount signed`, incl. PRs) and found no existing report of this defect. The closest relative is #2597 (fast-math `T.__exp`/`__log`/… on integer/`float8`), a different op family in the same `intrin_rule_cuda.cc` whose mangler likewise returns an empty name for out-of-domain dtypes; that one manifests as silent-identity or an undefined symbol, whereas popcount is an unconditional ICE for every signed integer.

**Reach.** The trigger is `T.popcount` on any integer dtype other than `uint32`/`uint64` — i.e. all signed widths and `uint8`/`uint16`. `popcount` is documented with a signedness-agnostic contract ("count the number of set bits", parameter typed `PrimExpr`/`_T -> _T`), and `int32` is the default integer dtype a user reaches for, so hitting the abort is the natural case (only the two 32/64-bit unsigned spellings dodge it). It appears in no shipped code: `grep -rn popcount examples/ testing/` returns 0 sites in each, which is why CI is green despite every non-`uint32`/`uint64` invocation failing.

**Impact.** The trigger is narrow — `T.popcount` on any signed (or sub-word) integer dtype, which no shipped kernel currently uses — but when it fires it is a compile-time `InternalError`: it aborts the whole compile deterministically for every such input, is caught immediately, and corrupts no data (nothing reaches the device silently). Fixing it closes the signed/sub-word dtype class for `popcount` and removes the disagreement with its own signedness-agnostic docstring and with the sibling `T.clz`, which already accepts signed integers; it does not unblock any existing workload.

Start in 3rdparty/tvm/src/target/cuda/intrin_rule_cuda.cc at CUDAPopcount and compare its dtype dispatch with the signed-aware CLZ path; the similar src/cuda/codegen/intrin_rule_cuda.cc copy is unregistered. Reproduce the provided dtype matrix, then verify that supported integer widths compile without an unresolved tirx.popcount and return the expected set-bit counts.

Written by the indexing model from the issue text.