test/CodeGen/PTX/ld.ll

   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 i16 @t1_u16(i16* %p) {
  65 entry:
  66 ;CHECK: ld.global.u16 rh0, [r1];
  67 ;CHECK-NEXT: ret;
  68   %x = load i16* %p
  69   ret i16 %x
  70 }
  71
  72 define ptx_device i32 @t1_u32(i32* %p) {
  73 entry:
  74 ;CHECK: ld.global.u32 r0, [r1];
  75 ;CHECK-NEXT: ret;
  76   %x = load i32* %p
  77   ret i32 %x
  78 }
  79
  80 define ptx_device i64 @t1_u64(i64* %p) {
  81 entry:
  82 ;CHECK: ld.global.u64 rd0, [r1];
  83 ;CHECK-NEXT: ret;
  84   %x = load i64* %p
  85   ret i64 %x
  86 }
  87
  88 define ptx_device float @t1_f32(float* %p) {
  89 entry:
  90 ;CHECK: ld.global.f32 f0, [r1];
  91 ;CHECK-NEXT: ret;
  92   %x = load float* %p
  93   ret float %x
  94 }
  95
  96 define ptx_device double @t1_f64(double* %p) {
  97 entry:
  98 ;CHECK: ld.global.f64 fd0, [r1];
  99 ;CHECK-NEXT: ret;
 100   %x = load double* %p
 101   ret double %x
 102 }
 103
 104 define ptx_device i16 @t2_u16(i16* %p) {
 105 entry:
 106 ;CHECK: ld.global.u16 rh0, [r1+2];
 107 ;CHECK-NEXT: ret;
 108   %i = getelementptr i16* %p, i32 1
 109   %x = load i16* %i
 110   ret i16 %x
 111 }
 112
 113 define ptx_device i32 @t2_u32(i32* %p) {
 114 entry:
 115 ;CHECK: ld.global.u32 r0, [r1+4];
 116 ;CHECK-NEXT: ret;
 117   %i = getelementptr i32* %p, i32 1
 118   %x = load i32* %i
 119   ret i32 %x
 120 }
 121
 122 define ptx_device i64 @t2_u64(i64* %p) {
 123 entry:
 124 ;CHECK: ld.global.u64 rd0, [r1+8];
 125 ;CHECK-NEXT: ret;
 126   %i = getelementptr i64* %p, i32 1
 127   %x = load i64* %i
 128   ret i64 %x
 129 }
 130
 131 define ptx_device float @t2_f32(float* %p) {
 132 entry:
 133 ;CHECK: ld.global.f32 f0, [r1+4];
 134 ;CHECK-NEXT: ret;
 135   %i = getelementptr float* %p, i32 1
 136   %x = load float* %i
 137   ret float %x
 138 }
 139
 140 define ptx_device double @t2_f64(double* %p) {
 141 entry:
 142 ;CHECK: ld.global.f64 fd0, [r1+8];
 143 ;CHECK-NEXT: ret;
 144   %i = getelementptr double* %p, i32 1
 145   %x = load double* %i
 146   ret double %x
 147 }
 148
 149 define ptx_device i16 @t3_u16(i16* %p, i32 %q) {
 150 entry:
 151 ;CHECK: shl.b32 r0, r2, 1;
 152 ;CHECK-NEXT: add.u32 r0, r1, r0;
 153 ;CHECK-NEXT: ld.global.u16 rh0, [r0];
 154   %i = getelementptr i16* %p, i32 %q
 155   %x = load i16* %i
 156   ret i16 %x
 157 }
 158
 159 define ptx_device i32 @t3_u32(i32* %p, i32 %q) {
 160 entry:
 161 ;CHECK: shl.b32 r0, r2, 2;
 162 ;CHECK-NEXT: add.u32 r0, r1, r0;
 163 ;CHECK-NEXT: ld.global.u32 r0, [r0];
 164   %i = getelementptr i32* %p, i32 %q
 165   %x = load i32* %i
 166   ret i32 %x
 167 }
 168
 169 define ptx_device i64 @t3_u64(i64* %p, i32 %q) {
 170 entry:
 171 ;CHECK: shl.b32 r0, r2, 3;
 172 ;CHECK-NEXT: add.u32 r0, r1, r0;
 173 ;CHECK-NEXT: ld.global.u64 rd0, [r0];
 174   %i = getelementptr i64* %p, i32 %q
 175   %x = load i64* %i
 176   ret i64 %x
 177 }
 178
 179 define ptx_device float @t3_f32(float* %p, i32 %q) {
 180 entry:
 181 ;CHECK: shl.b32 r0, r2, 2;
 182 ;CHECK-NEXT: add.u32 r0, r1, r0;
 183 ;CHECK-NEXT: ld.global.f32 f0, [r0];
 184   %i = getelementptr float* %p, i32 %q
 185   %x = load float* %i
 186   ret float %x
 187 }
 188
 189 define ptx_device double @t3_f64(double* %p, i32 %q) {
 190 entry:
 191 ;CHECK: shl.b32 r0, r2, 3;
 192 ;CHECK-NEXT: add.u32 r0, r1, r0;
 193 ;CHECK-NEXT: ld.global.f64 fd0, [r0];
 194   %i = getelementptr double* %p, i32 %q
 195   %x = load double* %i
 196   ret double %x
 197 }
 198
 199 define ptx_device i16 @t4_global_u16() {
 200 entry:
 201 ;CHECK: mov.u32 r0, array_i16;
 202 ;CHECK-NEXT: ld.global.u16 rh0, [r0];
 203 ;CHECK-NEXT: ret;
 204   %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 0
 205   %x = load i16* %i
 206   ret i16 %x
 207 }
 208
 209 define ptx_device i32 @t4_global_u32() {
 210 entry:
 211 ;CHECK: mov.u32 r0, array_i32;
 212 ;CHECK-NEXT: ld.global.u32 r0, [r0];
 213 ;CHECK-NEXT: ret;
 214   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 0
 215   %x = load i32* %i
 216   ret i32 %x
 217 }
 218
 219 define ptx_device i64 @t4_global_u64() {
 220 entry:
 221 ;CHECK: mov.u32 r0, array_i64;
 222 ;CHECK-NEXT: ld.global.u64 rd0, [r0];
 223 ;CHECK-NEXT: ret;
 224   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 0
 225   %x = load i64* %i
 226   ret i64 %x
 227 }
 228
 229 define ptx_device float @t4_global_f32() {
 230 entry:
 231 ;CHECK: mov.u32 r0, array_float;
 232 ;CHECK-NEXT: ld.global.f32 f0, [r0];
 233 ;CHECK-NEXT: ret;
 234   %i = getelementptr [10 x float]* @array_float, i32 0, i32 0
 235   %x = load float* %i
 236   ret float %x
 237 }
 238
 239 define ptx_device double @t4_global_f64() {
 240 entry:
 241 ;CHECK: mov.u32 r0, array_double;
 242 ;CHECK-NEXT: ld.global.f64 fd0, [r0];
 243 ;CHECK-NEXT: ret;
 244   %i = getelementptr [10 x double]* @array_double, i32 0, i32 0
 245   %x = load double* %i
 246   ret double %x
 247 }
 248
 249 define ptx_device i16 @t4_const_u16() {
 250 entry:
 251 ;CHECK: mov.u32 r0, array_constant_i16;
 252 ;CHECK-NEXT: ld.const.u16 rh0, [r0];
 253 ;CHECK-NEXT: ret;
 254   %i = getelementptr [10 x i16] addrspace(1)* @array_constant_i16, i32 0, i32 0
 255   %x = load i16 addrspace(1)* %i
 256   ret i16 %x
 257 }
 258
 259 define ptx_device i32 @t4_const_u32() {
 260 entry:
 261 ;CHECK: mov.u32 r0, array_constant_i32;
 262 ;CHECK-NEXT: ld.const.u32 r0, [r0];
 263 ;CHECK-NEXT: ret;
 264   %i = getelementptr [10 x i32] addrspace(1)* @array_constant_i32, i32 0, i32 0
 265   %x = load i32 addrspace(1)* %i
 266   ret i32 %x
 267 }
 268
 269 define ptx_device i64 @t4_const_u64() {
 270 entry:
 271 ;CHECK: mov.u32 r0, array_constant_i64;
 272 ;CHECK-NEXT: ld.const.u64 rd0, [r0];
 273 ;CHECK-NEXT: ret;
 274   %i = getelementptr [10 x i64] addrspace(1)* @array_constant_i64, i32 0, i32 0
 275   %x = load i64 addrspace(1)* %i
 276   ret i64 %x
 277 }
 278
 279 define ptx_device float @t4_const_f32() {
 280 entry:
 281 ;CHECK: mov.u32 r0, array_constant_float;
 282 ;CHECK-NEXT: ld.const.f32 f0, [r0];
 283 ;CHECK-NEXT: ret;
 284   %i = getelementptr [10 x float] addrspace(1)* @array_constant_float, i32 0, i32 0
 285   %x = load float addrspace(1)* %i
 286   ret float %x
 287 }
 288
 289 define ptx_device double @t4_const_f64() {
 290 entry:
 291 ;CHECK: mov.u32 r0, array_constant_double;
 292 ;CHECK-NEXT: ld.const.f64 fd0, [r0];
 293 ;CHECK-NEXT: ret;
 294   %i = getelementptr [10 x double] addrspace(1)* @array_constant_double, i32 0, i32 0
 295   %x = load double addrspace(1)* %i
 296   ret double %x
 297 }
 298
 299 define ptx_device i16 @t4_local_u16() {
 300 entry:
 301 ;CHECK: mov.u32 r0, array_local_i16;
 302 ;CHECK-NEXT: ld.local.u16 rh0, [r0];
 303 ;CHECK-NEXT: ret;
 304   %i = getelementptr [10 x i16] addrspace(2)* @array_local_i16, i32 0, i32 0
 305   %x = load i16 addrspace(2)* %i
 306   ret i16 %x
 307 }
 308
 309 define ptx_device i32 @t4_local_u32() {
 310 entry:
 311 ;CHECK: mov.u32 r0, array_local_i32;
 312 ;CHECK-NEXT: ld.local.u32 r0, [r0];
 313 ;CHECK-NEXT: ret;
 314   %i = getelementptr [10 x i32] addrspace(2)* @array_local_i32, i32 0, i32 0
 315   %x = load i32 addrspace(2)* %i
 316   ret i32 %x
 317 }
 318
 319 define ptx_device i64 @t4_local_u64() {
 320 entry:
 321 ;CHECK: mov.u32 r0, array_local_i64;
 322 ;CHECK-NEXT: ld.local.u64 rd0, [r0];
 323 ;CHECK-NEXT: ret;
 324   %i = getelementptr [10 x i64] addrspace(2)* @array_local_i64, i32 0, i32 0
 325   %x = load i64 addrspace(2)* %i
 326   ret i64 %x
 327 }
 328
 329 define ptx_device float @t4_local_f32() {
 330 entry:
 331 ;CHECK: mov.u32 r0, array_local_float;
 332 ;CHECK-NEXT: ld.local.f32 f0, [r0];
 333 ;CHECK-NEXT: ret;
 334   %i = getelementptr [10 x float] addrspace(2)* @array_local_float, i32 0, i32 0
 335   %x = load float addrspace(2)* %i
 336   ret float %x
 337 }
 338
 339 define ptx_device double @t4_local_f64() {
 340 entry:
 341 ;CHECK: mov.u32 r0, array_local_double;
 342 ;CHECK-NEXT: ld.local.f64 fd0, [r0];
 343 ;CHECK-NEXT: ret;
 344   %i = getelementptr [10 x double] addrspace(2)* @array_local_double, i32 0, i32 0
 345   %x = load double addrspace(2)* %i
 346   ret double %x
 347 }
 348
 349 define ptx_device i16 @t4_shared_u16() {
 350 entry:
 351 ;CHECK: mov.u32 r0, array_shared_i16;
 352 ;CHECK-NEXT: ld.shared.u16 rh0, [r0];
 353 ;CHECK-NEXT: ret;
 354   %i = getelementptr [10 x i16] addrspace(4)* @array_shared_i16, i32 0, i32 0
 355   %x = load i16 addrspace(4)* %i
 356   ret i16 %x
 357 }
 358
 359 define ptx_device i32 @t4_shared_u32() {
 360 entry:
 361 ;CHECK: mov.u32 r0, array_shared_i32;
 362 ;CHECK-NEXT: ld.shared.u32 r0, [r0];
 363 ;CHECK-NEXT: ret;
 364   %i = getelementptr [10 x i32] addrspace(4)* @array_shared_i32, i32 0, i32 0
 365   %x = load i32 addrspace(4)* %i
 366   ret i32 %x
 367 }
 368
 369 define ptx_device i64 @t4_shared_u64() {
 370 entry:
 371 ;CHECK: mov.u32 r0, array_shared_i64;
 372 ;CHECK-NEXT: ld.shared.u64 rd0, [r0];
 373 ;CHECK-NEXT: ret;
 374   %i = getelementptr [10 x i64] addrspace(4)* @array_shared_i64, i32 0, i32 0
 375   %x = load i64 addrspace(4)* %i
 376   ret i64 %x
 377 }
 378
 379 define ptx_device float @t4_shared_f32() {
 380 entry:
 381 ;CHECK: mov.u32 r0, array_shared_float;
 382 ;CHECK-NEXT: ld.shared.f32 f0, [r0];
 383 ;CHECK-NEXT: ret;
 384   %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0
 385   %x = load float addrspace(4)* %i
 386   ret float %x
 387 }
 388
 389 define ptx_device double @t4_shared_f64() {
 390 entry:
 391 ;CHECK: mov.u32 r0, array_shared_double;
 392 ;CHECK-NEXT: ld.shared.f64 fd0, [r0];
 393 ;CHECK-NEXT: ret;
 394   %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0
 395   %x = load double addrspace(4)* %i
 396   ret double %x
 397 }
 398
 399 define ptx_device i16 @t5_u16() {
 400 entry:
 401 ;CHECK: mov.u32 r0, array_i16;
 402 ;CHECK-NEXT: ld.global.u16 rh0, [r0+2];
 403 ;CHECK-NEXT: ret;
 404   %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 1
 405   %x = load i16* %i
 406   ret i16 %x
 407 }
 408
 409 define ptx_device i32 @t5_u32() {
 410 entry:
 411 ;CHECK: mov.u32 r0, array_i32;
 412 ;CHECK-NEXT: ld.global.u32 r0, [r0+4];
 413 ;CHECK-NEXT: ret;
 414   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 1
 415   %x = load i32* %i
 416   ret i32 %x
 417 }
 418
 419 define ptx_device i64 @t5_u64() {
 420 entry:
 421 ;CHECK: mov.u32 r0, array_i64;
 422 ;CHECK-NEXT: ld.global.u64 rd0, [r0+8];
 423 ;CHECK-NEXT: ret;
 424   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 1
 425   %x = load i64* %i
 426   ret i64 %x
 427 }
 428
 429 define ptx_device float @t5_f32() {
 430 entry:
 431 ;CHECK: mov.u32 r0, array_float;
 432 ;CHECK-NEXT: ld.global.f32 f0, [r0+4];
 433 ;CHECK-NEXT: ret;
 434   %i = getelementptr [10 x float]* @array_float, i32 0, i32 1
 435   %x = load float* %i
 436   ret float %x
 437 }
 438
 439 define ptx_device double @t5_f64() {
 440 entry:
 441 ;CHECK: mov.u32 r0, array_double;
 442 ;CHECK-NEXT: ld.global.f64 fd0, [r0+8];
 443 ;CHECK-NEXT: ret;
 444   %i = getelementptr [10 x double]* @array_double, i32 0, i32 1
 445   %x = load double* %i
 446   ret double %x
 447 }