test/CodeGen/PTX/st.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 void @t1_u16(i16* %p, i16 %x) {
  65 entry:
  66 ;CHECK: st.global.u16 [r1], rh1;
  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 [r1], r2;
  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 [r1], rd1;
  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 [r1], f1;
  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 [r1], fd1;
  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 [r1+2], rh1;
 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 [r1+4], r2;
 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 [r1+8], rd1;
 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 [r1+4], f1;
 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 [r1+8], fd1;
 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 r0, r2, 1;
 152 ;CHECK-NEXT: add.u32 r0, r1, r0;
 153 ;CHECK-NEXT: st.global.u16 [r0], rh1;
 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 r0, r2, 2;
 163 ;CHECK-NEXT: add.u32 r0, r1, r0;
 164 ;CHECK-NEXT: st.global.u32 [r0], r3;
 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 r0, r2, 3;
 174 ;CHECK-NEXT: add.u32 r0, r1, r0;
 175 ;CHECK-NEXT: st.global.u64 [r0], rd1;
 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 r0, r2, 2;
 185 ;CHECK-NEXT: add.u32 r0, r1, r0;
 186 ;CHECK-NEXT: st.global.f32 [r0], f1;
 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 r0, r2, 3;
 196 ;CHECK-NEXT: add.u32 r0, r1, r0;
 197 ;CHECK-NEXT: st.global.f64 [r0], fd1;
 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 r0, array_i16;
 207 ;CHECK-NEXT: st.global.u16 [r0], rh1;
 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 r0, array_i32;
 217 ;CHECK-NEXT: st.global.u32 [r0], r1;
 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 r0, array_i64;
 227 ;CHECK-NEXT: st.global.u64 [r0], rd1;
 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 r0, array_float;
 237 ;CHECK-NEXT: st.global.f32 [r0], f1;
 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 r0, array_double;
 247 ;CHECK-NEXT: st.global.f64 [r0], fd1;
 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 r0, array_local_i16;
 257 ;CHECK-NEXT: st.local.u16 [r0], rh1;
 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 r0, array_local_i32;
 267 ;CHECK-NEXT: st.local.u32 [r0], r1;
 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 r0, array_local_i64;
 277 ;CHECK-NEXT: st.local.u64 [r0], rd1;
 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 r0, array_local_float;
 287 ;CHECK-NEXT: st.local.f32 [r0], f1;
 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 r0, array_local_double;
 297 ;CHECK-NEXT: st.local.f64 [r0], fd1;
 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 r0, array_shared_i16;
 307 ;CHECK-NEXT: st.shared.u16 [r0], rh1;
 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 r0, array_shared_i32;
 317 ;CHECK-NEXT: st.shared.u32 [r0], r1;
 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 r0, array_shared_i64;
 327 ;CHECK-NEXT: st.shared.u64 [r0], rd1;
 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 r0, array_shared_float;
 337 ;CHECK-NEXT: st.shared.f32 [r0], f1;
 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 r0, array_shared_double;
 347 ;CHECK-NEXT: st.shared.f64 [r0], fd1;
 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 r0, array_i16;
 357 ;CHECK-NEXT: st.global.u16 [r0+2], rh1;
 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 r0, array_i32;
 367 ;CHECK-NEXT: st.global.u32 [r0+4], r1;
 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 r0, array_i64;
 377 ;CHECK-NEXT: st.global.u64 [r0+8], rd1;
 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 r0, array_float;
 387 ;CHECK-NEXT: st.global.f32 [r0+4], f1;
 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 r0, array_double;
 397 ;CHECK-NEXT: st.global.f64 [r0+8], fd1;
 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 }