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