1 ; RUN: llc < %s -march=nvptx -mcpu=sm_10 | FileCheck %s --check-prefix=PTX32 2 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 3 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_10 | FileCheck %s --check-prefix=PTX64 4 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64 5 6 7 ;; i8 8 9 define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) { 10 ; PTX32: st.global.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}} 11 ; PTX32: ret 12 ; PTX64: st.global.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}} 13 ; PTX64: ret 14 store i8 %a, i8 addrspace(1)* %ptr 15 ret void 16 } 17 18 define void @st_shared_i8(i8 addrspace(3)* %ptr, i8 %a) { 19 ; PTX32: st.shared.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}} 20 ; PTX32: ret 21 ; PTX64: st.shared.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}} 22 ; PTX64: ret 23 store i8 %a, i8 addrspace(3)* %ptr 24 ret void 25 } 26 27 define void @st_local_i8(i8 addrspace(5)* %ptr, i8 %a) { 28 ; PTX32: st.local.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}} 29 ; PTX32: ret 30 ; PTX64: st.local.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}} 31 ; PTX64: ret 32 store i8 %a, i8 addrspace(5)* %ptr 33 ret void 34 } 35 36 ;; i16 37 38 define void @st_global_i16(i16 addrspace(1)* %ptr, i16 %a) { 39 ; PTX32: st.global.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}} 40 ; PTX32: ret 41 ; PTX64: st.global.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}} 42 ; PTX64: ret 43 store i16 %a, i16 addrspace(1)* %ptr 44 ret void 45 } 46 47 define void @st_shared_i16(i16 addrspace(3)* %ptr, i16 %a) { 48 ; PTX32: st.shared.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}} 49 ; PTX32: ret 50 ; PTX64: st.shared.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}} 51 ; PTX64: ret 52 store i16 %a, i16 addrspace(3)* %ptr 53 ret void 54 } 55 56 define void @st_local_i16(i16 addrspace(5)* %ptr, i16 %a) { 57 ; PTX32: st.local.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}} 58 ; PTX32: ret 59 ; PTX64: st.local.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}} 60 ; PTX64: ret 61 store i16 %a, i16 addrspace(5)* %ptr 62 ret void 63 } 64 65 ;; i32 66 67 define void @st_global_i32(i32 addrspace(1)* %ptr, i32 %a) { 68 ; PTX32: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}} 69 ; PTX32: ret 70 ; PTX64: st.global.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}} 71 ; PTX64: ret 72 store i32 %a, i32 addrspace(1)* %ptr 73 ret void 74 } 75 76 define void @st_shared_i32(i32 addrspace(3)* %ptr, i32 %a) { 77 ; PTX32: st.shared.u32 [%r{{[0-9]+}}], %r{{[0-9]+}} 78 ; PTX32: ret 79 ; PTX64: st.shared.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}} 80 ; PTX64: ret 81 store i32 %a, i32 addrspace(3)* %ptr 82 ret void 83 } 84 85 define void @st_local_i32(i32 addrspace(5)* %ptr, i32 %a) { 86 ; PTX32: st.local.u32 [%r{{[0-9]+}}], %r{{[0-9]+}} 87 ; PTX32: ret 88 ; PTX64: st.local.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}} 89 ; PTX64: ret 90 store i32 %a, i32 addrspace(5)* %ptr 91 ret void 92 } 93 94 ;; i64 95 96 define void @st_global_i64(i64 addrspace(1)* %ptr, i64 %a) { 97 ; PTX32: st.global.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}} 98 ; PTX32: ret 99 ; PTX64: st.global.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}} 100 ; PTX64: ret 101 store i64 %a, i64 addrspace(1)* %ptr 102 ret void 103 } 104 105 define void @st_shared_i64(i64 addrspace(3)* %ptr, i64 %a) { 106 ; PTX32: st.shared.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}} 107 ; PTX32: ret 108 ; PTX64: st.shared.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}} 109 ; PTX64: ret 110 store i64 %a, i64 addrspace(3)* %ptr 111 ret void 112 } 113 114 define void @st_local_i64(i64 addrspace(5)* %ptr, i64 %a) { 115 ; PTX32: st.local.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}} 116 ; PTX32: ret 117 ; PTX64: st.local.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}} 118 ; PTX64: ret 119 store i64 %a, i64 addrspace(5)* %ptr 120 ret void 121 } 122 123 ;; f32 124 125 define void @st_global_f32(float addrspace(1)* %ptr, float %a) { 126 ; PTX32: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}} 127 ; PTX32: ret 128 ; PTX64: st.global.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}} 129 ; PTX64: ret 130 store float %a, float addrspace(1)* %ptr 131 ret void 132 } 133 134 define void @st_shared_f32(float addrspace(3)* %ptr, float %a) { 135 ; PTX32: st.shared.f32 [%r{{[0-9]+}}], %f{{[0-9]+}} 136 ; PTX32: ret 137 ; PTX64: st.shared.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}} 138 ; PTX64: ret 139 store float %a, float addrspace(3)* %ptr 140 ret void 141 } 142 143 define void @st_local_f32(float addrspace(5)* %ptr, float %a) { 144 ; PTX32: st.local.f32 [%r{{[0-9]+}}], %f{{[0-9]+}} 145 ; PTX32: ret 146 ; PTX64: st.local.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}} 147 ; PTX64: ret 148 store float %a, float addrspace(5)* %ptr 149 ret void 150 } 151 152 ;; f64 153 154 define void @st_global_f64(double addrspace(1)* %ptr, double %a) { 155 ; PTX32: st.global.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}} 156 ; PTX32: ret 157 ; PTX64: st.global.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}} 158 ; PTX64: ret 159 store double %a, double addrspace(1)* %ptr 160 ret void 161 } 162 163 define void @st_shared_f64(double addrspace(3)* %ptr, double %a) { 164 ; PTX32: st.shared.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}} 165 ; PTX32: ret 166 ; PTX64: st.shared.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}} 167 ; PTX64: ret 168 store double %a, double addrspace(3)* %ptr 169 ret void 170 } 171 172 define void @st_local_f64(double addrspace(5)* %ptr, double %a) { 173 ; PTX32: st.local.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}} 174 ; PTX32: ret 175 ; PTX64: st.local.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}} 176 ; PTX64: ret 177 store double %a, double addrspace(5)* %ptr 178 ret void 179 } 180
