diff --git a/llvm/lib/Analysis/GlobalsModRef.cpp b/llvm/lib/Analysis/GlobalsModRef.cpp --- a/llvm/lib/Analysis/GlobalsModRef.cpp +++ b/llvm/lib/Analysis/GlobalsModRef.cpp @@ -951,10 +951,16 @@ if (GV->hasLocalLinkage() && !UnknownFunctionsWithLocalLinkage) if (const Function *F = Call->getCalledFunction()) if (NonAddressTakenGlobals.count(GV)) - if (const FunctionInfo *FI = getFunctionInfo(F)) - Known = unionModRef(FI->getModRefInfoForGlobal(*GV), - getModRefInfoForArgument(Call, GV, AAQI)); - + if (const FunctionInfo *FI = getFunctionInfo(F)) { + ModRefInfo SideEffects = + Call->hasFnAttr(Attribute::NoSync) ? ModRefInfo::NoModRef + : Call->onlyReadsMemory() ? ModRefInfo::Ref + : ModRefInfo::ModRef; + Known = unionModRef( + FI->getModRefInfoForGlobal(*GV), + unionModRef(SideEffects, + getModRefInfoForArgument(Call, GV, AAQI))); + } if (!isModOrRefSet(Known)) return ModRefInfo::NoModRef; // No need to query other mod/ref analyses return intersectModRef(Known, AAResultBase::getModRefInfo(Call, Loc, AAQI)); diff --git a/llvm/test/Analysis/GlobalsModRef/dead-uses.ll b/llvm/test/Analysis/GlobalsModRef/dead-uses.ll --- a/llvm/test/Analysis/GlobalsModRef/dead-uses.ll +++ b/llvm/test/Analysis/GlobalsModRef/dead-uses.ll @@ -36,7 +36,7 @@ } ; Function Attrs: nounwind -define internal void @f() { +define internal void @f() nosync { entry: %tobool = icmp ne i32 0, 0 br i1 %tobool, label %if.then, label %if.end diff --git a/llvm/test/Analysis/GlobalsModRef/func-memattributes.ll b/llvm/test/Analysis/GlobalsModRef/func-memattributes.ll --- a/llvm/test/Analysis/GlobalsModRef/func-memattributes.ll +++ b/llvm/test/Analysis/GlobalsModRef/func-memattributes.ll @@ -27,5 +27,5 @@ declare i32 @func_readonly() #0 declare i32 @func_read_argmem_only() #1 -attributes #0 = { readonly nounwind } -attributes #1 = { readonly argmemonly nounwind } +attributes #0 = { readonly nounwind nosync } +attributes #1 = { readonly argmemonly nounwind nosync } diff --git a/llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll b/llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll new file mode 100644 --- /dev/null +++ b/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" + +@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 } diff --git a/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken1.ll b/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken1.ll --- a/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken1.ll +++ b/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken1.ll @@ -8,7 +8,7 @@ ; CHECK-LABEL: @main() define dso_local i32 @main() { entry: - %tmp0 = call i8* @llvm.objc.autoreleasePoolPush() #1 + %tmp0 = call i8* @llvm.stacksave() #1 %tmp6 = load i8, i8* @deallocCalled, align 1 %tobool = icmp ne i8 %tmp6, 0 br i1 %tobool, label %if.else, label %if.end @@ -18,10 +18,10 @@ unreachable ; CHECK-LABEL: if.end: -; CHECK-NEXT: call void @llvm.objc.autoreleasePoolPop +; CHECK-NEXT: call void @llvm.stackrestore ; CHECK-NOT: load i8, i8* @deallocCalled if.end: ; preds = %entry - call void @llvm.objc.autoreleasePoolPop(i8* %tmp0) + call void @llvm.stackrestore(i8* %tmp0) %tmp7 = load i8, i8* @deallocCalled, align 1 %tobool3 = icmp ne i8 %tmp7, 0 br i1 %tobool3, label %if.end6, label %if.else5 @@ -35,10 +35,10 @@ ret i32 0 } -declare i8* @llvm.objc.autoreleasePoolPush() #1 -declare void @llvm.objc.autoreleasePoolPop(i8*) #1 +declare i8* @llvm.stacksave() #1 +declare void @llvm.stackrestore(i8*) #1 declare dso_local void @__assert_fail() #0 -attributes #0 = { noreturn nounwind } -attributes #1 = { nounwind } +attributes #0 = { noreturn nosync nounwind } +attributes #1 = { nosync nounwind } diff --git a/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken2.ll b/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken2.ll --- a/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken2.ll +++ b/llvm/test/Analysis/GlobalsModRef/intrinsic_addressnottaken2.ll @@ -14,7 +14,7 @@ ; CHECK-LABEL: @main() define dso_local i32 @main() { entry: - %tmp0 = call i8* @llvm.objc.autoreleasePoolPush() #1 + %tmp0 = call i8* @llvm.stacksave() #1 %tmp6 = load i8, i8* @deallocCalled, align 1 %tobool = icmp ne i8 %tmp6, 0 br i1 %tobool, label %if.else, label %if.end @@ -24,10 +24,10 @@ unreachable ; CHECK-LABEL: if.end: -; CHECK-NEXT: call void @llvm.objc.autoreleasePoolPop +; CHECK-NEXT: call void @llvm.stackrestore ; CHECK-NOT: load i8, i8* @deallocCalled if.end: ; preds = %entry - call void @llvm.objc.autoreleasePoolPop(i8* %tmp0) + call void @llvm.stackrestore(i8* %tmp0) %tmp7 = load i8, i8* @deallocCalled, align 1 %tobool3 = icmp ne i8 %tmp7, 0 br i1 %tobool3, label %if.end6, label %if.else5 @@ -41,10 +41,10 @@ ret i32 0 } -declare i8* @llvm.objc.autoreleasePoolPush() #1 -declare void @llvm.objc.autoreleasePoolPop(i8*) #1 +declare i8* @llvm.stacksave() #1 +declare void @llvm.stackrestore(i8*) #1 declare dso_local void @__assert_fail() #0 -attributes #0 = { noreturn nounwind } -attributes #1 = { nounwind } +attributes #0 = { noreturn nosync nounwind } +attributes #1 = { nosync nounwind } diff --git a/llvm/test/Analysis/GlobalsModRef/intrinsic_addresstaken.ll b/llvm/test/Analysis/GlobalsModRef/intrinsic_addresstaken.ll --- a/llvm/test/Analysis/GlobalsModRef/intrinsic_addresstaken.ll +++ b/llvm/test/Analysis/GlobalsModRef/intrinsic_addresstaken.ll @@ -18,7 +18,7 @@ ; CHECK-LABEL: @main() define dso_local i32 @main() { entry: - %tmp0 = call i8* @llvm.objc.autoreleasePoolPush() #1 + %tmp0 = call i8* @llvm.stacksave() #1 %tmp6 = load i8, i8* @deallocCalled, align 1 %tobool = icmp ne i8 %tmp6, 0 br i1 %tobool, label %if.else, label %if.end @@ -28,9 +28,9 @@ unreachable ; CHECK-LABEL: if.end: -; CHECK-NEXT: call void @llvm.objc.autoreleasePoolPop +; CHECK-NEXT: call void @llvm.stackrestore if.end: ; preds = %entry - call void @llvm.objc.autoreleasePoolPop(i8* %tmp0) + call void @llvm.stackrestore(i8* %tmp0) %tmp7 = load i8, i8* @deallocCalled, align 1 %tobool3 = icmp ne i8 %tmp7, 0 br i1 %tobool3, label %if.end6, label %if.else5 @@ -44,8 +44,8 @@ ret i32 0 } -declare i8* @llvm.objc.autoreleasePoolPush() #1 -declare void @llvm.objc.autoreleasePoolPop(i8*) #1 +declare i8* @llvm.stacksave() #1 +declare void @llvm.stackrestore(i8*) #1 declare dso_local void @__assert_fail() #0 attributes #0 = { noreturn nounwind } diff --git a/llvm/test/Analysis/GlobalsModRef/no-escape.ll b/llvm/test/Analysis/GlobalsModRef/no-escape.ll --- a/llvm/test/Analysis/GlobalsModRef/no-escape.ll +++ b/llvm/test/Analysis/GlobalsModRef/no-escape.ll @@ -8,14 +8,14 @@ @d = common global i32 0, align 4 @e = common global i32* null, align 4 -define void @foo(i32* %P) noinline { +define void @foo(i32* %P) noinline nosync { ; CHECK: define void @foo %loadp = load i32, i32* %P, align 4 store i32 %loadp, i32* @d, align 4 ret void } -define void @bar() noinline { +define void @bar() noinline nosync { ; CHECK: define void @bar %loadp = load i32, i32* @d, align 4 store i32 %loadp, i32* @c, align 4