test/CodeGen/NVPTX/intrinsic-old.ll

   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)