1*9880d681SAndroid Build Coastguard Worker; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 2*9880d681SAndroid Build Coastguard Worker; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64 3*9880d681SAndroid Build Coastguard Worker 4*9880d681SAndroid Build Coastguard Worker 5*9880d681SAndroid Build Coastguard Worker;; i8 6*9880d681SAndroid Build Coastguard Workerdefine i8 @ld_global_i8(i8 addrspace(1)* %ptr) { 7*9880d681SAndroid Build Coastguard Worker; PTX32: ld.global.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] 8*9880d681SAndroid Build Coastguard Worker; PTX32: ret 9*9880d681SAndroid Build Coastguard Worker; PTX64: ld.global.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 10*9880d681SAndroid Build Coastguard Worker; PTX64: ret 11*9880d681SAndroid Build Coastguard Worker %a = load i8, i8 addrspace(1)* %ptr 12*9880d681SAndroid Build Coastguard Worker ret i8 %a 13*9880d681SAndroid Build Coastguard Worker} 14*9880d681SAndroid Build Coastguard Worker 15*9880d681SAndroid Build Coastguard Workerdefine i8 @ld_shared_i8(i8 addrspace(3)* %ptr) { 16*9880d681SAndroid Build Coastguard Worker; PTX32: ld.shared.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] 17*9880d681SAndroid Build Coastguard Worker; PTX32: ret 18*9880d681SAndroid Build Coastguard Worker; PTX64: ld.shared.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 19*9880d681SAndroid Build Coastguard Worker; PTX64: ret 20*9880d681SAndroid Build Coastguard Worker %a = load i8, i8 addrspace(3)* %ptr 21*9880d681SAndroid Build Coastguard Worker ret i8 %a 22*9880d681SAndroid Build Coastguard Worker} 23*9880d681SAndroid Build Coastguard Worker 24*9880d681SAndroid Build Coastguard Workerdefine i8 @ld_local_i8(i8 addrspace(5)* %ptr) { 25*9880d681SAndroid Build Coastguard Worker; PTX32: ld.local.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] 26*9880d681SAndroid Build Coastguard Worker; PTX32: ret 27*9880d681SAndroid Build Coastguard Worker; PTX64: ld.local.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 28*9880d681SAndroid Build Coastguard Worker; PTX64: ret 29*9880d681SAndroid Build Coastguard Worker %a = load i8, i8 addrspace(5)* %ptr 30*9880d681SAndroid Build Coastguard Worker ret i8 %a 31*9880d681SAndroid Build Coastguard Worker} 32*9880d681SAndroid Build Coastguard Worker 33*9880d681SAndroid Build Coastguard Worker;; i16 34*9880d681SAndroid Build Coastguard Workerdefine i16 @ld_global_i16(i16 addrspace(1)* %ptr) { 35*9880d681SAndroid Build Coastguard Worker; PTX32: ld.global.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}] 36*9880d681SAndroid Build Coastguard Worker; PTX32: ret 37*9880d681SAndroid Build Coastguard Worker; PTX64: ld.global.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 38*9880d681SAndroid Build Coastguard Worker; PTX64: ret 39*9880d681SAndroid Build Coastguard Worker %a = load i16, i16 addrspace(1)* %ptr 40*9880d681SAndroid Build Coastguard Worker ret i16 %a 41*9880d681SAndroid Build Coastguard Worker} 42*9880d681SAndroid Build Coastguard Worker 43*9880d681SAndroid Build Coastguard Workerdefine i16 @ld_shared_i16(i16 addrspace(3)* %ptr) { 44*9880d681SAndroid Build Coastguard Worker; PTX32: ld.shared.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}] 45*9880d681SAndroid Build Coastguard Worker; PTX32: ret 46*9880d681SAndroid Build Coastguard Worker; PTX64: ld.shared.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 47*9880d681SAndroid Build Coastguard Worker; PTX64: ret 48*9880d681SAndroid Build Coastguard Worker %a = load i16, i16 addrspace(3)* %ptr 49*9880d681SAndroid Build Coastguard Worker ret i16 %a 50*9880d681SAndroid Build Coastguard Worker} 51*9880d681SAndroid Build Coastguard Worker 52*9880d681SAndroid Build Coastguard Workerdefine i16 @ld_local_i16(i16 addrspace(5)* %ptr) { 53*9880d681SAndroid Build Coastguard Worker; PTX32: ld.local.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}] 54*9880d681SAndroid Build Coastguard Worker; PTX32: ret 55*9880d681SAndroid Build Coastguard Worker; PTX64: ld.local.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 56*9880d681SAndroid Build Coastguard Worker; PTX64: ret 57*9880d681SAndroid Build Coastguard Worker %a = load i16, i16 addrspace(5)* %ptr 58*9880d681SAndroid Build Coastguard Worker ret i16 %a 59*9880d681SAndroid Build Coastguard Worker} 60*9880d681SAndroid Build Coastguard Worker 61*9880d681SAndroid Build Coastguard Worker;; i32 62*9880d681SAndroid Build Coastguard Workerdefine i32 @ld_global_i32(i32 addrspace(1)* %ptr) { 63*9880d681SAndroid Build Coastguard Worker; PTX32: ld.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}] 64*9880d681SAndroid Build Coastguard Worker; PTX32: ret 65*9880d681SAndroid Build Coastguard Worker; PTX64: ld.global.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 66*9880d681SAndroid Build Coastguard Worker; PTX64: ret 67*9880d681SAndroid Build Coastguard Worker %a = load i32, i32 addrspace(1)* %ptr 68*9880d681SAndroid Build Coastguard Worker ret i32 %a 69*9880d681SAndroid Build Coastguard Worker} 70*9880d681SAndroid Build Coastguard Worker 71*9880d681SAndroid Build Coastguard Workerdefine i32 @ld_shared_i32(i32 addrspace(3)* %ptr) { 72*9880d681SAndroid Build Coastguard Worker; PTX32: ld.shared.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}] 73*9880d681SAndroid Build Coastguard Worker; PTX32: ret 74*9880d681SAndroid Build Coastguard Worker; PTX64: ld.shared.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 75*9880d681SAndroid Build Coastguard Worker; PTX64: ret 76*9880d681SAndroid Build Coastguard Worker %a = load i32, i32 addrspace(3)* %ptr 77*9880d681SAndroid Build Coastguard Worker ret i32 %a 78*9880d681SAndroid Build Coastguard Worker} 79*9880d681SAndroid Build Coastguard Worker 80*9880d681SAndroid Build Coastguard Workerdefine i32 @ld_local_i32(i32 addrspace(5)* %ptr) { 81*9880d681SAndroid Build Coastguard Worker; PTX32: ld.local.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}] 82*9880d681SAndroid Build Coastguard Worker; PTX32: ret 83*9880d681SAndroid Build Coastguard Worker; PTX64: ld.local.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}] 84*9880d681SAndroid Build Coastguard Worker; PTX64: ret 85*9880d681SAndroid Build Coastguard Worker %a = load i32, i32 addrspace(5)* %ptr 86*9880d681SAndroid Build Coastguard Worker ret i32 %a 87*9880d681SAndroid Build Coastguard Worker} 88*9880d681SAndroid Build Coastguard Worker 89*9880d681SAndroid Build Coastguard Worker;; i64 90*9880d681SAndroid Build Coastguard Workerdefine i64 @ld_global_i64(i64 addrspace(1)* %ptr) { 91*9880d681SAndroid Build Coastguard Worker; PTX32: ld.global.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}] 92*9880d681SAndroid Build Coastguard Worker; PTX32: ret 93*9880d681SAndroid Build Coastguard Worker; PTX64: ld.global.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}] 94*9880d681SAndroid Build Coastguard Worker; PTX64: ret 95*9880d681SAndroid Build Coastguard Worker %a = load i64, i64 addrspace(1)* %ptr 96*9880d681SAndroid Build Coastguard Worker ret i64 %a 97*9880d681SAndroid Build Coastguard Worker} 98*9880d681SAndroid Build Coastguard Worker 99*9880d681SAndroid Build Coastguard Workerdefine i64 @ld_shared_i64(i64 addrspace(3)* %ptr) { 100*9880d681SAndroid Build Coastguard Worker; PTX32: ld.shared.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}] 101*9880d681SAndroid Build Coastguard Worker; PTX32: ret 102*9880d681SAndroid Build Coastguard Worker; PTX64: ld.shared.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}] 103*9880d681SAndroid Build Coastguard Worker; PTX64: ret 104*9880d681SAndroid Build Coastguard Worker %a = load i64, i64 addrspace(3)* %ptr 105*9880d681SAndroid Build Coastguard Worker ret i64 %a 106*9880d681SAndroid Build Coastguard Worker} 107*9880d681SAndroid Build Coastguard Worker 108*9880d681SAndroid Build Coastguard Workerdefine i64 @ld_local_i64(i64 addrspace(5)* %ptr) { 109*9880d681SAndroid Build Coastguard Worker; PTX32: ld.local.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}] 110*9880d681SAndroid Build Coastguard Worker; PTX32: ret 111*9880d681SAndroid Build Coastguard Worker; PTX64: ld.local.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}] 112*9880d681SAndroid Build Coastguard Worker; PTX64: ret 113*9880d681SAndroid Build Coastguard Worker %a = load i64, i64 addrspace(5)* %ptr 114*9880d681SAndroid Build Coastguard Worker ret i64 %a 115*9880d681SAndroid Build Coastguard Worker} 116*9880d681SAndroid Build Coastguard Worker 117*9880d681SAndroid Build Coastguard Worker;; f32 118*9880d681SAndroid Build Coastguard Workerdefine float @ld_global_f32(float addrspace(1)* %ptr) { 119*9880d681SAndroid Build Coastguard Worker; PTX32: ld.global.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}] 120*9880d681SAndroid Build Coastguard Worker; PTX32: ret 121*9880d681SAndroid Build Coastguard Worker; PTX64: ld.global.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}] 122*9880d681SAndroid Build Coastguard Worker; PTX64: ret 123*9880d681SAndroid Build Coastguard Worker %a = load float, float addrspace(1)* %ptr 124*9880d681SAndroid Build Coastguard Worker ret float %a 125*9880d681SAndroid Build Coastguard Worker} 126*9880d681SAndroid Build Coastguard Worker 127*9880d681SAndroid Build Coastguard Workerdefine float @ld_shared_f32(float addrspace(3)* %ptr) { 128*9880d681SAndroid Build Coastguard Worker; PTX32: ld.shared.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}] 129*9880d681SAndroid Build Coastguard Worker; PTX32: ret 130*9880d681SAndroid Build Coastguard Worker; PTX64: ld.shared.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}] 131*9880d681SAndroid Build Coastguard Worker; PTX64: ret 132*9880d681SAndroid Build Coastguard Worker %a = load float, float addrspace(3)* %ptr 133*9880d681SAndroid Build Coastguard Worker ret float %a 134*9880d681SAndroid Build Coastguard Worker} 135*9880d681SAndroid Build Coastguard Worker 136*9880d681SAndroid Build Coastguard Workerdefine float @ld_local_f32(float addrspace(5)* %ptr) { 137*9880d681SAndroid Build Coastguard Worker; PTX32: ld.local.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}] 138*9880d681SAndroid Build Coastguard Worker; PTX32: ret 139*9880d681SAndroid Build Coastguard Worker; PTX64: ld.local.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}] 140*9880d681SAndroid Build Coastguard Worker; PTX64: ret 141*9880d681SAndroid Build Coastguard Worker %a = load float, float addrspace(5)* %ptr 142*9880d681SAndroid Build Coastguard Worker ret float %a 143*9880d681SAndroid Build Coastguard Worker} 144*9880d681SAndroid Build Coastguard Worker 145*9880d681SAndroid Build Coastguard Worker;; f64 146*9880d681SAndroid Build Coastguard Workerdefine double @ld_global_f64(double addrspace(1)* %ptr) { 147*9880d681SAndroid Build Coastguard Worker; PTX32: ld.global.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}] 148*9880d681SAndroid Build Coastguard Worker; PTX32: ret 149*9880d681SAndroid Build Coastguard Worker; PTX64: ld.global.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}] 150*9880d681SAndroid Build Coastguard Worker; PTX64: ret 151*9880d681SAndroid Build Coastguard Worker %a = load double, double addrspace(1)* %ptr 152*9880d681SAndroid Build Coastguard Worker ret double %a 153*9880d681SAndroid Build Coastguard Worker} 154*9880d681SAndroid Build Coastguard Worker 155*9880d681SAndroid Build Coastguard Workerdefine double @ld_shared_f64(double addrspace(3)* %ptr) { 156*9880d681SAndroid Build Coastguard Worker; PTX32: ld.shared.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}] 157*9880d681SAndroid Build Coastguard Worker; PTX32: ret 158*9880d681SAndroid Build Coastguard Worker; PTX64: ld.shared.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}] 159*9880d681SAndroid Build Coastguard Worker; PTX64: ret 160*9880d681SAndroid Build Coastguard Worker %a = load double, double addrspace(3)* %ptr 161*9880d681SAndroid Build Coastguard Worker ret double %a 162*9880d681SAndroid Build Coastguard Worker} 163*9880d681SAndroid Build Coastguard Worker 164*9880d681SAndroid Build Coastguard Workerdefine double @ld_local_f64(double addrspace(5)* %ptr) { 165*9880d681SAndroid Build Coastguard Worker; PTX32: ld.local.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}] 166*9880d681SAndroid Build Coastguard Worker; PTX32: ret 167*9880d681SAndroid Build Coastguard Worker; PTX64: ld.local.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}] 168*9880d681SAndroid Build Coastguard Worker; PTX64: ret 169*9880d681SAndroid Build Coastguard Worker %a = load double, double addrspace(5)* %ptr 170*9880d681SAndroid Build Coastguard Worker ret double %a 171*9880d681SAndroid Build Coastguard Worker} 172