[BUG][Fuzzer][wrong-code] CuTeDSL codegen sign-extends a `T.Cast` from an unsigned narrow int (uint8/uint16/uint32) to a wider int or float, instead of zero-extending

### Required prerequisites

- [x] I have read the documentation .
- [x] I have searched the Issue Tracker that this hasn't already been reported.

### What version of TileLang are you using?

0.1.13 (reproduced on 0.1.13 with nvidia-cutlass-dsl 4.3.1).

### System information

NVIDIA L40S (`sm_89`), CUDA 12.8, Python 3.13, tilelang 0.1.13, nvidia-cutlass-dsl 4.3.1. The defect is in the target-independent CuTeDSL scalar-cast codegen, so it is not specific to one GPU.

### Problem description

On the **CuTeDSL backend** (`target="cutedsl"`), a `T.Cast` from an **unsigned** narrow integer buffer element (`uint8`, `uint16`, `uint32`) to a wider integer (`int16`/`int32`/`int64`) or to a float **sign-extends** the source value instead of zero-extending it. Any source value with the top bit set (e.g. a `uint8` >= 128) is silently turned negative. No error, no warning — wrong numbers.

Expected vs actual:

| source (uint8) | expected int32 | actual (CuTeDSL) |
|---|---|---|
| 128 | 128 | **-128** |
| 200 | 200 | **-56** |
| 255 | 255 | **-1** |

Blast radius — same root across widths and targets, plus the two controls that isolate it (all confirmed on sm_89, 0.1.13)

| cast | wrong? |
|---|---|
| uint8 -> int32 | yes (200 -> -56) |
| uint8 -> int16 | yes (131 -> -125) |
| uint8 -> **float32** | yes (128 -> **-128.0**) |
| uint16 -> int32 | yes (38400 -> -27136) |
| uint32 -> int64 | yes (2^31 -> -2^31) |
| uint8 -> uint8 (copy, control) | **no** — load alone is fine |
| int8 -> int32 (signed source, control) | **no** — sign-extend is correct here |

### Reproducible example code

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

@T.prim_func
def cast_u8_i32(A: T.Tensor((256,), "uint8"), C: T.Tensor((256,), "int32")):
with T.Kernel(1, threads=256) as bx:
i = T.get_thread_binding()
C[i] = T.Cast("int32", A[i])

