llvm/test/Analysis/GlobalsModRef/functions_without_nosync.ll

   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
   7 target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
   8 target 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
  16 define 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
  23 declare 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
  29 define 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
  36 declare void @llvm.nvvm.barrier0() #0
  37
  38 attributes #0 = { convergent nounwind }