test/CodeGen/PTX/ld.ll

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