test/CodeGen/NVPTX/intrinsic-old.ll

   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)