[PATCH] D115302: GlobalsModRef should treat functions w/o nosync conservatively.
Artem Belevich via Phabricator via llvm-commits
llvm-commits at lists.llvm.org
Wed Dec 8 12:15:35 PST 2021
tra updated this revision to Diff 392867.
tra added a comment.
Renamed the test.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D115302/new/
https://reviews.llvm.org/D115302
Files:
llvm/lib/Analysis/GlobalsModRef.cpp
llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll
Index: llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll
===================================================================
--- /dev/null
+++ llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll
@@ -0,0 +1,37 @@
+; RUN: opt -globals-aa -gvn -S < %s | FileCheck %s
+;
+; Functions w/o `nosync` attribute may communicate via memory and must be
+; treated conservatively.
+
+target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+ at s = internal local_unnamed_addr addrspace(3) global i32 undef, align 4
+
+; CHECK-LABEL: @bar_sync
+; CHECK: store
+; CHECK: tail call void @llvm.nvvm.bar.sync(i32 0)
+; CHECK: load
+define dso_local i32 @bar_sync(i32 %0) local_unnamed_addr {
+ store i32 %0, i32* addrspacecast (i32 addrspace(3)* @s to i32*), align 4
+ tail call void @llvm.nvvm.bar.sync(i32 0)
+ %2 = load i32, i32* addrspacecast (i32 addrspace(3)* @s to i32*), 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, i32* addrspacecast (i32 addrspace(3)* @s to i32*), align 4
+ tail call void @llvm.nvvm.barrier0()
+ %2 = load i32, i32* addrspacecast (i32 addrspace(3)* @s to i32*), align 4
+ ret i32 %2
+}
+
+declare void @llvm.nvvm.barrier0() #0
+
+attributes #0 = { convergent nounwind }
Index: llvm/lib/Analysis/GlobalsModRef.cpp
===================================================================
--- llvm/lib/Analysis/GlobalsModRef.cpp
+++ llvm/lib/Analysis/GlobalsModRef.cpp
@@ -913,6 +913,10 @@
ModRefInfo ConservativeResult =
Call->onlyReadsMemory() ? ModRefInfo::Ref : ModRefInfo::ModRef;
+ if (Function *F = Call->getCalledFunction())
+ if (!F->hasNoSync())
+ return ConservativeResult;
+
// Iterate through all the arguments to the called function. If any argument
// is based on GV, return the conservative result.
for (auto &A : Call->args()) {
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D115302.392867.patch
Type: text/x-patch
Size: 2061 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20211208/470540d6/attachment.bin>
More information about the llvm-commits
mailing list