llvm / llvm/llvm-project

[llvm][LoopUnroll][nvvm] Loop peeling breaks convergence of aligned CTA barriers

Open
#224,146 0 comments 0 reactions 0 assignees View on GitHub
loopoptim
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

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.