diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -415,8 +415,8 @@
 
 def int_objc_autorelease                    : Intrinsic<[llvm_ptr_ty],
                                                         [llvm_ptr_ty]>;
-def int_objc_autoreleasePoolPop             : Intrinsic<[], [llvm_ptr_ty]>;
-def int_objc_autoreleasePoolPush            : Intrinsic<[llvm_ptr_ty], []>;
+def int_objc_autoreleasePoolPop             : Intrinsic<[], [llvm_ptr_ty], [IntrNoSync]>;
+def int_objc_autoreleasePoolPush            : Intrinsic<[llvm_ptr_ty], [], [IntrNoSync]>;
 def int_objc_autoreleaseReturnValue         : Intrinsic<[llvm_ptr_ty],
                                                         [llvm_ptr_ty]>;
 def int_objc_copyWeak                       : Intrinsic<[],
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
@@ -913,6 +913,9 @@
   ModRefInfo ConservativeResult =
       Call->onlyReadsMemory() ? ModRefInfo::Ref : ModRefInfo::ModRef;
 
+  if (!Call->hasFnAttr(Attribute::NoSync))
+    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()) {
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
@@ -39,6 +39,6 @@
 declare void @llvm.objc.autoreleasePoolPop(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
@@ -45,6 +45,6 @@
 declare void @llvm.objc.autoreleasePoolPop(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/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
diff --git a/llvm/test/Transforms/ObjCARC/basic.ll b/llvm/test/Transforms/ObjCARC/basic.ll
--- a/llvm/test/Transforms/ObjCARC/basic.ll
+++ b/llvm/test/Transforms/ObjCARC/basic.ll
@@ -8,8 +8,8 @@
 declare i8* @llvm.objc.unsafeClaimAutoreleasedReturnValue(i8*)
 declare void @llvm.objc.release(i8*)
 declare i8* @llvm.objc.autorelease(i8*)
-declare i8* @llvm.objc.autoreleaseReturnValue(i8*)
-declare void @llvm.objc.autoreleasePoolPop(i8*)
+declare i8* @llvm.objc.autoreleaseReturnValue(i8*) #2
+declare void @llvm.objc.autoreleasePoolPop(i8*) #2
 declare i8* @llvm.objc.autoreleasePoolPush()
 declare i8* @llvm.objc.retainBlock(i8*)
 
@@ -3073,6 +3073,7 @@
 !4 = !DIFile(filename: "path/to/file", directory: "/path/to/dir")
 !5 = !{i32 2, !"Debug Info Version", i32 3}
 
-; CHECK: attributes [[NUW]] = { nounwind }
-; CHECK: attributes #1 = { nofree nosync nounwind readnone speculatable willreturn }
+; CHECK-DAG: attributes [[NUW]] = { nounwind }
+; CHECK-DAG: attributes #{{.*}} = { nosync nounwind }
+; CHECK-DAG: attributes #{{.*}} = { nofree nosync nounwind readnone speculatable willreturn }
 ; CHECK: ![[RELEASE]] = !{}