diff options
| author | Che-Liang Chiou <clchiou@gmail.com> | 2011-03-05 14:17:37 +0000 |
|---|---|---|
| committer | Che-Liang Chiou <clchiou@gmail.com> | 2011-03-05 14:17:37 +0000 |
| commit | 0df2c50c2bbf7b53f9abe519d655b70e92f65864 (patch) | |
| tree | d870c42485b6a314aa1b74bfe61d04e06cf0970f /test/CodeGen/PTX | |
| parent | afc7d235e91a769f74d87bbe745558ed1b692ff7 (diff) | |
| download | external_llvm-0df2c50c2bbf7b53f9abe519d655b70e92f65864.zip external_llvm-0df2c50c2bbf7b53f9abe519d655b70e92f65864.tar.gz external_llvm-0df2c50c2bbf7b53f9abe519d655b70e92f65864.tar.bz2 | |
ptx: add basic intrinsic support
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@127084 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGen/PTX')
| -rw-r--r-- | test/CodeGen/PTX/intrinsic.ll | 43 |
1 files changed, 43 insertions, 0 deletions
diff --git a/test/CodeGen/PTX/intrinsic.ll b/test/CodeGen/PTX/intrinsic.ll new file mode 100644 index 0000000..24c328b --- /dev/null +++ b/test/CodeGen/PTX/intrinsic.ll @@ -0,0 +1,43 @@ +; RUN: llc < %s -march=ptx | FileCheck %s + +define ptx_device i16 @tid_x() { +; CHECK: mov.u16 rh0, tid.x; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.tid.x() + ret i16 %x +} + +define ptx_device i16 @tid_y() { +; CHECK: mov.u16 rh0, tid.y; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.tid.y() + ret i16 %x +} + +define ptx_device i16 @tid_z() { +; CHECK: mov.u16 rh0, tid.z; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.tid.z() + ret i16 %x +} + +define ptx_device i16 @tid_w() { +; CHECK: mov.u16 rh0, tid.w; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.tid.w() + ret i16 %x +} + +define ptx_device void @bar_sync() { +; CHECK: bar.sync 0 +; CHECK-NEXT: ret; + call void @llvm.ptx.bar.sync(i32 0) + ret void +} + +declare i16 @llvm.ptx.read.tid.x() +declare i16 @llvm.ptx.read.tid.y() +declare i16 @llvm.ptx.read.tid.z() +declare i16 @llvm.ptx.read.tid.w() + +declare void @llvm.ptx.bar.sync(i32 %i) |