Home | History | Annotate | Download | only in PTX 
      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 .shared .b8 array_shared_i16[20];
     10 @array_shared_i16 = external addrspace(4) global [10 x i16]
     11 
     12 ;CHECK: .extern .global .b8 array_i32[40];
     13 @array_i32 = external global [10 x i32]
     14 
     15 ;CHECK: .extern .const .b8 array_constant_i32[40];
     16 @array_constant_i32 = external addrspace(1) constant [10 x i32]
     17 
     18 ;CHECK: .extern .shared .b8 array_shared_i32[40];
     19 @array_shared_i32 = external addrspace(4) global [10 x i32]
     20 
     21 ;CHECK: .extern .global .b8 array_i64[80];
     22 @array_i64 = external global [10 x i64]
     23 
     24 ;CHECK: .extern .const .b8 array_constant_i64[80];
     25 @array_constant_i64 = external addrspace(1) constant [10 x i64]
     26 
     27 ;CHECK: .extern .shared .b8 array_shared_i64[80];
     28 @array_shared_i64 = external addrspace(4) global [10 x i64]
     29 
     30 ;CHECK: .extern .global .b8 array_float[40];
     31 @array_float = external global [10 x float]
     32 
     33 ;CHECK: .extern .const .b8 array_constant_float[40];
     34 @array_constant_float = external addrspace(1) constant [10 x float]
     35 
     36 ;CHECK: .extern .shared .b8 array_shared_float[40];
     37 @array_shared_float = external addrspace(4) global [10 x float]
     38 
     39 ;CHECK: .extern .global .b8 array_double[80];
     40 @array_double = external global [10 x double]
     41 
     42 ;CHECK: .extern .const .b8 array_constant_double[80];
     43 @array_constant_double = external addrspace(1) constant [10 x double]
     44 
     45 ;CHECK: .extern .shared .b8 array_shared_double[80];
     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 }
    338