1 ; RUN: llc < %s -march=ptx32 | FileCheck %s 2 3 ;CHECK: .extern .global .b8 array_i16[20]; 4 @array_i16 = external global [10 x i16] 5 6 ;CHECK: .extern .const .b8 array_constant_i16[20]; 7 @array_constant_i16 = external addrspace(1) constant [10 x i16] 8 9 ;CHECK: .extern .local .b8 array_local_i16[20]; 10 @array_local_i16 = external addrspace(2) global [10 x i16] 11 12 ;CHECK: .extern .shared .b8 array_shared_i16[20]; 13 @array_shared_i16 = external addrspace(4) global [10 x i16] 14 15 ;CHECK: .extern .global .b8 array_i32[40]; 16 @array_i32 = external global [10 x i32] 17 18 ;CHECK: .extern .const .b8 array_constant_i32[40]; 19 @array_constant_i32 = external addrspace(1) constant [10 x i32] 20 21 ;CHECK: .extern .local .b8 array_local_i32[40]; 22 @array_local_i32 = external addrspace(2) global [10 x i32] 23 24 ;CHECK: .extern .shared .b8 array_shared_i32[40]; 25 @array_shared_i32 = external addrspace(4) global [10 x i32] 26 27 ;CHECK: .extern .global .b8 array_i64[80]; 28 @array_i64 = external global [10 x i64] 29 30 ;CHECK: .extern .const .b8 array_constant_i64[80]; 31 @array_constant_i64 = external addrspace(1) constant [10 x i64] 32 33 ;CHECK: .extern .local .b8 array_local_i64[80]; 34 @array_local_i64 = external addrspace(2) global [10 x i64] 35 36 ;CHECK: .extern .shared .b8 array_shared_i64[80]; 37 @array_shared_i64 = external addrspace(4) global [10 x i64] 38 39 ;CHECK: .extern .global .b8 array_float[40]; 40 @array_float = external global [10 x float] 41 42 ;CHECK: .extern .const .b8 array_constant_float[40]; 43 @array_constant_float = external addrspace(1) constant [10 x float] 44 45 ;CHECK: .extern .local .b8 array_local_float[40]; 46 @array_local_float = external addrspace(2) global [10 x float] 47 48 ;CHECK: .extern .shared .b8 array_shared_float[40]; 49 @array_shared_float = external addrspace(4) global [10 x float] 50 51 ;CHECK: .extern .global .b8 array_double[80]; 52 @array_double = external global [10 x double] 53 54 ;CHECK: .extern .const .b8 array_constant_double[80]; 55 @array_constant_double = external addrspace(1) constant [10 x double] 56 57 ;CHECK: .extern .local .b8 array_local_double[80]; 58 @array_local_double = external addrspace(2) global [10 x double] 59 60 ;CHECK: .extern .shared .b8 array_shared_double[80]; 61 @array_shared_double = external addrspace(4) global [10 x double] 62 63 64 define ptx_device void @t1_u16(i16* %p, i16 %x) { 65 entry: 66 ;CHECK: st.global.u16 [r{{[0-9]+}}], rh{{[0-9]+}}; 67 ;CHECK-NEXT: ret; 68 store i16 %x, i16* %p 69 ret void 70 } 71 72 define ptx_device void @t1_u32(i32* %p, i32 %x) { 73 entry: 74 ;CHECK: st.global.u32 [r{{[0-9]+}}], r{{[0-9]+}}; 75 ;CHECK-NEXT: ret; 76 store i32 %x, i32* %p 77 ret void 78 } 79 80 define ptx_device void @t1_u64(i64* %p, i64 %x) { 81 entry: 82 ;CHECK: st.global.u64 [r{{[0-9]+}}], rd{{[0-9]+}}; 83 ;CHECK-NEXT: ret; 84 store i64 %x, i64* %p 85 ret void 86 } 87 88 define ptx_device void @t1_f32(float* %p, float %x) { 89 entry: 90 ;CHECK: st.global.f32 [r{{[0-9]+}}], r{{[0-9]+}}; 91 ;CHECK-NEXT: ret; 92 store float %x, float* %p 93 ret void 94 } 95 96 define ptx_device void @t1_f64(double* %p, double %x) { 97 entry: 98 ;CHECK: st.global.f64 [r{{[0-9]+}}], rd{{[0-9]+}}; 99 ;CHECK-NEXT: ret; 100 store double %x, double* %p 101 ret void 102 } 103 104 define ptx_device void @t2_u16(i16* %p, i16 %x) { 105 entry: 106 ;CHECK: st.global.u16 [r{{[0-9]+}}+2], rh{{[0-9]+}}; 107 ;CHECK-NEXT: ret; 108 %i = getelementptr i16* %p, i32 1 109 store i16 %x, i16* %i 110 ret void 111 } 112 113 define ptx_device void @t2_u32(i32* %p, i32 %x) { 114 entry: 115 ;CHECK: st.global.u32 [r{{[0-9]+}}+4], r{{[0-9]+}}; 116 ;CHECK-NEXT: ret; 117 %i = getelementptr i32* %p, i32 1 118 store i32 %x, i32* %i 119 ret void 120 } 121 122 define ptx_device void @t2_u64(i64* %p, i64 %x) { 123 entry: 124 ;CHECK: st.global.u64 [r{{[0-9]+}}+8], rd{{[0-9]+}}; 125 ;CHECK-NEXT: ret; 126 %i = getelementptr i64* %p, i32 1 127 store i64 %x, i64* %i 128 ret void 129 } 130 131 define ptx_device void @t2_f32(float* %p, float %x) { 132 entry: 133 ;CHECK: st.global.f32 [r{{[0-9]+}}+4], r{{[0-9]+}}; 134 ;CHECK-NEXT: ret; 135 %i = getelementptr float* %p, i32 1 136 store float %x, float* %i 137 ret void 138 } 139 140 define ptx_device void @t2_f64(double* %p, double %x) { 141 entry: 142 ;CHECK: st.global.f64 [r{{[0-9]+}}+8], rd{{[0-9]+}}; 143 ;CHECK-NEXT: ret; 144 %i = getelementptr double* %p, i32 1 145 store double %x, double* %i 146 ret void 147 } 148 149 define ptx_device void @t3_u16(i16* %p, i32 %q, i16 %x) { 150 entry: 151 ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 1; 152 ;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]]; 153 ;CHECK-NEXT: st.global.u16 [r[[R0]]], rh{{[0-9]+}}; 154 ;CHECK-NEXT: ret; 155 %i = getelementptr i16* %p, i32 %q 156 store i16 %x, i16* %i 157 ret void 158 } 159 160 define ptx_device void @t3_u32(i32* %p, i32 %q, i32 %x) { 161 entry: 162 ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 2; 163 ;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]]; 164 ;CHECK-NEXT: st.global.u32 [r[[R0]]], r{{[0-9]+}}; 165 ;CHECK-NEXT: ret; 166 %i = getelementptr i32* %p, i32 %q 167 store i32 %x, i32* %i 168 ret void 169 } 170 171 define ptx_device void @t3_u64(i64* %p, i32 %q, i64 %x) { 172 entry: 173 ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 3; 174 ;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]]; 175 ;CHECK-NEXT: st.global.u64 [r[[R0]]], rd{{[0-9]+}}; 176 ;CHECK-NEXT: ret; 177 %i = getelementptr i64* %p, i32 %q 178 store i64 %x, i64* %i 179 ret void 180 } 181 182 define ptx_device void @t3_f32(float* %p, i32 %q, float %x) { 183 entry: 184 ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 2; 185 ;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]]; 186 ;CHECK-NEXT: st.global.f32 [r[[R0]]], r{{[0-9]+}}; 187 ;CHECK-NEXT: ret; 188 %i = getelementptr float* %p, i32 %q 189 store float %x, float* %i 190 ret void 191 } 192 193 define ptx_device void @t3_f64(double* %p, i32 %q, double %x) { 194 entry: 195 ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 3; 196 ;CHECK-NEXT: add.u32 r[[R0]], r{{[0-9]+}}, r[[R0]]; 197 ;CHECK-NEXT: st.global.f64 [r[[R0]]], rd{{[0-9]+}}; 198 ;CHECK-NEXT: ret; 199 %i = getelementptr double* %p, i32 %q 200 store double %x, double* %i 201 ret void 202 } 203 204 define ptx_device void @t4_global_u16(i16 %x) { 205 entry: 206 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16; 207 ;CHECK-NEXT: st.global.u16 [r[[R0]]], rh{{[0-9]+}}; 208 ;CHECK-NEXT: ret; 209 %i = getelementptr [10 x i16]* @array_i16, i16 0, i16 0 210 store i16 %x, i16* %i 211 ret void 212 } 213 214 define ptx_device void @t4_global_u32(i32 %x) { 215 entry: 216 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i32; 217 ;CHECK-NEXT: st.global.u32 [r[[R0]]], r{{[0-9]+}}; 218 ;CHECK-NEXT: ret; 219 %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 0 220 store i32 %x, i32* %i 221 ret void 222 } 223 224 define ptx_device void @t4_global_u64(i64 %x) { 225 entry: 226 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i64; 227 ;CHECK-NEXT: st.global.u64 [r[[R0]]], rd{{[0-9]+}}; 228 ;CHECK-NEXT: ret; 229 %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 0 230 store i64 %x, i64* %i 231 ret void 232 } 233 234 define ptx_device void @t4_global_f32(float %x) { 235 entry: 236 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_float; 237 ;CHECK-NEXT: st.global.f32 [r[[R0]]], r{{[0-9]+}}; 238 ;CHECK-NEXT: ret; 239 %i = getelementptr [10 x float]* @array_float, i32 0, i32 0 240 store float %x, float* %i 241 ret void 242 } 243 244 define ptx_device void @t4_global_f64(double %x) { 245 entry: 246 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_double; 247 ;CHECK-NEXT: st.global.f64 [r[[R0]]], rd{{[0-9]+}}; 248 ;CHECK-NEXT: ret; 249 %i = getelementptr [10 x double]* @array_double, i32 0, i32 0 250 store double %x, double* %i 251 ret void 252 } 253 254 define ptx_device void @t4_local_u16(i16 %x) { 255 entry: 256 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i16; 257 ;CHECK-NEXT: st.local.u16 [r[[R0]]], rh{{[0-9]+}}; 258 ;CHECK-NEXT: ret; 259 %i = getelementptr [10 x i16] addrspace(2)* @array_local_i16, i32 0, i32 0 260 store i16 %x, i16 addrspace(2)* %i 261 ret void 262 } 263 264 define ptx_device void @t4_local_u32(i32 %x) { 265 entry: 266 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i32; 267 ;CHECK-NEXT: st.local.u32 [r[[R0]]], r{{[0-9]+}}; 268 ;CHECK-NEXT: ret; 269 %i = getelementptr [10 x i32] addrspace(2)* @array_local_i32, i32 0, i32 0 270 store i32 %x, i32 addrspace(2)* %i 271 ret void 272 } 273 274 define ptx_device void @t4_local_u64(i64 %x) { 275 entry: 276 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i64; 277 ;CHECK-NEXT: st.local.u64 [r[[R0]]], rd{{[0-9]+}}; 278 ;CHECK-NEXT: ret; 279 %i = getelementptr [10 x i64] addrspace(2)* @array_local_i64, i32 0, i32 0 280 store i64 %x, i64 addrspace(2)* %i 281 ret void 282 } 283 284 define ptx_device void @t4_local_f32(float %x) { 285 entry: 286 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_float; 287 ;CHECK-NEXT: st.local.f32 [r[[R0]]], r{{[0-9]+}}; 288 ;CHECK-NEXT: ret; 289 %i = getelementptr [10 x float] addrspace(2)* @array_local_float, i32 0, i32 0 290 store float %x, float addrspace(2)* %i 291 ret void 292 } 293 294 define ptx_device void @t4_local_f64(double %x) { 295 entry: 296 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_double; 297 ;CHECK-NEXT: st.local.f64 [r[[R0]]], rd{{[0-9]+}}; 298 ;CHECK-NEXT: ret; 299 %i = getelementptr [10 x double] addrspace(2)* @array_local_double, i32 0, i32 0 300 store double %x, double addrspace(2)* %i 301 ret void 302 } 303 304 define ptx_device void @t4_shared_u16(i16 %x) { 305 entry: 306 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i16; 307 ;CHECK-NEXT: st.shared.u16 [r[[R0]]], rh{{[0-9]+}}; 308 ;CHECK-NEXT: ret; 309 %i = getelementptr [10 x i16] addrspace(4)* @array_shared_i16, i32 0, i32 0 310 store i16 %x, i16 addrspace(4)* %i 311 ret void 312 } 313 314 define ptx_device void @t4_shared_u32(i32 %x) { 315 entry: 316 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i32; 317 ;CHECK-NEXT: st.shared.u32 [r[[R0]]], r{{[0-9]+}}; 318 ;CHECK-NEXT: ret; 319 %i = getelementptr [10 x i32] addrspace(4)* @array_shared_i32, i32 0, i32 0 320 store i32 %x, i32 addrspace(4)* %i 321 ret void 322 } 323 324 define ptx_device void @t4_shared_u64(i64 %x) { 325 entry: 326 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i64; 327 ;CHECK-NEXT: st.shared.u64 [r[[R0]]], rd{{[0-9]+}}; 328 ;CHECK-NEXT: ret; 329 %i = getelementptr [10 x i64] addrspace(4)* @array_shared_i64, i32 0, i32 0 330 store i64 %x, i64 addrspace(4)* %i 331 ret void 332 } 333 334 define ptx_device void @t4_shared_f32(float %x) { 335 entry: 336 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_float; 337 ;CHECK-NEXT: st.shared.f32 [r[[R0]]], r{{[0-9]+}}; 338 ;CHECK-NEXT: ret; 339 %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0 340 store float %x, float addrspace(4)* %i 341 ret void 342 } 343 344 define ptx_device void @t4_shared_f64(double %x) { 345 entry: 346 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_double; 347 ;CHECK-NEXT: st.shared.f64 [r[[R0]]], rd{{[0-9]+}}; 348 ;CHECK-NEXT: ret; 349 %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0 350 store double %x, double addrspace(4)* %i 351 ret void 352 } 353 354 define ptx_device void @t5_u16(i16 %x) { 355 entry: 356 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16; 357 ;CHECK-NEXT: st.global.u16 [r[[R0]]+2], rh{{[0-9]+}}; 358 ;CHECK-NEXT: ret; 359 %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 1 360 store i16 %x, i16* %i 361 ret void 362 } 363 364 define ptx_device void @t5_u32(i32 %x) { 365 entry: 366 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i32; 367 ;CHECK-NEXT: st.global.u32 [r[[R0]]+4], r{{[0-9]+}}; 368 ;CHECK-NEXT: ret; 369 %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 1 370 store i32 %x, i32* %i 371 ret void 372 } 373 374 define ptx_device void @t5_u64(i64 %x) { 375 entry: 376 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i64; 377 ;CHECK-NEXT: st.global.u64 [r[[R0]]+8], rd{{[0-9]+}}; 378 ;CHECK-NEXT: ret; 379 %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 1 380 store i64 %x, i64* %i 381 ret void 382 } 383 384 define ptx_device void @t5_f32(float %x) { 385 entry: 386 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_float; 387 ;CHECK-NEXT: st.global.f32 [r[[R0]]+4], r{{[0-9]+}}; 388 ;CHECK-NEXT: ret; 389 %i = getelementptr [10 x float]* @array_float, i32 0, i32 1 390 store float %x, float* %i 391 ret void 392 } 393 394 define ptx_device void @t5_f64(double %x) { 395 entry: 396 ;CHECK: mov.u32 r[[R0:[0-9]+]], array_double; 397 ;CHECK-NEXT: st.global.f64 [r[[R0]]+8], rd{{[0-9]+}}; 398 ;CHECK-NEXT: ret; 399 %i = getelementptr [10 x double]* @array_double, i32 0, i32 1 400 store double %x, double* %i 401 ret void 402 } 403
