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