[BUG][Fuzzer][ice-on-valid-code] `int2`/`uint2` (any integer width ∉ `{1,4,8,16,32,64}`) aborts CUDA codegen with `Cannot convert type int2` instead of compiling

### Required prerequisites

- [x] I have read the documentation .
- [x] I have searched the [issue tracker](https://github.com/tile-ai/tilelang/issues) and could not find an existing report of this defect.

### What version of TileLang are you using?

0.1.13 (also present on `main` — the codegen switch below is unchanged there).

### System information

TileLang 0.1.13; CUDA 12.8; NVIDIA L40S (sm_89). The failure is in host-side C++ codegen, so it is architecture-independent (no kernel is launched).

### Problem description

A scalar cast to the 2-bit integer dtype `int2` (or `uint2`) aborts compilation with an internal fatal error, `Cannot convert type int2 to CUDA type`, even though the frontend accepts the dtype and the immediately neighbouring sub-byte integer widths — `int1`, `int4`, `int8` — all compile from the same kernel.

The CUDA type printer's integer branch is a `switch (t.bits())` with cases for `1, 4, 8, 16, 32, 64` and a `default` that sets `fail = true`. A 2-bit integer has `bits() == 2`, which matches no case, falls through to `default`, and reaches the final `LOG(FATAL)`. There is no case for width 2. The same hole is width-general: any frontend-admitted integer width not in that case set (`int2`, `int3`, `int24`, …) hits the identical fatal — confirmed by running `int3` and `int24` (both abort the same way). `int2`/`uint2` is the meaningful instance because it is a real sub-byte dtype the project already uses (as a bitnet quantization storage tag).

The value that reaches the printer is legal C++: the sibling case for width 4 already maps a *scalar* `int4` onto `signed char` (a widening storage type), so an in-register 2-bit integer has the same natural target (`signed char`) — the printer simply never handles that width.

No traceback field would show a kernel fault; this is a compile-time `InternalError` raised before any code is emitted.

### Reproducible example code

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

N = 128

@tilelang.jit(out_idx=[1])
def prog(dt):
@T.prim_func
def main(A: T.Tensor((N,), "int8"), C: T.Tensor((N,), "int8")):
with T.Kernel(1, threads=N) as bx:
i = T.get_thread_binding(0)
C[i] = T.Cast("int8", T.Cast(dt, A[i])) # narrow to dt, widen back
return main

a = torch.arange(N, dtype=torch.int8).cuda()
for dt in ["int4", "int2", "uint2"]: # int4 = handled sibling; int2/uint2 = bits()==2
try:
prog(dt)(a)
print(f"{dt:6s} compiles")
except Exception as e:
print(f"{dt:6s} {type(e).__name__}: {str(e).splitlines()[0]}")
```

Output:

```
int4 compiles
int2 InternalError: Cannot convert type int2 to CUDA type
uint2 InternalError: Cannot convert type uint2 to CUDA type
```

### Traceback

No traceback from a running kernel — compilation aborts. The error is `tvm.error.InternalError: Cannot convert type int2 to CUDA type`, raised from the `LOG(FATAL)` at the end of `CodeGenTileLangCUDA::PrintType`.

### Expected behavior

`int2`/`uint2` reaches the printer through the same path as its neighbours `int1`/`int4`/`int8`, so it would either be printed (a `case 2` mapping the scalar onto a widening storage type such as `signed char`, as the existing `case 4` scalar branch already does), or — if a 2-bit scalar element is not intended to be materialised — rejected earlier with a clean, user-facing diagnostic rather than an internal `LOG(FATAL)` inside codegen. The current behaviour is an internal abort on a dtype the frontend admits and the adjacent widths accept.

### Additional context

**Root cause.** `CodeGenTileLangCUDA::PrintType` has no case for a 2-bit integer: its integer branch dispatches on `switch (t.bits())` with cases `1, 4, 8, 16, 32, 64` only, so `bits() == 2` falls through `default` to the terminal fatal.

Where it falls through (v0.1.13 source)

The integer branch opens the switch at [`src/cuda/codegen/codegen_cuda.cc#L934`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_cuda.cc#L934). The scalar `case 4:` maps a sub-byte scalar onto `signed char` at [`#L952-L956`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_cuda.cc#L952-L956). A 2-bit width matches none of the cases and hits [`default: fail = true;`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_cuda.cc#L1078-L1080), then the final `LOG(FATAL) << "Cannot convert type " << t << " to CUDA type"` at [`#L1090`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_cuda.cc#L1090). On `main` the same switch still has cases `1, 4, 8, 16, 32, 64` and no case for 2.

**Suggested fix.** Add a `case 2:` alongside the existing `case 4:` in `PrintType`'s integer switch, emitting the same scalar widening (`signed char` / `char` for the unsigned form), so `int2`/`uint2` follows the `int1`/`int4`/`int8` path. If a 2-bit scalar is deliberately out of scope, reject it at the frontend so the failure is a clean diagnostic rather than a codegen `LOG(FATAL)` (not verified end-to-end).

**Provenance.** The integer `switch (t.bits())` is inherited from the TVM-derived CUDA type printer and has never carried a width-2 case; blame on the current lines resolves only to relocation/namespace refactors (e.g. #2166, #2297), not to the origin of the case set. Introduced no later than the fork's initial CUDA codegen; the origin before that is in upstream TVM and is unverified. Not a regression — the width was never handled.

**Dedup.** I searched the open and closed tracker and found no existing report of a scalar `int2`/`uint2` codegen abort. The related sub-byte `PrintType` issue #2480 is a distinct mechanism (a 32-lane `int8`/`uint8` *vector* whose `make_longlong4` arity is wrong, since it is stored as `longlong4`); this report is the scalar path where `bits()==2` matches no switch case at all.

**Reach.** The `int2` dtype string parses and the neighbouring integer widths `int1`/`int4`/`int8` are codegen-handled — the dtype is admitted by the frontend and its siblings are printed. In the shipped examples `int2` appears only as a quantization *storage tag* — `examples/bitnet-1.58b/utils_quant.py:55` sets `W_dtype="int2"`, and the bitnet kernels pack four 2-bit weights into an `int8` byte and decode them to `int8` on-chip. I ran the shipped example verbatim on 0.1.13 — `examples/bitnet-1.58b/kernel_benchmark/tilelang_bitnet_158_int8xint2_decode.py` (its `__main__` calls `assert_bitnet_158_int8xint2_decode_correctness(1, 256, 256, int8, int32, int32)`): it **compiles and runs to completion (exit 0, `torch.testing.assert_close` passes)** and never hits the printer. The reason it dodges the bug is explicit in the kernel: `storage_dtype = T.int8`, `num_bits = 2`, `B_shape = (N, K // 8 * 2)` — the 2-bit weights are packed into `int8` bytes and the only `int2` in the generated CUDA is inside a C-level decode helper (`int2b_t → int8b_t`), so no TIR value ever carries `bits()==2` into `PrintType`. No test exercises a scalar `int2`/`uint2` (grep of `testing/` at v0.1.13: zero `int2`/`uint2` sites; examples: only the `utils_quant.py` storage tag). The crash therefore surfaces when a 2-bit integer is used as an actual element type rather than a packed storage width, which the shipped examples avoid — which is why CI is green.

**Impact.** The trigger is narrow: a scalar element of exactly 2-bit integer width (`int2`/`uint2`), which the shipped examples never materialise (they pack `int2` into `int8` bytes). When it does fire the failure is a compile-time `LOG(FATAL)` inside codegen — loud, deterministic, and caught immediately before any code is emitted, so it corrupts no output and cannot reach a running kernel silently; the cost is that this one valid kernel refuses to build with an internal error rather than a clean diagnostic. Fixing it closes a sub-byte width gap so `int2`/`uint2` follows the same printer path as its `int1`/`int4`/`int8` neighbours.

§17 Generalization record

**Two-level root.**
- **Source-level root:** `CodeGenTileLangCUDA::PrintType`'s integer branch — the `switch (t.bits())` at `src/cuda/codegen/codegen_cuda.cc#L934`, whose cases are `{1, 4, 8, 16, 32, 64}`. Any other width hits `default: fail = true` (`#L1078`) → terminal `LOG(FATAL)` (`#L1090`). The fragility is a whitelist-of-widths with a fatal default.
- **Operator-level root:** any op that lets a `DataType` whose `bits()` is not in that whitelist reach `PrintType` as an element type (scalar or vector). Here that op is `T.Cast(dt, …)`; the frontend admits the dtype string, so the width reaches codegen unfiltered.

**Example-run result (PART 1).** Ran the cited shipped example verbatim on 0.1.13: `examples/bitnet-1.58b/kernel_benchmark/tilelang_bitnet_158_int8xint2_decode.py` → **exit 0, compiles and runs, `assert_close` passes, printer never reached.** It DODGES the bug because `int2` there is a packed storage tag (`storage_dtype=T.int8`, 4×2-bit packed per byte); the only `int2` in the emitted CUDA is inside a C-level `int2b_t→int8b_t` decode helper, never a TIR element dtype. `utils_quant.py:55` (`W_dtype="int2"`) is likewise a BitBLAS matmul-config tag, not a TIR element. No `testing/` site uses `int2`/`uint2` at all. So the inherited "int2 appears in examples" claim is TRUE but the examples never materialise a 2-bit *element* → they cannot trip the printer.

**4-axis sweep** (each cell its own fresh process + fresh `TILELANG_CACHE_DIR`; kernel `C[i] = Cast("int8", Cast(dt, A[i]))` unless noted; observed, not asserted):

| axis | cell tested | observed result | same-root? |
|---|---|---|---|
| related-type | `uint2` (bits==2, unsigned) | `InternalError: Cannot convert type uint2 to CUDA type` | **same** — hole is the *width*, signedness irrelevant |
| related-type | `int1` / `uint1` | both **compile** | same family (case 1 present) — bounds the hole |
| related-type | `int4` / `uint4` | both **compile** | same family (case 4 present; its scalar `signed char` branch is what `case 2` should mirror) |
| related-type | `int8` / `int16` / `bool` | all **compile** | same family — crash is a missing case, not a broad sub-byte defect |
| related-type | `int3` | `InternalError: Cannot convert type int3 to CUDA type` | **same** — width-general, not int2-specific |
| related-type | `int24` | `InternalError: Cannot convert type int24 to CUDA type` | **same** — any width ∉ {1,4,8,16,32,64} |
| related-operator | vector `int2x8` (`T.vectorized(8)` cast) | `InternalError: Cannot convert type int2x8 to CUDA type` | **same** — no `case 2` for any lane count |
| related-source | float branch of the SAME `PrintType`, `switch(t.bits())` cases `{16,32,64}` (`#L808`) | not reachably triggerable — sub-32 float dtypes (fp16/bf16/fp8/fp4) are all caught by earlier `is_float8`/`is_float4`/`is_bfloat16` branches before this switch | **distinct-adjacent** (same fragile idiom, hole not reachable for shipped dtypes) |
| similar-logic | `CUDAMath`/`CUDAFastMath` int `switch(t.bits())` cases `{32,64}` (`#L124/#L150`) | on an uncased width returns `""` → SILENT identity passthrough, NOT a fatal | **distinct mechanism** (same whitelist idiom, but silent not loud — already tracked separately as the math-intrin empty-mangle passthrough) |

**Class hypothesis (confirmed & broadened):** "missing-case in `PrintType`'s integer `switch(t.bits())`" — a legal, frontend-admitted integer *width* matches no case and falls through to the terminal fatal. Not int2-specific: `int2`/`uint2`/`int3`/`int24` and vector `int2x8` all reproduce; `int2`/`uint2` is the meaningful trigger (a real dtype the project uses as a storage tag). **Reframed** title + Problem to the width-general class while keeping `int2`/`uint2` as the concrete instance.

**Distinct-adjacent, NOT filed here:**
- The float-branch whitelist (`#L808`) shares the fragile idiom but has no reachable hole for shipped float dtypes.
- The `CUDAMath` mangler whitelist is a *silent* passthrough (returns `""`), a different failure mode; already tracked as a separate finding, not merged.
- #2480 is a 32-lane `int8`/`uint8` **vector** with wrong `make_longlong4` arity — a wrong-value/arity defect in an *existing* case, not a missing case. Different mechanism.
- B113/B045 are the `int4x2` **lane** hole (a lane count of a width the switch *does* handle); this B114 is the integer *width* hole. Same class ("missing-case in `PrintType`"), different missing entry — do **not** merge.

**No new distinct bug found while sweeping** — every reproducing cell (`uint2`/`int3`/`int24`/`int2x8`) is the same missing-width root as the headline. The only newly-verified fact is that the crash is width-general, not `int2`-specific (folded into the headline, not a separate bug).

Start in src/cuda/codegen/codegen_cuda.cc at CodeGenTileLangCUDA::PrintType, especially the integer switch around lines 934-1090 and the existing scalar int4 handling. Run the provided Cast reproducer for int2 and uint2, then verify that the affected widths compile without the InternalError and that existing int1, int4, and int8 behavior remains unchanged.

Written by the indexing model from the issue text.