48 lines · plain
1; RUN: opt -passes=aa-eval -aa-pipeline=nvptx-aa,basic-aa -print-all-alias-modref-info < %s -disable-output 2>&1 \2; RUN: | FileCheck %s --check-prefixes CHECK-ALIAS3 4target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"5target triple = "nvptx64-nvidia-cuda"6 7;;CHECK-ALIAS-LABEL: Function: test_sideeffect8;;CHECK-ALIAS: Both ModRef: Ptr: i32* %0 <-> call9define void @test_sideeffect(ptr %out) {10entry:11 %0 = addrspacecast ptr %out to ptr addrspace(1)12 call void asm sideeffect "membar.gl;", ""()13 store i32 5, ptr addrspace(1) %0, align 414 ret void15}16 17;;CHECK-ALIAS-LABEL: Function: test_indirect18;;CHECK-ALIAS: Both ModRef: Ptr: i32* %0 <-> %1 = call19define i32 @test_indirect(ptr %out) {20entry:21 %0 = addrspacecast ptr %out to ptr addrspace(1)22 store i32 0, ptr addrspace(1) %0, align 423 %1 = call i32 asm "ld.global.u32 $0, [$1];", "=r,*m"(ptr addrspace(1) elementtype(i32) %0)24 store i32 0, ptr addrspace(1) %0, align 425 ret i32 %126}27 28;;CHECK-ALIAS-LABEL: Function: test_memory29;;CHECK-ALIAS: Both ModRef: Ptr: i32* %0 <-> %1 = call30define i32 @test_memory(ptr %out) {31entry:32 %0 = addrspacecast ptr %out to ptr addrspace(1)33 store i32 0, ptr addrspace(1) %0, align 434 %1 = call i32 asm "ld.global.u32 $0, [$1];", "=r,l,~{memory}"(ptr addrspace(1) %0)35 store i32 0, ptr addrspace(1) %0, align 436 ret i32 %137}38 39;;CHECK-ALIAS-LABEL: Function: test_no_sideeffect40;;CHECK-ALIAS: NoModRef: Ptr: i32* %0 <-> %1 = call41define void @test_no_sideeffect(ptr %in, ptr %out) {42entry:43 %0 = addrspacecast ptr %out to ptr addrspace(1)44 %1 = call i32 asm "cvt.u32.u64 $0, $1;", "=r,l"(ptr %in)45 store i32 %1, ptr addrspace(1) %0, align 446 ret void47}48