Home | History | Annotate | Download | only in NVPTX 
      1 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
      2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
      3 
      4 define ptx_device i32 @test_tid_x() {
      5 ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.x;
      6 ; CHECK: ret;
      7 	%x = call i32 @llvm.ptx.read.tid.x()
      8 	ret i32 %x
      9 }
     10 
     11 define ptx_device i32 @test_tid_y() {
     12 ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.y;
     13 ; CHECK: ret;
     14 	%x = call i32 @llvm.ptx.read.tid.y()
     15 	ret i32 %x
     16 }
     17 
     18 define ptx_device i32 @test_tid_z() {
     19 ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.z;
     20 ; CHECK: ret;
     21 	%x = call i32 @llvm.ptx.read.tid.z()
     22 	ret i32 %x
     23 }
     24 
     25 define ptx_device i32 @test_tid_w() {
     26 ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.w;
     27 ; CHECK: ret;
     28 	%x = call i32 @llvm.ptx.read.tid.w()
     29 	ret i32 %x
     30 }
     31 
     32 define ptx_device i32 @test_ntid_x() {
     33 ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.x;
     34 ; CHECK: ret;
     35 	%x = call i32 @llvm.ptx.read.ntid.x()
     36 	ret i32 %x
     37 }
     38 
     39 define ptx_device i32 @test_ntid_y() {
     40 ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.y;
     41 ; CHECK: ret;
     42 	%x = call i32 @llvm.ptx.read.ntid.y()
     43 	ret i32 %x
     44 }
     45 
     46 define ptx_device i32 @test_ntid_z() {
     47 ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.z;
     48 ; CHECK: ret;
     49 	%x = call i32 @llvm.ptx.read.ntid.z()
     50 	ret i32 %x
     51 }
     52 
     53 define ptx_device i32 @test_ntid_w() {
     54 ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.w;
     55 ; CHECK: ret;
     56 	%x = call i32 @llvm.ptx.read.ntid.w()
     57 	ret i32 %x
     58 }
     59 
     60 define ptx_device i32 @test_laneid() {
     61 ; CHECK: mov.u32 %r{{[0-9]+}}, %laneid;
     62 ; CHECK: ret;
     63 	%x = call i32 @llvm.ptx.read.laneid()
     64 	ret i32 %x
     65 }
     66 
     67 define ptx_device i32 @test_warpid() {
     68 ; CHECK: mov.u32 %r{{[0-9]+}}, %warpid;
     69 ; CHECK: ret;
     70 	%x = call i32 @llvm.ptx.read.warpid()
     71 	ret i32 %x
     72 }
     73 
     74 define ptx_device i32 @test_nwarpid() {
     75 ; CHECK: mov.u32 %r{{[0-9]+}}, %nwarpid;
     76 ; CHECK: ret;
     77 	%x = call i32 @llvm.ptx.read.nwarpid()
     78 	ret i32 %x
     79 }
     80 
     81 define ptx_device i32 @test_ctaid_x() {
     82 ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.x;
     83 ; CHECK: ret;
     84 	%x = call i32 @llvm.ptx.read.ctaid.x()
     85 	ret i32 %x
     86 }
     87 
     88 define ptx_device i32 @test_ctaid_y() {
     89 ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.y;
     90 ; CHECK: ret;
     91 	%x = call i32 @llvm.ptx.read.ctaid.y()
     92 	ret i32 %x
     93 }
     94 
     95 define ptx_device i32 @test_ctaid_z() {
     96 ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.z;
     97 ; CHECK: ret;
     98 	%x = call i32 @llvm.ptx.read.ctaid.z()
     99 	ret i32 %x
    100 }
    101 
    102 define ptx_device i32 @test_ctaid_w() {
    103 ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.w;
    104 ; CHECK: ret;
    105 	%x = call i32 @llvm.ptx.read.ctaid.w()
    106 	ret i32 %x
    107 }
    108 
    109 define ptx_device i32 @test_nctaid_x() {
    110 ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;
    111 ; CHECK: ret;
    112 	%x = call i32 @llvm.ptx.read.nctaid.x()
    113 	ret i32 %x
    114 }
    115 
    116 define ptx_device i32 @test_nctaid_y() {
    117 ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.y;
    118 ; CHECK: ret;
    119 	%x = call i32 @llvm.ptx.read.nctaid.y()
    120 	ret i32 %x
    121 }
    122 
    123 define ptx_device i32 @test_nctaid_z() {
    124 ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.z;
    125 ; CHECK: ret;
    126 	%x = call i32 @llvm.ptx.read.nctaid.z()
    127 	ret i32 %x
    128 }
    129 
    130 define ptx_device i32 @test_nctaid_w() {
    131 ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.w;
    132 ; CHECK: ret;
    133 	%x = call i32 @llvm.ptx.read.nctaid.w()
    134 	ret i32 %x
    135 }
    136 
    137 define ptx_device i32 @test_smid() {
    138 ; CHECK: mov.u32 %r{{[0-9]+}}, %smid;
    139 ; CHECK: ret;
    140 	%x = call i32 @llvm.ptx.read.smid()
    141 	ret i32 %x
    142 }
    143 
    144 define ptx_device i32 @test_nsmid() {
    145 ; CHECK: mov.u32 %r{{[0-9]+}}, %nsmid;
    146 ; CHECK: ret;
    147 	%x = call i32 @llvm.ptx.read.nsmid()
    148 	ret i32 %x
    149 }
    150 
    151 define ptx_device i32 @test_gridid() {
    152 ; CHECK: mov.u32 %r{{[0-9]+}}, %gridid;
    153 ; CHECK: ret;
    154 	%x = call i32 @llvm.ptx.read.gridid()
    155 	ret i32 %x
    156 }
    157 
    158 define ptx_device i32 @test_lanemask_eq() {
    159 ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_eq;
    160 ; CHECK: ret;
    161 	%x = call i32 @llvm.ptx.read.lanemask.eq()
    162 	ret i32 %x
    163 }
    164 
    165 define ptx_device i32 @test_lanemask_le() {
    166 ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_le;
    167 ; CHECK: ret;
    168 	%x = call i32 @llvm.ptx.read.lanemask.le()
    169 	ret i32 %x
    170 }
    171 
    172 define ptx_device i32 @test_lanemask_lt() {
    173 ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_lt;
    174 ; CHECK: ret;
    175 	%x = call i32 @llvm.ptx.read.lanemask.lt()
    176 	ret i32 %x
    177 }
    178 
    179 define ptx_device i32 @test_lanemask_ge() {
    180 ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_ge;
    181 ; CHECK: ret;
    182 	%x = call i32 @llvm.ptx.read.lanemask.ge()
    183 	ret i32 %x
    184 }
    185 
    186 define ptx_device i32 @test_lanemask_gt() {
    187 ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_gt;
    188 ; CHECK: ret;
    189 	%x = call i32 @llvm.ptx.read.lanemask.gt()
    190 	ret i32 %x
    191 }
    192 
    193 define ptx_device i32 @test_clock() {
    194 ; CHECK: mov.u32 %r{{[0-9]+}}, %clock;
    195 ; CHECK: ret;
    196 	%x = call i32 @llvm.ptx.read.clock()
    197 	ret i32 %x
    198 }
    199 
    200 define ptx_device i64 @test_clock64() {
    201 ; CHECK: mov.u64 %rd{{[0-9]+}}, %clock64;
    202 ; CHECK: ret;
    203 	%x = call i64 @llvm.ptx.read.clock64()
    204 	ret i64 %x
    205 }
    206 
    207 define ptx_device i32 @test_pm0() {
    208 ; CHECK: mov.u32 %r{{[0-9]+}}, %pm0;
    209 ; CHECK: ret;
    210 	%x = call i32 @llvm.ptx.read.pm0()
    211 	ret i32 %x
    212 }
    213 
    214 define ptx_device i32 @test_pm1() {
    215 ; CHECK: mov.u32 %r{{[0-9]+}}, %pm1;
    216 ; CHECK: ret;
    217 	%x = call i32 @llvm.ptx.read.pm1()
    218 	ret i32 %x
    219 }
    220 
    221 define ptx_device i32 @test_pm2() {
    222 ; CHECK: mov.u32 %r{{[0-9]+}}, %pm2;
    223 ; CHECK: ret;
    224 	%x = call i32 @llvm.ptx.read.pm2()
    225 	ret i32 %x
    226 }
    227 
    228 define ptx_device i32 @test_pm3() {
    229 ; CHECK: mov.u32 %r{{[0-9]+}}, %pm3;
    230 ; CHECK: ret;
    231 	%x = call i32 @llvm.ptx.read.pm3()
    232 	ret i32 %x
    233 }
    234 
    235 define ptx_device void @test_bar_sync() {
    236 ; CHECK: bar.sync 0
    237 ; CHECK: ret;
    238 	call void @llvm.ptx.bar.sync(i32 0)
    239 	ret void
    240 }
    241 
    242 declare i32 @llvm.ptx.read.tid.x()
    243 declare i32 @llvm.ptx.read.tid.y()
    244 declare i32 @llvm.ptx.read.tid.z()
    245 declare i32 @llvm.ptx.read.tid.w()
    246 declare i32 @llvm.ptx.read.ntid.x()
    247 declare i32 @llvm.ptx.read.ntid.y()
    248 declare i32 @llvm.ptx.read.ntid.z()
    249 declare i32 @llvm.ptx.read.ntid.w()
    250 
    251 declare i32 @llvm.ptx.read.laneid()
    252 declare i32 @llvm.ptx.read.warpid()
    253 declare i32 @llvm.ptx.read.nwarpid()
    254 
    255 declare i32 @llvm.ptx.read.ctaid.x()
    256 declare i32 @llvm.ptx.read.ctaid.y()
    257 declare i32 @llvm.ptx.read.ctaid.z()
    258 declare i32 @llvm.ptx.read.ctaid.w()
    259 declare i32 @llvm.ptx.read.nctaid.x()
    260 declare i32 @llvm.ptx.read.nctaid.y()
    261 declare i32 @llvm.ptx.read.nctaid.z()
    262 declare i32 @llvm.ptx.read.nctaid.w()
    263 
    264 declare i32 @llvm.ptx.read.smid()
    265 declare i32 @llvm.ptx.read.nsmid()
    266 declare i32 @llvm.ptx.read.gridid()
    267 
    268 declare i32 @llvm.ptx.read.lanemask.eq()
    269 declare i32 @llvm.ptx.read.lanemask.le()
    270 declare i32 @llvm.ptx.read.lanemask.lt()
    271 declare i32 @llvm.ptx.read.lanemask.ge()
    272 declare i32 @llvm.ptx.read.lanemask.gt()
    273 
    274 declare i32 @llvm.ptx.read.clock()
    275 declare i64 @llvm.ptx.read.clock64()
    276 
    277 declare i32 @llvm.ptx.read.pm0()
    278 declare i32 @llvm.ptx.read.pm1()
    279 declare i32 @llvm.ptx.read.pm2()
    280 declare i32 @llvm.ptx.read.pm3()
    281 
    282 declare void @llvm.ptx.bar.sync(i32 %i)
    283