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