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 7 define 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 } 15 define 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 } 23 define 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 33 define 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 } 41 define 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 } 49 define 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 59 define 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 } 67 define 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 } 75 define 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 85 define 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 } 93 define 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 } 101 define 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 111 define 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 } 119 define 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 } 127 define 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 137 define 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 } 145 define 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 } 153 define 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
