tile-ai / tile-ai/tilelang

[BUG][Fuzzer][ice-on-valid-code] int8 GEMM through an int16 shared intermediate hangs the compiler (ptxas spins indefinitely) instead of compiling — an int32 intermediate compiles in ~11s

Open
#2,397 1 comment 0 reactions 0 assignees View on GitHub
bug
Dominant language
Python
Stars
7.4k
Forks
745
Avg merge
1d 1h
Merged PRs (30d)
104

Description

### 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.11+cuda.gita4399e4c

### System information

NVIDIA A10G (sm_86), CUDA 12.8 (nvcc V12.8.93), PyTorch 2.8.0, Python 3.11. `execution_backend="cython"`.

### Problem description

An int8 GEMM whose int32 accumulator fragment is copied through an **int16 shared intermediate** (then back out to int32) makes the CUDA compile hang: ptxas runs at 100% CPU and never finishes. Changing only the intermediate dtype from `int16` to `int32` compiles in ~11s. The hang is gated entirely on the int16 narrowing of a GEMM-output fragment.

Trigger boundary:

| config | result |
|---|---|
| int8 GEMM → **int16** shared mid → int32 store | **HANG** (ptxas at 100% CPU, no completion within 150s) |
| int8 GEMM → int32 shared mid → int32 store | compiles ~11s, result equals the full int64 matmul (maxabs=0) |
| no GEMM (global int32 load) → int32→int16→int32 copies | compiles ~10s |

Two things are not the trigger:
- **The software pipeline.** A plain serial K-loop GEMM (no `T.Pipelined`) with the int16 intermediate hangs identically.
- **The int16 shared store/load on its own.** Replacing the GEMM accumulator with a plain global int32 load — same int16 shared intermediate, same int32→int16→int32 narrowing/widening — compiles fine.

What remains is specifically a **GEMM-output (int32) fragment → int16 shared narrowing copy**.

### Reproducible example code

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

M = N = K = 256
bM = bN = 128
bK = 32
MID_DT = "int16" # <-- the trigger; change to "int32" and it compiles in ~11s

@T.prim_func
def main(A: T.Tensor((M, K), "int8"), B: T.Tensor((K, N), "int8"), C: T.Tensor((M, N), "int32")):
with T.Kernel(T.ceildiv(N, bN), T.ceildiv(M, bM), threads=128) as (bx, by):
As = T.alloc_shared((bM, bK), "int8")
Bs = T.alloc_shared((bK, bN), "int8")
Cl = T.alloc_fragment((bM, bN), "int32")
Cmid = T.alloc_shared((bM, bN), MID_DT)
T.clear(Cl)
for k in T.Pipelined(T.ceildiv(K, bK), num_stages=2):
T.copy(A[by * bM, k * bK], As)
T.copy(B[k * bK, bx * bN], Bs)
T.gemm(As, Bs, Cl)
T.copy(Cl, Cmid) # int32 -> int16 (narrowing)
T.copy(Cmid, C[by * bM, bx * bN]) # int16 -> int32 (widening)

tilelang.compile(main, out_idx=[2]) # never returns — hangs in ptxas
```
(A plain `T.serial` K-loop instead of `T.Pipelined` hangs identically.)

[A12_hang.txt](https://github.com/user-attachments/files/28922297/A12_hang.txt)

### Traceback

```pytb
None — the build hangs in ptxas; there is no error.
```

### Expected behavior

Compile in bounded time (as the int32-intermediate version does), or reject this pattern with a clear error — rather than hang.

### Additional context

Isolating the toolchain stage shows the hang is in a `ptxas` optimization pass, on a legal, normal-sized
PTX:

| step | result |
|---|---|
| `nvcc --ptx` (int16 kernel) | fast — produces a legal 937-line PTX |
| `ptxas -O0` (that PTX) | compiles in seconds |
| `ptxas -O1` / `-O2` / `-O3` (that PTX) | hangs (>90 s, no output) |
| `ptxas` (int32-intermediate PTX, any `-O`) | fast |

So it's not a TileLang codegen blowup or an nvcc-frontend issue — the generated CUDA/PTX is ordinary, and
ptxas only hangs with optimization enabled. The int16 PTX differs from the int32 one almost only in
16-bit ops (254 `cvt` + vectorized `st.shared.v2.u16`/`ld.shared.v4.u16`, vs 4 in the int32 version),
which come from lowering the int32→int16 narrowing as per-element conversions. Since the int32 path
compiles, a TileLang-side workaround looks possible; the ptxas hang on legal PTX may also be worth
reporting to NVIDIA.

The hanging PTX is attached (`A12_hang.txt`, `.version 8.7` / `.target sm_86`); it reproduces standalone
without TileLang:
```
ptxas -O1 -arch=sm_86 A12_hang.txt -o /dev/null # hangs; -O0 finishes in seconds
```
ptxas: `release 12.8, V12.8.93`. (Whether it hangs is a property of ptxas, not the PTX `.version`.)

[A12_hang.txt](https://github.com/user-attachments/files/28922309/A12_hang.txt)

Contributor guide

Open the contributing guide

Research direction

Start with the reproducible Python example and the attached A12_hang.txt PTX, then compare the int16 and int32 intermediate cases by running ptxas at -O0 and -O1. Done means the narrowing GEMM-output path compiles in bounded time or rejects the pattern with a clear error, without introducing incorrect results.

Written by the indexing model from the issue text.

Assessment

Tech stack
python
Domain
compilers
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Quiet
Clarity
Mostly clear
Newbie friendliness
48/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.