1*9880d681SAndroid Build Coastguard Worker; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s 2*9880d681SAndroid Build Coastguard Worker 3*9880d681SAndroid Build Coastguard Worker 4*9880d681SAndroid Build Coastguard Workerdeclare i8 @llvm.nvvm.ldu.global.i.i8.p1i8(i8 addrspace(1)* %ptr, i32 %align) 5*9880d681SAndroid Build Coastguard Workerdeclare i32 @llvm.nvvm.ldu.global.i.i32.p1i32(i32 addrspace(1)* %ptr, i32 %align) 6*9880d681SAndroid Build Coastguard Workerdeclare i8 @llvm.nvvm.ldg.global.i.i8.p1i8(i8 addrspace(1)* %ptr, i32 %align) 7*9880d681SAndroid Build Coastguard Workerdeclare i32 @llvm.nvvm.ldg.global.i.i32.p1i32(i32 addrspace(1)* %ptr, i32 %align) 8*9880d681SAndroid Build Coastguard Worker 9*9880d681SAndroid Build Coastguard Worker 10*9880d681SAndroid Build Coastguard Worker; CHECK: func0 11*9880d681SAndroid Build Coastguard Workerdefine i8 @func0(i8 addrspace(1)* %ptr) { 12*9880d681SAndroid Build Coastguard Worker; ldu.global.u8 13*9880d681SAndroid Build Coastguard Worker %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1i8(i8 addrspace(1)* %ptr, i32 4) 14*9880d681SAndroid Build Coastguard Worker ret i8 %val 15*9880d681SAndroid Build Coastguard Worker} 16*9880d681SAndroid Build Coastguard Worker 17*9880d681SAndroid Build Coastguard Worker; CHECK: func1 18*9880d681SAndroid Build Coastguard Workerdefine i32 @func1(i32 addrspace(1)* %ptr) { 19*9880d681SAndroid Build Coastguard Worker; ldu.global.u32 20*9880d681SAndroid Build Coastguard Worker %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1i32(i32 addrspace(1)* %ptr, i32 4) 21*9880d681SAndroid Build Coastguard Worker ret i32 %val 22*9880d681SAndroid Build Coastguard Worker} 23*9880d681SAndroid Build Coastguard Worker 24*9880d681SAndroid Build Coastguard Worker; CHECK: func2 25*9880d681SAndroid Build Coastguard Workerdefine i8 @func2(i8 addrspace(1)* %ptr) { 26*9880d681SAndroid Build Coastguard Worker; ld.global.nc.u8 27*9880d681SAndroid Build Coastguard Worker %val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1i8(i8 addrspace(1)* %ptr, i32 4) 28*9880d681SAndroid Build Coastguard Worker ret i8 %val 29*9880d681SAndroid Build Coastguard Worker} 30*9880d681SAndroid Build Coastguard Worker 31*9880d681SAndroid Build Coastguard Worker; CHECK: func3 32*9880d681SAndroid Build Coastguard Workerdefine i32 @func3(i32 addrspace(1)* %ptr) { 33*9880d681SAndroid Build Coastguard Worker; ld.global.nc.u32 34*9880d681SAndroid Build Coastguard Worker %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1i32(i32 addrspace(1)* %ptr, i32 4) 35*9880d681SAndroid Build Coastguard Worker ret i32 %val 36*9880d681SAndroid Build Coastguard Worker} 37