test/CodeGen/PTX/st.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 void @t1_u16(i16* %p, i16 %x) {
  50 entry:
  51 ;CHECK: st.global.u16 [%r{{[0-9]+}}], %rh{{[0-9]+}};
  52 ;CHECK: ret;
  53   store i16 %x, i16* %p
  54   ret void
  55 }
  56
  57 define ptx_device void @t1_u32(i32* %p, i32 %x) {
  58 entry:
  59 ;CHECK: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}};
  60 ;CHECK: ret;
  61   store i32 %x, i32* %p
  62   ret void
  63 }
  64
  65 define ptx_device void @t1_u64(i64* %p, i64 %x) {
  66 entry:
  67 ;CHECK: st.global.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}};
  68 ;CHECK: ret;
  69   store i64 %x, i64* %p
  70   ret void
  71 }
  72
  73 define ptx_device void @t1_f32(float* %p, float %x) {
  74 entry:
  75 ;CHECK: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}};
  76 ;CHECK: ret;
  77   store float %x, float* %p
  78   ret void
  79 }
  80
  81 define ptx_device void @t1_f64(double* %p, double %x) {
  82 entry:
  83 ;CHECK: st.global.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}};
  84 ;CHECK: ret;
  85   store double %x, double* %p
  86   ret void
  87 }
  88
  89 define ptx_device void @t2_u16(i16* %p, i16 %x) {
  90 entry:
  91 ;CHECK: st.global.u16 [%r{{[0-9]+}}+2], %rh{{[0-9]+}};
  92 ;CHECK: ret;
  93   %i = getelementptr i16* %p, i32 1
  94   store i16 %x, i16* %i
  95   ret void
  96 }
  97
  98 define ptx_device void @t2_u32(i32* %p, i32 %x) {
  99 entry:
 100 ;CHECK: st.global.u32 [%r{{[0-9]+}}+4], %r{{[0-9]+}};
 101 ;CHECK: ret;
 102   %i = getelementptr i32* %p, i32 1
 103   store i32 %x, i32* %i
 104   ret void
 105 }
 106
 107 define ptx_device void @t2_u64(i64* %p, i64 %x) {
 108 entry:
 109 ;CHECK: st.global.u64 [%r{{[0-9]+}}+8], %rd{{[0-9]+}};
 110 ;CHECK: ret;
 111   %i = getelementptr i64* %p, i32 1
 112   store i64 %x, i64* %i
 113   ret void
 114 }
 115
 116 define ptx_device void @t2_f32(float* %p, float %x) {
 117 entry:
 118 ;CHECK: st.global.f32 [%r{{[0-9]+}}+4], %f{{[0-9]+}};
 119 ;CHECK: ret;
 120   %i = getelementptr float* %p, i32 1
 121   store float %x, float* %i
 122   ret void
 123 }
 124
 125 define ptx_device void @t2_f64(double* %p, double %x) {
 126 entry:
 127 ;CHECK: st.global.f64 [%r{{[0-9]+}}+8], %fd{{[0-9]+}};
 128 ;CHECK: ret;
 129   %i = getelementptr double* %p, i32 1
 130   store double %x, double* %i
 131   ret void
 132 }
 133
 134 define ptx_device void @t3_u16(i16* %p, i32 %q, i16 %x) {
 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: st.global.u16 [%r{{[0-9]+}}], %rh{{[0-9]+}};
 139 ;CHECK: ret;
 140   %i = getelementptr i16* %p, i32 %q
 141   store i16 %x, i16* %i
 142   ret void
 143 }
 144
 145 define ptx_device void @t3_u32(i32* %p, i32 %q, i32 %x) {
 146 entry:
 147 ;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 2;
 148 ;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]];
 149 ;CHECK: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}};
 150 ;CHECK: ret;
 151   %i = getelementptr i32* %p, i32 %q
 152   store i32 %x, i32* %i
 153   ret void
 154 }
 155
 156 define ptx_device void @t3_u64(i64* %p, i32 %q, i64 %x) {
 157 entry:
 158 ;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 3;
 159 ;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]];
 160 ;CHECK: st.global.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}};
 161 ;CHECK: ret;
 162   %i = getelementptr i64* %p, i32 %q
 163   store i64 %x, i64* %i
 164   ret void
 165 }
 166
 167 define ptx_device void @t3_f32(float* %p, i32 %q, float %x) {
 168 entry:
 169 ;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 2;
 170 ;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]];
 171 ;CHECK: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}};
 172 ;CHECK: ret;
 173   %i = getelementptr float* %p, i32 %q
 174   store float %x, float* %i
 175   ret void
 176 }
 177
 178 define ptx_device void @t3_f64(double* %p, i32 %q, double %x) {
 179 entry:
 180 ;CHECK: shl.b32 %r[[R0:[0-9]+]], %r{{[0-9]+}}, 3;
 181 ;CHECK: add.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r[[R0]];
 182 ;CHECK: st.global.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}};
 183 ;CHECK: ret;
 184   %i = getelementptr double* %p, i32 %q
 185   store double %x, double* %i
 186   ret void
 187 }
 188
 189 define ptx_device void @t4_global_u16(i16 %x) {
 190 entry:
 191 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i16;
 192 ;CHECK: st.global.u16 [%r[[R0]]], %rh{{[0-9]+}};
 193 ;CHECK: ret;
 194   %i = getelementptr [10 x i16]* @array_i16, i16 0, i16 0
 195   store i16 %x, i16* %i
 196   ret void
 197 }
 198
 199 define ptx_device void @t4_global_u32(i32 %x) {
 200 entry:
 201 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i32;
 202 ;CHECK: st.global.u32 [%r[[R0]]], %r{{[0-9]+}};
 203 ;CHECK: ret;
 204   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 0
 205   store i32 %x, i32* %i
 206   ret void
 207 }
 208
 209 define ptx_device void @t4_global_u64(i64 %x) {
 210 entry:
 211 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i64;
 212 ;CHECK: st.global.u64 [%r[[R0]]], %rd{{[0-9]+}};
 213 ;CHECK: ret;
 214   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 0
 215   store i64 %x, i64* %i
 216   ret void
 217 }
 218
 219 define ptx_device void @t4_global_f32(float %x) {
 220 entry:
 221 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_float;
 222 ;CHECK: st.global.f32 [%r[[R0]]], %f{{[0-9]+}};
 223 ;CHECK: ret;
 224   %i = getelementptr [10 x float]* @array_float, i32 0, i32 0
 225   store float %x, float* %i
 226   ret void
 227 }
 228
 229 define ptx_device void @t4_global_f64(double %x) {
 230 entry:
 231 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_double;
 232 ;CHECK: st.global.f64 [%r[[R0]]], %fd{{[0-9]+}};
 233 ;CHECK: ret;
 234   %i = getelementptr [10 x double]* @array_double, i32 0, i32 0
 235   store double %x, double* %i
 236   ret void
 237 }
 238
 239 define ptx_device void @t4_shared_u16(i16 %x) {
 240 entry:
 241 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_i16;
 242 ;CHECK: st.shared.u16 [%r[[R0]]], %rh{{[0-9]+}};
 243 ;CHECK: ret;
 244   %i = getelementptr [10 x i16] addrspace(4)* @array_shared_i16, i32 0, i32 0
 245   store i16 %x, i16 addrspace(4)* %i
 246   ret void
 247 }
 248
 249 define ptx_device void @t4_shared_u32(i32 %x) {
 250 entry:
 251 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_i32;
 252 ;CHECK: st.shared.u32 [%r[[R0]]], %r{{[0-9]+}};
 253 ;CHECK: ret;
 254   %i = getelementptr [10 x i32] addrspace(4)* @array_shared_i32, i32 0, i32 0
 255   store i32 %x, i32 addrspace(4)* %i
 256   ret void
 257 }
 258
 259 define ptx_device void @t4_shared_u64(i64 %x) {
 260 entry:
 261 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_i64;
 262 ;CHECK: st.shared.u64 [%r[[R0]]], %rd{{[0-9]+}};
 263 ;CHECK: ret;
 264   %i = getelementptr [10 x i64] addrspace(4)* @array_shared_i64, i32 0, i32 0
 265   store i64 %x, i64 addrspace(4)* %i
 266   ret void
 267 }
 268
 269 define ptx_device void @t4_shared_f32(float %x) {
 270 entry:
 271 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_float;
 272 ;CHECK: st.shared.f32 [%r[[R0]]], %f{{[0-9]+}};
 273 ;CHECK: ret;
 274   %i = getelementptr [10 x float] addrspace(4)* @array_shared_float, i32 0, i32 0
 275   store float %x, float addrspace(4)* %i
 276   ret void
 277 }
 278
 279 define ptx_device void @t4_shared_f64(double %x) {
 280 entry:
 281 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_shared_double;
 282 ;CHECK: st.shared.f64 [%r[[R0]]], %fd{{[0-9]+}};
 283 ;CHECK: ret;
 284   %i = getelementptr [10 x double] addrspace(4)* @array_shared_double, i32 0, i32 0
 285   store double %x, double addrspace(4)* %i
 286   ret void
 287 }
 288
 289 define ptx_device void @t5_u16(i16 %x) {
 290 entry:
 291 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i16;
 292 ;CHECK: st.global.u16 [%r[[R0]]+2], %rh{{[0-9]+}};
 293 ;CHECK: ret;
 294   %i = getelementptr [10 x i16]* @array_i16, i32 0, i32 1
 295   store i16 %x, i16* %i
 296   ret void
 297 }
 298
 299 define ptx_device void @t5_u32(i32 %x) {
 300 entry:
 301 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i32;
 302 ;CHECK: st.global.u32 [%r[[R0]]+4], %r{{[0-9]+}};
 303 ;CHECK: ret;
 304   %i = getelementptr [10 x i32]* @array_i32, i32 0, i32 1
 305   store i32 %x, i32* %i
 306   ret void
 307 }
 308
 309 define ptx_device void @t5_u64(i64 %x) {
 310 entry:
 311 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_i64;
 312 ;CHECK: st.global.u64 [%r[[R0]]+8], %rd{{[0-9]+}};
 313 ;CHECK: ret;
 314   %i = getelementptr [10 x i64]* @array_i64, i32 0, i32 1
 315   store i64 %x, i64* %i
 316   ret void
 317 }
 318
 319 define ptx_device void @t5_f32(float %x) {
 320 entry:
 321 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_float;
 322 ;CHECK: st.global.f32 [%r[[R0]]+4], %f{{[0-9]+}};
 323 ;CHECK: ret;
 324   %i = getelementptr [10 x float]* @array_float, i32 0, i32 1
 325   store float %x, float* %i
 326   ret void
 327 }
 328
 329 define ptx_device void @t5_f64(double %x) {
 330 entry:
 331 ;CHECK: mov.u32 %r[[R0:[0-9]+]], array_double;
 332 ;CHECK: st.global.f64 [%r[[R0]]+8], %fd{{[0-9]+}};
 333 ;CHECK: ret;
 334   %i = getelementptr [10 x double]* @array_double, i32 0, i32 1
 335   store double %x, double* %i
 336   ret void
 337 }