[BugFix] Decode an unsigned sub-word without a zero-point bias - #3348

Refs #3301 (F163).

Why

_tir_packed_uint_to_uint_to_float masks its field and then subtracts 2^(nbit - 1) - 1, the zero-point of an asymmetric signed decode, even though both the storage and the destination are unsigned:

max_int_value = (1 << (nbit - 1)) - 1
return ((val >> (pos.astype(T.uint32) * tirx.const(nbit, T.uint32))) & tirx.const(
    (1 << nbit) - 1, "uint32")).astype(dtype) - tirx.const(max_int_value, dtype)

Every decoded value comes back shifted by that constant. Decoding the nibbles 0..7 of one uint32:

uint_to_uint decode : [-7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0]
expected (unsigned) : [ 0.0,  1.0,  2.0,  3.0,  4.0,  5.0,  6.0, 7.0]
int_to_int control  : [ 0.0,  1.0,  2.0,  3.0,  4.0,  5.0,  6.0, 7.0]

The signed sibling _tir_packed_int_to_int_to_float handles its sign by sign-extending rather than biasing, so the two decoders disagreed on the same payload — the bias does not belong in an unsigned decode.

What changed

Mask and shift, with no bias. A docstring states the contract, since the name and the behaviour are what disagreed.

The helper is not referenced anywhere in the tree, so nothing that currently runs changes behaviour. It lives in a module that four test files import from, and the same function ships under tilelang/quantize/quantization.py in released wheels, so it is worth leaving correct.

Verification

• Three cases added to testing/python/issue/test_tilelang_u32_signed_decode.py, which already tests a sibling decoder from the same module and whose packing helper and kernel shape are reused here. All three fail on an unmodified tree with Mismatched elements: 4 / 4 (100.0%) and similar; the three existing signed cases pass in both states.
• test_tilelang_u32_signed_decode.py, test_tilelang_issue_2947.py, test_tilelang_bf16x2_pack_special_values.py and testing/python/quantize/: 39 passed, 1 skipped.
• The new cases are marked requires_cuda and skip cleanly without a device.
• The kernel cache was cleared before each run.
• pre-commit run --files <changed files>: all 12 applicable hooks pass.

Summary

• Updated _tir_packed_uint_to_uint_to_float to decode packed fields as unsigned values without subtracting a zero-point bias.
• Added CUDA-parametrized tests for unsigned 2-, 4-, and 8-bit values packed in uint32. The tests check that decoded float32 values match the inputs.

Testing

The author reports 39 passed and 1 skipped across the listed test files and quantization tests. The author also reports that all 12 applicable pre-commit hooks passed. These results were not independently verified.

Verified end to end, together with the other contributions from this batch: VERIFICATION.md