[BUG][Fuzzer][wrong-code] `T.copy` of a packed sub-byte buffer (`int4`/`uint4`/`float4_e2m1fn`) silently drops ~half the elements via a non-atomic cross-thread byte read-modify-write

### 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 this session on the released `0.1.13` wheel). Source on `main` re-checked at `d36ec37c596ff8e587fc67fcf1e88bcb90cc60d6`.

### System information

NVIDIA L40S (`sm_89`), CUDA 12.8, Python 3.13, tilelang 0.1.13. The store helper and the copy loop are target-independent CUDA codegen/templates (no `sm_XX` gate), so the defect is not specific to one GPU; it fires wherever the copy loop assigns the two nibbles of a byte to different threads.

### Problem description

A `T.copy` of a 128-element **sub-byte packed** buffer round-tripped through shared memory (`global → shared → global`) at `threads=128` silently returns a result where ~half the elements are wrong — ~56 of the 64 packed bytes come back corrupted — even though a copy of a buffer back to itself must be byte-for-byte identical. Nothing errors; the wrong output is deterministic. This affects **every 2-elements-per-byte dtype TileLang packs**: `int4`, `uint4`, and `float4_e2m1fn` (fp4) all reproduce identically at runtime. The same copy at `threads=32`, and the same copy on the byte-wide `int8` at `threads=128`, both return the exact input.

Two adjacent sub-byte elements share one storage byte (2 nibbles/byte). The scalar store lowers to a packed helper (`tl_int4_packed_store` / `tl_uint4_packed_store` / `tl_fp4_packed_store`) that is a **non-atomic** read-modify-write of that byte. When the copy loop hands element `2i` and element `2i+1` to different threads (which is what `threads=128` over a 128-element tile does — thread `t` handles logical index `t`), both threads read the same byte, each rewrites its own nibble, and each writes the whole byte back with no atomicity and no barrier between them; one thread's whole-byte store clobbers the other thread's nibble. About half the elements are lost.

Observed mismatch — three dtypes, runtime-verified; and the sub-byte boundary

```
TRIGGER (threads=128, 128-elem identity copy, packed bytes must be identical):
int4 -> MISMATCH ~56–57/64 bytes (runtime)
uint4 -> MISMATCH ~56–57/64 bytes (runtime)
float4_e2m1fn -> MISMATCH ~56–57/64 bytes (runtime)

CONTROL (same kernel, one variable changed):
int4 threads=32 -> byte-identity (vectorized: both nibbles on one thread)
int8 threads=128 -> byte-identity (one byte per element, no sharing)
```

Boundary — which sub-byte dtypes reach this path:
- **2 elements/byte → hits the race:** `int4`, `uint4`, `float4_e2m1fn` (dispatched to a packed byte-RMW store at `codegen_cuda.cc` — int4/uint4 at L5297, fp4 at L5316/L5329). All three runtime-confirmed with the identical ~56–57/64 signature.
- **>2 elements/byte → compile-rejected before the store:** `int2`/`uint2` (4/byte) and `uint1` (8/byte) fail the storage-lowering guard `data_bits % 8 == 0` ("Need to load/store by multiple of bytes"), so they never reach the racy path.
- `int1` is int32-backed and `bool` is one-byte-per-element — neither is 2+-per-byte packed, both safe.

So the family is exactly the three 2-per-byte packed dtypes; the denser packings are already gated out at storage lowering.

The generated CUDA for the trigger contains exactly this per-element scalar store (no vectorization) and **no** `__syncthreads`, atomic, `cp.async.bulk`, `mbarrier`, or TMA descriptor — it is the plain SIMT scalar-store path:

```
tl_int4_packed_store((signed char*)As, ((int)threadIdx.x), tl_int4_packed_load((const signed char*)A, ((int)threadIdx.x)));
tl_int4_packed_store((signed char*)B, ((int)threadIdx.x), tl_int4_packed_load((const signed char*)As, ((int)threadIdx.x)));
```

The index is `threadIdx.x`, so thread `t` writes logical `int4` index `t`; threads `2i` and `2i+1` both read-modify-write shared byte `As[i]` (`idx >> 1`). The loser's whole-byte store overwrites the winner's nibble → the surviving byte holds one nibble and a zero in the other, matching `first@[8] got=0 exp=16`.

This is not a regression — the mechanism has been present since the `int4` packed store was introduced (see Provenance).

### Reproducible example code

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

@T.prim_func
def k(A: T.Tensor((128,), "int4"), B: T.Tensor((128,), "int4")):
with T.Kernel(1, threads=128):
As = T.alloc_shared((128,), "int4")
T.copy(A[0:128], As) # g -> s
T.copy(As, B[0:128]) # s -> g ; B must equal A byte-for-byte
# threads=128 over 128 elements => thread t handles int4 index t,
# so threads 2i and 2i+1 both RMW the shared byte As[i].