k = tilelang.compile(cast_u8_i32, target="cutedsl", out_idx=-1)
a = (torch.arange(256, dtype=torch.int32, device="cuda") % 256).to(torch.uint8)
c = k(a)
# A[128]=128 -> got -128 (expected 128); 128 of 256 values wrong (all >= 128)
print((c != a.to(torch.int32)).sum().item(), "mismatches") # -> 128
```

Emitted CuTeDSL source (the whole kernel body):

```python
C.iterator[tid] = cutlass.Int32(A.iterator[tid])
```

### Traceback

No traceback — the kernel compiles and runs to completion; the result is silently wrong and deterministic.

### Expected behavior

A `T.Cast` from an unsigned narrow integer to a wider integer or to float should **zero-extend**, so `uint8` 200 becomes `int32` 200, not -56 — matching the source dtype's unsigned value. Sign-extension is correct only for a signed source (`int8 -> int32`), which the control confirms is already handled correctly.

### Additional context

**Root cause.** A scalar `T.Cast` widening an unsigned integer takes two steps, and both drop the unsignedness:

1. The unsigned buffer element is loaded as a *signed* MLIR element. The backend does this on purpose and says so — the `BroadcastNode` visitor at [`codegen_cutedsl.cc#L388-L400`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_cutedsl.cc#L388-L400) comments: *"CuTeDSL/MLIR normalizes unsigned integer tensor loads to signed types (e.g., Uint8 pointer -> i8 tensor elements)"* and prints `Int(bits)` for a `uint` source. So a `uint8` 200 is already the signed value -56 before the cast runs.
2. The scalar cast then falls through to the Python base ([`codegen_cutedsl.cc#L571-L572`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_cutedsl.cc#L571-L572) → [`codegen_py.cc` `CastFromTo_`, L589-L597](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_py.cc#L589-L597)), which emits `PrintType(target)(value)` — i.e. `cutlass.Int32(value)` — with **no signedness handling**. Widening the already-signed -56 sign-extends it.

Masking the value to its unsigned range *before* the widening cast makes the result correct (`C[i] = T.Cast("int32", A[i]) & 255` → all 256 values correct, verified), which localizes the root to the signed reinterpretation of the uint load, not a zero-masked load. Note the same visitor *does* have an integer-widening special case just above ([L480-L502](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/cuda/codegen/codegen_cutedsl.cc#L480-L502)), but it only fires when the cast's operand is a `shift_left`/`shift_right` call — a plain buffer load does not match, so this cast is unguarded.

**Suggested fix.** When the *source* dtype is an unsigned integer narrower than the target, the CuTeDSL `CastNode` scalar path (or `CastFromTo_`) should zero-extend: mask the value to `(1 << from.bits()) - 1` before constructing the target type, or route through the matching unsigned CuTeDSL type first (e.g. `cutlass.Int32(cutlass.Uint8(value))`). The `BroadcastNode` uint->signed remap is the same hazard from the other direction and may warrant auditing together.

**Provenance.** The signedness-free `CastFromTo_` emit and the scalar-cast delegation are present in the 0.1.13 sources (permalinks above) and reproduced at runtime on 0.1.13 this session (`uint8` 200 → -56, 128/256 values wrong); origin not bisected below 0.1.13.

**Dedup.** I searched the open and closed Issue Tracker and found no existing report of this CuTeDSL scalar unsigned-widening sign-extension defect. It is distinct from #2482 (the CUDA `_tir_u32_to_int_to_float` intrinsic) and #2489 (sm_100 int32 broadcast).

**Same root, wider scope — comparisons and min/max, no cast needed.** The unsigned→signed load normalization affects *every* op whose result depends on signed-vs-unsigned interpretation, not just cast. I enumerated that set on 0.1.13 (`uint8` operands `200`, `50`):

Full signedness-scope on 0.1.13 (uint8 200 = top-bit-set, other = 50)

| op | expected | CuTeDSL | affected? |
|---|---|---|---|
| `T.Cast("int32", a)` | 200 | **-56** | yes (this issue) |
| `a < b` (200<50) | False | **True** | **yes** |
| `a > b` (200>50) | True | **False** | **yes** |
| `a <= b` (200<=50) | False | **True** | **yes** |
| `T.min(a,b)` | 50 | **-56** | **yes** |
| `T.max(a,b)` | 200 | **50** | **yes** |
| `a >> 1` (200>>1) | 100 | 100 | no |
| `a // 3`, `a % 3` (uint8→uint8) | 66, 2 | 66, 2 | no |
| `a + b`, `a - b`, `a * b`, `a == b`, `a & b`, `a \| b` | — | correct | no |

The comparison and min/max cases are **more severe than the cast because no cast is needed to trigger them** — a plain `T.max`/`T.min` or `<`/`>`/`<=` on a `uint8` buffer with values above 127 silently uses the signed interpretation (e.g. `max(200,50)` emits `tl.max(, ...)` → 50). The unaffected ops confirm the boundary: `+ - * == & |` are signedness-independent at the bit level, and `>>`/`//`/`%` on `uint8→uint8` stay in the range where signed and unsigned agree. So the full defect scope is **widening cast + comparison + min/max**; a signedness-aware fix covers all of them. All reproduced this session on 0.1.13.

**Reach.** Unsigned narrow-int inputs widened to a wider int or to float are common in quantize/dequantize and byte-processing kernels on the CuTeDSL backend. The trigger is any `T.Cast` from an unsigned narrow int whose value has the top bit set (a `uint8` ≥ 128, etc.); values below the signed range are unaffected. The comparison / min / max faces above share this one root and need no cast to trigger.

**Impact.** The trigger needs one ingredient — an unsigned narrow-int value with the top bit set — which is ordinary in quantize/byte kernels where such values are the common case, not an edge case. When it fires there is no error and the result is deterministic: every affected value flips to a negative number. Because the same signed reinterpretation also reaches comparisons, `min`, and `max` with no cast at all, ordinary clamp/compare logic on unsigned data silently takes the wrong branch.

Run the provided uint8-to-int32 CuTeDSL reproducer first, then inspect BroadcastNode at src/cuda/codegen/codegen_cutedsl.cc#L388-L400 and the scalar cast path around L480-L502 and L571-L572, plus CastFromTo_ in src/cuda/codegen/codegen_py.cc#L589-L597. Trace how unsigned loads reach widening casts, comparisons, and min/max, and verify completion against the listed uint8/uint16/uint32 cases and signed-source controls.

Written by the indexing model from the issue text.