diff options
| author | Alex MacLean <amaclean@nvidia.com> | 2025-05-22 19:38:10 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2025-05-22 19:38:10 -0700 |
| commit | 3a84a4e55d896e573fe175f34c793b0c294dec6b (patch) | |
| tree | 04d4e220b75daf5ddeb6cf93a41a5f7094d4f666 /llvm/test/Analysis/GlobalsModRef | |
| parent | 5d76555f935cfe75d1e45bd56544786cc56718c3 (diff) | |
Reland "[NVPTX] Unify and extend barrier{.cta} intrinsic support" (#141143)
Note: This relands #140615 adding a ".count" suffix to the non-".all"
variants.
Our current intrinsic support for barrier intrinsics is confusing and
incomplete, with multiple intrinsics mapping to the same instruction and
intrinsic names not clearly conveying intrinsic semantics. Further, we
lack support for some variants. This change unifies the IR
representation to a single consistently named set of intrinsics.
- llvm.nvvm.barrier.cta.sync.aligned.all(i32)
- llvm.nvvm.barrier.cta.sync.aligned.count(i32, i32)
- llvm.nvvm.barrier.cta.arrive.aligned.count(i32, i32)
- llvm.nvvm.barrier.cta.sync.all(i32)
- llvm.nvvm.barrier.cta.sync.count(i32, i32)
- llvm.nvvm.barrier.cta.arrive.count(i32, i32)
The following Auto-Upgrade rules are used to maintain compatibility with
IR using the legacy intrinsics:
* llvm.nvvm.barrier0 --> llvm.nvvm.barrier.cta.sync.aligned.all(0)
* llvm.nvvm.barrier.n --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
* llvm.nvvm.bar.sync --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
* llvm.nvvm.barrier --> llvm.nvvm.barrier.cta.sync.aligned.count(x, y)
* llvm.nvvm.barrier.sync --> llvm.nvvm.barrier.cta.sync.all(x)
* llvm.nvvm.barrier.sync.cnt --> llvm.nvvm.barrier.cta.sync.count(x, y)
Diffstat (limited to 'llvm/test/Analysis/GlobalsModRef')
| -rw-r--r-- | llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll | 19 |
1 files changed, 3 insertions, 16 deletions
diff --git a/llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll b/llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll index e92a45807ed9..7019694439bb 100644 --- a/llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll +++ b/llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll @@ -11,28 +11,15 @@ target triple = "nvptx64-nvidia-cuda" ; CHECK-LABEL: @bar_sync ; CHECK: store -; CHECK: tail call void @llvm.nvvm.bar.sync(i32 0) +; CHECK: tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0) ; CHECK: load define dso_local i32 @bar_sync(i32 %0) local_unnamed_addr { store i32 %0, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 - tail call void @llvm.nvvm.bar.sync(i32 0) + tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0) %2 = load i32, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 ret i32 %2 } -declare void @llvm.nvvm.bar.sync(i32) #0 - -; CHECK-LABEL: @barrier0 -; CHECK: store -; CHECK: tail call void @llvm.nvvm.barrier0() -; CHECK: load -define dso_local i32 @barrier0(i32 %0) local_unnamed_addr { - store i32 %0, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 - tail call void @llvm.nvvm.barrier0() - %2 = load i32, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 - ret i32 %2 -} - -declare void @llvm.nvvm.barrier0() #0 +declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #0 attributes #0 = { convergent nounwind } |