1; RUN: opt -aa-pipeline=globals-aa -passes='require<globals-aa>,gvn' -S < %s | FileCheck %s 2; RUN: opt -aa-pipeline=basic-aa,globals-aa -passes='require<globals-aa>,gvn' -S < %s | FileCheck %s 3; 4; Functions w/o `nosync` attribute may communicate via memory and must be 5; treated conservatively. Taken from https://reviews.llvm.org/D115302. 6 7target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64" 8target triple = "nvptx64-nvidia-cuda" 9 10@s = internal local_unnamed_addr addrspace(3) global i32 undef, align 4 11 12; CHECK-LABEL: @bar_sync 13; CHECK: store 14; CHECK: tail call void @llvm.nvvm.bar.sync(i32 0) 15; CHECK: load 16define dso_local i32 @bar_sync(i32 %0) local_unnamed_addr { 17 store i32 %0, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 18 tail call void @llvm.nvvm.bar.sync(i32 0) 19 %2 = load i32, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 20 ret i32 %2 21} 22 23declare void @llvm.nvvm.bar.sync(i32) #0 24 25; CHECK-LABEL: @barrier0 26; CHECK: store 27; CHECK: tail call void @llvm.nvvm.barrier0() 28; CHECK: load 29define dso_local i32 @barrier0(i32 %0) local_unnamed_addr { 30 store i32 %0, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 31 tail call void @llvm.nvvm.barrier0() 32 %2 = load i32, ptr addrspacecast (ptr addrspace(3) @s to ptr), align 4 33 ret i32 %2 34} 35 36declare void @llvm.nvvm.barrier0() #0 37 38attributes #0 = { convergent nounwind } 39