summaryrefslogtreecommitdiff
path: root/test/CodeGen
diff options
context:
space:
mode:
authorJustin Holewinski <justin.holewinski@gmail.com>2011-03-08 14:10:18 +0000
committerJustin Holewinski <justin.holewinski@gmail.com>2011-03-08 14:10:18 +0000
commitc6f24f4086c786e8ee8e74193ea9a08e90853e0b (patch)
tree474f4ba654fff8f991fc88672efcc0b13a878b7b /test/CodeGen
parent86e4eeb6c79baf1c4bc4569a1e6ffc0c8f4b52b9 (diff)
downloadllvm-c6f24f4086c786e8ee8e74193ea9a08e90853e0b.tar.gz
llvm-c6f24f4086c786e8ee8e74193ea9a08e90853e0b.tar.bz2
llvm-c6f24f4086c786e8ee8e74193ea9a08e90853e0b.tar.xz
PTX: Add intrinsic support for ntid, ctaid, and nctaid registers
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@127246 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGen')
-rw-r--r--test/CodeGen/PTX/intrinsic.ll96
1 files changed, 96 insertions, 0 deletions
diff --git a/test/CodeGen/PTX/intrinsic.ll b/test/CodeGen/PTX/intrinsic.ll
index 24c328b773..804d8b5f95 100644
--- a/test/CodeGen/PTX/intrinsic.ll
+++ b/test/CodeGen/PTX/intrinsic.ll
@@ -28,6 +28,90 @@ define ptx_device i16 @tid_w() {
ret i16 %x
}
+define ptx_device i16 @ntid_x() {
+; CHECK: mov.u16 rh0, ntid.x;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ntid.x()
+ ret i16 %x
+}
+
+define ptx_device i16 @ntid_y() {
+; CHECK: mov.u16 rh0, ntid.y;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ntid.y()
+ ret i16 %x
+}
+
+define ptx_device i16 @ntid_z() {
+; CHECK: mov.u16 rh0, ntid.z;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ntid.z()
+ ret i16 %x
+}
+
+define ptx_device i16 @ntid_w() {
+; CHECK: mov.u16 rh0, ntid.w;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ntid.w()
+ ret i16 %x
+}
+
+define ptx_device i16 @ctaid_x() {
+; CHECK: mov.u16 rh0, ctaid.x;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ctaid.x()
+ ret i16 %x
+}
+
+define ptx_device i16 @ctaid_y() {
+; CHECK: mov.u16 rh0, ctaid.y;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ctaid.y()
+ ret i16 %x
+}
+
+define ptx_device i16 @ctaid_z() {
+; CHECK: mov.u16 rh0, ctaid.z;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ctaid.z()
+ ret i16 %x
+}
+
+define ptx_device i16 @ctaid_w() {
+; CHECK: mov.u16 rh0, ctaid.w;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.ctaid.w()
+ ret i16 %x
+}
+
+define ptx_device i16 @nctaid_x() {
+; CHECK: mov.u16 rh0, nctaid.x;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.nctaid.x()
+ ret i16 %x
+}
+
+define ptx_device i16 @nctaid_y() {
+; CHECK: mov.u16 rh0, nctaid.y;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.nctaid.y()
+ ret i16 %x
+}
+
+define ptx_device i16 @nctaid_z() {
+; CHECK: mov.u16 rh0, nctaid.z;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.nctaid.z()
+ ret i16 %x
+}
+
+define ptx_device i16 @nctaid_w() {
+; CHECK: mov.u16 rh0, nctaid.w;
+; CHECK-NEXT: ret;
+ %x = call i16 @llvm.ptx.read.nctaid.w()
+ ret i16 %x
+}
+
define ptx_device void @bar_sync() {
; CHECK: bar.sync 0
; CHECK-NEXT: ret;
@@ -39,5 +123,17 @@ 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 i16 @llvm.ptx.read.ntid.x()
+declare i16 @llvm.ptx.read.ntid.y()
+declare i16 @llvm.ptx.read.ntid.z()
+declare i16 @llvm.ptx.read.ntid.w()
+declare i16 @llvm.ptx.read.ctaid.x()
+declare i16 @llvm.ptx.read.ctaid.y()
+declare i16 @llvm.ptx.read.ctaid.z()
+declare i16 @llvm.ptx.read.ctaid.w()
+declare i16 @llvm.ptx.read.nctaid.x()
+declare i16 @llvm.ptx.read.nctaid.y()
+declare i16 @llvm.ptx.read.nctaid.z()
+declare i16 @llvm.ptx.read.nctaid.w()
declare void @llvm.ptx.bar.sync(i32 %i)