[llvm-commits] [llvm] r127246 - in /llvm/trunk: include/llvm/IntrinsicsPTX.td lib/Target/PTX/PTXIntrinsicInstrInfo.td test/CodeGen/PTX/intrinsic.ll

Justin Holewinski justin.holewinski at gmail.com
Tue Mar 8 06:10:18 PST 2011


Author: jholewinski
Date: Tue Mar  8 08:10:18 2011
New Revision: 127246

URL: http://llvm.org/viewvc/llvm-project?rev=127246&view=rev
Log:
PTX: Add intrinsic support for ntid, ctaid, and nctaid registers

Modified:
    llvm/trunk/include/llvm/IntrinsicsPTX.td
    llvm/trunk/lib/Target/PTX/PTXIntrinsicInstrInfo.td
    llvm/trunk/test/CodeGen/PTX/intrinsic.ll

Modified: llvm/trunk/include/llvm/IntrinsicsPTX.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IntrinsicsPTX.td?rev=127246&r1=127245&r2=127246&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IntrinsicsPTX.td (original)
+++ llvm/trunk/include/llvm/IntrinsicsPTX.td Tue Mar  8 08:10:18 2011
@@ -25,8 +25,14 @@
   }
 }
 
-defm int_ptx_read_tid : PTXReadSpecialRegisterIntrinsic;
-defm int_ptx_read_tid : PTXReadSpecialSubRegisterIntrinsic;
+defm int_ptx_read_tid    : PTXReadSpecialRegisterIntrinsic;
+defm int_ptx_read_tid    : PTXReadSpecialSubRegisterIntrinsic;
+defm int_ptx_read_ntid   : PTXReadSpecialRegisterIntrinsic;
+defm int_ptx_read_ntid   : PTXReadSpecialSubRegisterIntrinsic;
+defm int_ptx_read_ctaid  : PTXReadSpecialRegisterIntrinsic;
+defm int_ptx_read_ctaid  : PTXReadSpecialSubRegisterIntrinsic;
+defm int_ptx_read_nctaid : PTXReadSpecialRegisterIntrinsic;
+defm int_ptx_read_nctaid : PTXReadSpecialSubRegisterIntrinsic;
 
 let TargetPrefix = "ptx" in
   def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>;

Modified: llvm/trunk/lib/Target/PTX/PTXIntrinsicInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXIntrinsicInstrInfo.td?rev=127246&r1=127245&r2=127246&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXIntrinsicInstrInfo.td (original)
+++ llvm/trunk/lib/Target/PTX/PTXIntrinsicInstrInfo.td Tue Mar  8 08:10:18 2011
@@ -24,10 +24,28 @@
             [(set RRegu16:$d, (intop))]>;
 
 def PTX_READ_TID_R64 : PTX_READ_SPECIAL_REGISTER<"tid", int_ptx_read_tid_r64>;
-def PTX_READ_TID_X : PTX_READ_SPECIAL_SUB_REGISTER<"tid.x", int_ptx_read_tid_x>;
-def PTX_READ_TID_Y : PTX_READ_SPECIAL_SUB_REGISTER<"tid.y", int_ptx_read_tid_y>;
-def PTX_READ_TID_Z : PTX_READ_SPECIAL_SUB_REGISTER<"tid.z", int_ptx_read_tid_z>;
-def PTX_READ_TID_W : PTX_READ_SPECIAL_SUB_REGISTER<"tid.w", int_ptx_read_tid_w>;
+def PTX_READ_TID_X   : PTX_READ_SPECIAL_SUB_REGISTER<"tid.x", int_ptx_read_tid_x>;
+def PTX_READ_TID_Y   : PTX_READ_SPECIAL_SUB_REGISTER<"tid.y", int_ptx_read_tid_y>;
+def PTX_READ_TID_Z   : PTX_READ_SPECIAL_SUB_REGISTER<"tid.z", int_ptx_read_tid_z>;
+def PTX_READ_TID_W   : PTX_READ_SPECIAL_SUB_REGISTER<"tid.w", int_ptx_read_tid_w>;
+
+def PTX_READ_NTID_R64 : PTX_READ_SPECIAL_REGISTER<"ntid", int_ptx_read_ntid_r64>;
+def PTX_READ_NTID_X   : PTX_READ_SPECIAL_SUB_REGISTER<"ntid.x", int_ptx_read_ntid_x>;
+def PTX_READ_NTID_Y   : PTX_READ_SPECIAL_SUB_REGISTER<"ntid.y", int_ptx_read_ntid_y>;
+def PTX_READ_NTID_Z   : PTX_READ_SPECIAL_SUB_REGISTER<"ntid.z", int_ptx_read_ntid_z>;
+def PTX_READ_NTID_W   : PTX_READ_SPECIAL_SUB_REGISTER<"ntid.w", int_ptx_read_ntid_w>;
+
+def PTX_READ_CTAID_R64 : PTX_READ_SPECIAL_REGISTER<"ctaid", int_ptx_read_ctaid_r64>;
+def PTX_READ_CTAID_X   : PTX_READ_SPECIAL_SUB_REGISTER<"ctaid.x", int_ptx_read_ctaid_x>;
+def PTX_READ_CTAID_Y   : PTX_READ_SPECIAL_SUB_REGISTER<"ctaid.y", int_ptx_read_ctaid_y>;
+def PTX_READ_CTAID_Z   : PTX_READ_SPECIAL_SUB_REGISTER<"ctaid.z", int_ptx_read_ctaid_z>;
+def PTX_READ_CTAID_W   : PTX_READ_SPECIAL_SUB_REGISTER<"ctaid.w", int_ptx_read_ctaid_w>;
+
+def PTX_READ_NCTAID_R64 : PTX_READ_SPECIAL_REGISTER<"nctaid", int_ptx_read_nctaid_r64>;
+def PTX_READ_NCTAID_X   : PTX_READ_SPECIAL_SUB_REGISTER<"nctaid.x", int_ptx_read_nctaid_x>;
+def PTX_READ_NCTAID_Y   : PTX_READ_SPECIAL_SUB_REGISTER<"nctaid.y", int_ptx_read_nctaid_y>;
+def PTX_READ_NCTAID_Z   : PTX_READ_SPECIAL_SUB_REGISTER<"nctaid.z", int_ptx_read_nctaid_z>;
+def PTX_READ_NCTAID_W   : PTX_READ_SPECIAL_SUB_REGISTER<"nctaid.w", int_ptx_read_nctaid_w>;
 
 // PTX Parallel Synchronization and Communication Intrinsics
 

Modified: llvm/trunk/test/CodeGen/PTX/intrinsic.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/intrinsic.ll?rev=127246&r1=127245&r2=127246&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/intrinsic.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/intrinsic.ll Tue Mar  8 08:10:18 2011
@@ -28,6 +28,90 @@
 	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.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)





More information about the llvm-commits mailing list