flagos-ai / flagos-ai/FlagTree
[BUG][hygon] tl.load(..., volatile=True) causing cross-block spin-barrier deadlock on hygon
- Dominant language
- Python
- Stars
- 350
- Forks
- 149
- Avg merge
- 2d 19h
- Merged PRs (30d)
- 97
Description
# `tl.load(..., volatile=True)` Spin-Barrier Kernel deadlock on HCU
## Environment
- FlagTree: 0.6.0+hcu3.6
## Phenomenon
A cross-block synchronization (spin-barrier) kernel deadlocks on HCU when the spin-read uses `tl.load(ptr, volatile=True)`; the process stalls at device synchronization and never returns. The same kernel runs normally on NVIDIA (H20), Tianshu, and PPU platforms.
## Reproduction
Run `python test.py` to reproduce stably:
```python
import torch
import triton
import triton.language as tl
@triton.jit
def _barrier(ctr, off, NC):
tl.atomic_add(ctr + off, 1, sem="release")
while tl.load(ctr + off, volatile=True) < NC:
pass
@triton.jit
def _kern(ctr, part, out, NC: tl.constexpr, ROUNDS: tl.constexpr):
pid = tl.program_id(0)
for r in range(ROUNDS):
tl.atomic_add(part + pid, 1.0)
_barrier(ctr, r, NC)
tl.store(out + pid, tl.atomic_add(ctr + ROUNDS - 1, 0))
def main():
NC, ROUNDS = 64, 32
ctr = torch.zeros(ROUNDS, dtype=torch.int32, device="cuda")
part = torch.zeros(NC, dtype=torch.float32, device="cuda")
out = torch.zeros(NC, dtype=torch.float32, device="cuda")
_kern[(NC,)](ctr, part, out, NC=NC, ROUNDS=ROUNDS)
print("OK out[0]=", out[0].item(), flush=True)
if __name__ == "__main__":
main()
```
## Additional Information
The following are observations from the compilation artifacts of the hanging kernel (full dump files can be provided):
In TTIR/TTGIR, the spin-read carries the volatile flag:
```
%8 = tt.load %6 {isVolatile = true} : !tt.ptr
```
In LLIR, the corresponding load is an ordinary load (no volatile or cache-bypass flags), and the spin loop body is empty:
```llvm
%20 = load <1 x i32>, ptr addrspace(1) %16, align 4
```
Contributor guide
Research direction
Start by running the provided test.py reproduction on HCU and inspect the generated TTIR/TTGIR and LLIR artifacts. Trace how the volatile load and spin-loop body are lowered; done means the kernel no longer deadlocks and the LLIR preserves the required volatile behavior.
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
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100