flagos-ai / flagos-ai/FlagTree

[BUG][hygon] tl.load(..., volatile=True) causing cross-block spin-barrier deadlock on hygon

Open
#1,030 1 comment 0 reactions 0 assignees View on GitHub
hcu
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.