m = tilelang.compile(k, out_idx=[1])

# 64 packed bytes = 128 int4; every nibble is a valid int4 value in 0..7
# (no sign/range excuse). int4 storage is int8, so feed/check as uint8 bytes.
vals = [(((b // 8) % 8) << 4) | (b % 8) for b in range(64)]
a = torch.tensor(vals, dtype=torch.uint8, device="cuda").view(torch.int8)
b = m(a)
print("int4 threads=128 identity:", torch.equal(b.view(torch.uint8), a.view(torch.uint8)))
# -> False (56/64 bytes differ; first at byte 8: got 0x00, expected 0x10)

# CONTROL 1 — same kernel, only the thread count changes: correct.
@T.prim_func
def k32(A: T.Tensor((128,), "int4"), B: T.Tensor((128,), "int4")):
with T.Kernel(1, threads=32):
As = T.alloc_shared((128,), "int4")
T.copy(A[0:128], As); T.copy(As, B[0:128])
print("int4 threads=32 identity:", torch.equal(
tilelang.compile(k32, out_idx=[1])(a).view(torch.uint8), a.view(torch.uint8)))
# -> True

# CONTROL 2 — same structure on a full-byte dtype at threads=128: correct.
@T.prim_func
def k8(A: T.Tensor((128,), "int8"), B: T.Tensor((128,), "int8")):
with T.Kernel(1, threads=128):
As = T.alloc_shared((128,), "int8")
T.copy(A[0:128], As); T.copy(As, B[0:128])
a8 = torch.arange(128, dtype=torch.int32).remainder(127).to(torch.int8).cuda()
print("int8 threads=128 identity:", torch.equal(
tilelang.compile(k8, out_idx=[1])(a8).view(torch.uint8), a8.view(torch.uint8)))
# -> True
```

### Traceback

```
No traceback — the kernel compiles and runs to completion; the result is silently wrong and deterministic (56/64 packed bytes differ from the input).
```

### Expected behavior

A `global → shared → global` `T.copy` of a packed sub-byte buffer should return the input unchanged, exactly as the `int8` copy and the `threads=32` copy already do — a copy of a tensor back to itself is byte-identity. `int4`/`uint4`/`float4_e2m1fn` are first-class registered dtypes with dedicated packed load/store codegen, and `T.copy` places no dtype restriction, so the natural expectation is a correct copy regardless of the thread partition. Since the `threads=32` and `int8` siblings show the copy is computable, the primary resolution is to compute it correctly (layout inference keeps a byte's elements on one thread). If that is out of scope for an arbitrary partition, the acceptable fallback is a compile-time rejection of the byte-splitting layout — the same fail-loud treatment int2/uint2 already get at storage lowering — never a silent half-dropped result.

### Additional context

**Root cause.** The copy loop's layout inference partitions logical elements across threads without modelling a constraint that sub-byte dtypes impose: **two elements that share one physical storage byte must be owned by the same thread**, because the byte is the smallest writable unit and the packed store is a whole-byte read-modify-write. Just as fp16/fp32 carry an alignment constraint and MMA operands carry a distribution constraint, a 2-per-byte dtype carries a *byte-integrity* constraint on the store side — and layout inference does not account for it. So at `threads=128` it happily assigns logical `idx` and `idx+1` (sharing byte `idx>>1`) to different threads, and the two threads' non-atomic whole-byte stores clobber each other's nibble. The correct siblings satisfy the constraint by accident, not by design: `threads=32` vectorizes to a wider store that keeps both nibbles on one thread, and `int8` has one byte per element so there is nothing to share.

Mechanism (the lowering path that the missing constraint reaches)

`T.copy` lowers via [`CopyNode::MakeSIMTLoop`](https://github.com/tile-ai/tilelang/blob/d36ec37c596ff8e587fc67fcf1e88bcb90cc60d6/src/op/copy.cc#L452-L491), which emits [one `BufferStore` per logical element](https://github.com/tile-ai/tilelang/blob/d36ec37c596ff8e587fc67fcf1e88bcb90cc60d6/src/op/copy.cc#L487) with no sub-byte awareness (its only sub-byte guard is the FP4-unpack `LOG_FATAL` at [`copy.cc#L453-L458`](https://github.com/tile-ai/tilelang/blob/d36ec37c596ff8e587fc67fcf1e88bcb90cc60d6/src/op/copy.cc#L453-L458), which does not cover the 2-per-byte store path). Each scalar sub-byte store then lowers at [`codegen_cuda.cc#L5042-L5057`](https://github.com/tile-ai/tilelang/blob/d36ec37c596ff8e587fc67fcf1e88bcb90cc60d6/src/cuda/codegen/codegen_cuda.cc#L5042-L5057) to a packed helper — `tl_int4_packed_store` / `tl_uint4_packed_store` ([`common.h#L283-L299`](https://github.com/tile-ai/tilelang/blob/d36ec37c596ff8e587fc67fcf1e88bcb90cc60d6/src/tl_templates/cuda/common.h#L283-L299), `packed[idx>>1] = (byte & ~mask) | nibble;`) or `tl_fp4_packed_store` — each a non-atomic byte read-modify-write. When layout inference has split the two nibbles across threads, both RMW the same byte with no atomic and no barrier, so one whole-byte store clobbers the other's. All three faces (`int4`/`uint4`/`float4_e2m1fn`) are runtime-confirmed.

**Suggested fix.** Teach the copy-loop layout inference the sub-byte byte-integrity constraint: for a 2-per-byte dtype, the store-side partition must keep all elements of a physical byte on one thread (element-pair granularity), so the whole-byte write is uncontended — this is exactly what the correct `threads=32` sibling already achieves by vectorizing to a wider store, so the fix is to make layout inference *guarantee* that invariant rather than satisfy it by luck. If honouring it for an arbitrary partition is out of scope, the fallback is to **reject the layout** that splits a byte across threads at compile time (as int2/uint2 are already rejected at storage lowering) rather than emit a silent race; a last resort is making the packed store an atomic sub-word RMW, but that pays an atomic per element to paper over a partition that should not have been produced.

**Distinct from #2563.** #2563 is a compile-time `ICHECK(0)` abort on the Hopper **TMA-descriptor** path — [`to_CUtensorMapDataType`](https://github.com/tile-ai/tilelang/blob/d36ec37c596ff8e587fc67fcf1e88bcb90cc60d6/src/op/utils.cc#L189-L204) has no width case for 4-bit integers, so an `int4` TMA bulk copy aborts before running (`ice-on-valid-code`, still open). This report is a **silent runtime miscompile on the non-TMA SIMT scalar-store path** (the generated code contains zero TMA/descriptor/`cp.async.bulk`/`mbarrier`): a different root, a different path, and a different symptom. They are complementary — #2563's own suggested fix routes `int4` copies away from TMA and into exactly the SIMT path this report indicts.

**Provenance.** Introduced whole by PR [#2073](https://github.com/tile-ai/tilelang/pull/2073) ([`15309f5cba`](https://github.com/tile-ai/tilelang/commit/15309f5cba), 2026-04-22, ancestor of `main`): `git log -L` shows both the `tl_int4_packed_store` helper (`common.h`) and the `tl_int4_packed_store(...)` codegen emit line (`codegen_cuda.cc`) were added in that commit, already with the non-atomic byte RMW. There is no prior version that copied packed `int4` correctly, so this is not a regression — the race has been present since the packed `int4` store first shipped.

**Dedup.** I searched the open and closed tracker (`int4`, `nibble`, `packed`, `race`, `atomic`) and found no existing report of a cross-thread read-modify-write race in the `int4`/`uint4` packed store; the other `int4`-tagged issues are a different mechanism (an `int8` within-thread lane-pack mask bug, a 16-byte CUDA vector-type misaligned load, and sub-word sign-decode value bugs). Distinct from #2563 (see above).

**Reach.** The rare ingredient is a `T.copy` (or scalar store loop) over an `int4`/`uint4` buffer whose loop layout splits a byte across threads: `grep -rn "int4"` over `examples/` shows the shipped `int4` usage is `examples/gemm_int4` (`T.gemm`), whose shared loads vectorize the contiguous dimension within a thread and so keep both nibbles of a byte on one thread — that path does not race (consistent with #2563's own "RUN OK, bit-exact" M=N=K=1024 `int4` GEMM on the non-TMA copy path). Given a copy whose partition assigns adjacent nibbles to adjacent threads, the common default (one element per thread, e.g. `threads=128` over a 128-element tile) makes it fire on essentially every byte; no shipped example uses `T.copy` on a bare `int4` buffer at that layout, and no test exercises the sub-byte cross-thread store path, which is why CI is green.

Run the provided Python reproducer with int4 at threads=128, then inspect CopyNode::MakeSIMTLoop in src/op/copy.cc and the packed-store lowering in src/cuda/codegen/codegen_cuda.cc and src/tl_templates/cuda/common.h. Done means the int4, uint4, and float4_e2m1fn identity copies preserve every packed byte, or the byte-splitting layout is rejected at compile time; retain coverage for the threads=32 and int8 controls.

Written by the indexing model from the issue text.