[llvm][LoopUnroll][nvvm] Loop peeling breaks convergence of aligned CTA barriers
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
`LoopUnroll` peels the `%loop`/`%latch` natural loop below, splitting `llvm.nvvm.barrier.cta.sync.aligned.all` across thread-dependent paths despite its `convergent` attribute.
Originally, thread 0 reloads `%flag` and broadcasts it through shared memory; all threads execute the same two barrier sites on each retry. After peeling, thread 0 re-enters the peeled region through `%read`, while the other threads enter the residual loop. The two barrier sites become four, with different threads executing different copies. This violates [PTX's `.aligned` requirement](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar) and results in a hang.
Reproduced with unmodified LLVM 24.0.0git at `32dd4e7e8e375cf983ffcf49e72f3abd6158a1b1`:
## Reproducer
Godbolt: https://gcc.godbolt.org/z/b8rKzxvnn
Or locally:
```sh
opt -S -verify-each -passes=loop-unroll -pass-remarks=loop-unroll repro.ll -o unroll.ll
# remark: :0:0: peeled loop by 1 iterations
```
where `repro.ll` is
```llvm
target triple = "nvptx64-nvidia-cuda"
@scratch = internal addrspace(3) global i32 undef, align 4
define ptx_kernel void @poll(ptr addrspace(1) %flag) "nvvm.reqntid"="32" {
entry:
%tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
%leader = icmp eq i32 %tid, 0
br i1 %leader, label %read, label %loop
read:
%input = load atomic i32, ptr addrspace(1) %flag monotonic, align 4
br label %loop
loop:
%value = phi i32 [ %input, %read ], [ 0, %entry ], [ 0, %latch ]
call void asm sideeffect "@$2 st.shared.u32 [$0], $1;", "r,r,b,~{memory}"(ptr addrspace(3) @scratch, i32 %value, i1 %leader)
call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
%ready = load i32, ptr addrspace(3) @scratch, align 4
%done = icmp ne i32 %ready, 0
br i1 %done, label %exit, label %latch
latch:
call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
br i1 %leader, label %read, label %loop
exit:
ret void
}
```
Contributor guide
Research direction
Start with the repro.ll reproducer and run the documented opt command with loop-unroll and verify-each. Trace the LoopUnroll peeling behavior that transforms the two barrier sites into four thread-dependent paths. Done means the reproducer no longer breaks convergence of the aligned CTA barriers and the resulting unroll.ll remains valid.
Written by the indexing model from the issue text.
Assessment
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100