PTX: Add intrinsic support for ntid, ctaid, and nctaid registers

llvm-svn: 127246

1 files changed, 96 insertions, 0 deletions

diff --git a/llvm/test/CodeGen/PTX/intrinsic.ll b/llvm/test/CodeGen/PTX/intrinsic.ll
index 24c328b7736..804d8b5f95d 100644
--- a/llvm/test/CodeGen/PTX/intrinsic.ll
+++ b/llvm/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)