test/CodeGen/PTX/intrinsic.ll

   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)