test/CodeGen/PTX/intrinsic.ll

   1 ; RUN: llc < %s -march=ptx | FileCheck %s
   2
   3 define ptx_device i16 @tid_x() {
   4 ; CHECK: mov.u16 rh0, tid.x;
   5 ; CHECK-NEXT: ret;
   6         %x = call i16 @llvm.ptx.read.tid.x()
   7         ret i16 %x
   8 }
   9
  10 define ptx_device i16 @tid_y() {
  11 ; CHECK: mov.u16 rh0, tid.y;
  12 ; CHECK-NEXT: ret;
  13         %x = call i16 @llvm.ptx.read.tid.y()
  14         ret i16 %x
  15 }
  16
  17 define ptx_device i16 @tid_z() {
  18 ; CHECK: mov.u16 rh0, tid.z;
  19 ; CHECK-NEXT: ret;
  20         %x = call i16 @llvm.ptx.read.tid.z()
  21         ret i16 %x
  22 }
  23
  24 define ptx_device i16 @tid_w() {
  25 ; CHECK: mov.u16 rh0, tid.w;
  26 ; CHECK-NEXT: ret;
  27         %x = call i16 @llvm.ptx.read.tid.w()
  28         ret i16 %x
  29 }
  30
  31 define ptx_device void @bar_sync() {
  32 ; CHECK: bar.sync 0
  33 ; CHECK-NEXT: ret;
  34         call void @llvm.ptx.bar.sync(i32 0)
  35         ret void
  36 }
  37
  38 declare i16 @llvm.ptx.read.tid.x()
  39 declare i16 @llvm.ptx.read.tid.y()
  40 declare i16 @llvm.ptx.read.tid.z()
  41 declare i16 @llvm.ptx.read.tid.w()
  42
  43 declare void @llvm.ptx.bar.sync(i32 %i)