ptx: remove reg-reg addressing mode and st.const

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@122653 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Target/PTX/PTXISelDAGToDAG.cpp b/lib/Target/PTX/PTXISelDAGToDAG.cpp
index e09182de3737f7dd0f4b24ed8dbb695f64c00fb6..294b62f08f3770ec8b11790563ba2cd450d6bc18 100644 (file)
--- a/lib/Target/PTX/PTXISelDAGToDAG.cpp
+++ b/lib/Target/PTX/PTXISelDAGToDAG.cpp
@@ -67,8 +67,8 @@ bool PTXDAGToDAGISel::SelectADDRrr(SDValue &Addr, SDValue &R1, SDValue &R2) {
       isImm(Addr.getOperand(0)) || isImm(Addr.getOperand(1)))
     return false;
 
-  R1 = Addr.getOperand(0);
-  R2 = Addr.getOperand(1);
+  R1 = Addr;
+  R2 = CurDAG->getTargetConstant(0, MVT::i32);
   return true;
 }
 
@@ -76,17 +76,20 @@ bool PTXDAGToDAGISel::SelectADDRrr(SDValue &Addr, SDValue &R1, SDValue &R2) {
 bool PTXDAGToDAGISel::SelectADDRri(SDValue &Addr, SDValue &Base,
                                    SDValue &Offset) {
   if (Addr.getOpcode() != ISD::ADD) {
+    // let SelectADDRii handle the [imm] case
     if (isImm(Addr))
       return false;
-    // is [reg] but not [imm]
+    // it is [reg]
     Base = Addr;
     Offset = CurDAG->getTargetConstant(0, MVT::i32);
     return true;
   }
 
+  if (Addr.getNumOperands() < 2)
+    return false;
+
   // let SelectADDRii handle the [imm+imm] case
-  if (Addr.getNumOperands() >= 2 &&
-      isImm(Addr.getOperand(0)) && isImm(Addr.getOperand(1)))
+  if (isImm(Addr.getOperand(0)) && isImm(Addr.getOperand(1)))
     return false;
 
   // try [reg+imm] and [imm+reg]
@@ -96,13 +99,7 @@ bool PTXDAGToDAGISel::SelectADDRri(SDValue &Addr, SDValue &Base,
       return true;
     }
 
-  // either [reg+imm] and [imm+reg]
-  for (int i = 0; i < 2; i ++)
-    if (SelectImm(Addr.getOperand(1-i), Offset)) {
-      Base = Addr.getOperand(i);
-      return true;
-    }
-
+  // neither [reg+imm] nor [imm+reg]
   return false;
 }
 

diff --git a/lib/Target/PTX/PTXInstrInfo.td b/lib/Target/PTX/PTXInstrInfo.td
index 13b1d77f39721d5c0c433ec21e1401bf44be1230..8e4b720308458ebfd58665d7d1fdb1c225685ce4 100644 (file)
--- a/lib/Target/PTX/PTXInstrInfo.td
+++ b/lib/Target/PTX/PTXInstrInfo.td
@@ -76,16 +76,6 @@ def store_global
   return false;
 }]>;
 
-def store_constant
-  : PatFrag<(ops node:$d, node:$ptr), (store node:$d, node:$ptr), [{
-  const Value *Src;
-  const PointerType *PT;
-  if ((Src = cast<StoreSDNode>(N)->getSrcValue()) &&
-      (PT = dyn_cast<PointerType>(Src->getType())))
-    return PT->getAddressSpace() == PTX::CONSTANT;
-  return false;
-}]>;
-
 def store_local
   : PatFrag<(ops node:$d, node:$ptr), (store node:$d, node:$ptr), [{
   const Value *Src;
@@ -122,10 +112,6 @@ def ADDRri : ComplexPattern<i32, 2, "SelectADDRri", [], []>;
 def ADDRii : ComplexPattern<i32, 2, "SelectADDRii", [], []>;
 
 // Address operands
-def MEMrr : Operand<i32> {
-  let PrintMethod = "printMemOperand";
-  let MIOperandInfo = (ops RRegs32, RRegs32);
-}
 def MEMri : Operand<i32> {
   let PrintMethod = "printMemOperand";
   let MIOperandInfo = (ops RRegs32, i32imm);
@@ -182,7 +168,7 @@ multiclass INT3ntnc<string opcstr, SDNode opnode> {
 
 multiclass PTX_LD<string opstr, RegisterClass RC, PatFrag pat_load> {
   def rr : InstPTX<(outs RC:$d),
-                   (ins MEMrr:$a),
+                   (ins MEMri:$a),
                    !strconcat(opstr, ".%type\t$d, [$a]"),
                    [(set RC:$d, (pat_load ADDRrr:$a))]>;
   def ri : InstPTX<(outs RC:$d),
@@ -197,7 +183,7 @@ multiclass PTX_LD<string opstr, RegisterClass RC, PatFrag pat_load> {
 
 multiclass PTX_ST<string opstr, RegisterClass RC, PatFrag pat_store> {
   def rr : InstPTX<(outs),
-                   (ins RC:$d, MEMrr:$a),
+                   (ins RC:$d, MEMri:$a),
                    !strconcat(opstr, ".%type\t[$a], $d"),
                    [(pat_store RC:$d, ADDRrr:$a)]>;
   def ri : InstPTX<(outs),
@@ -251,7 +237,6 @@ defm LDp : PTX_LD<"ld.param",  RRegs32, load_parameter>;
 defm LDs : PTX_LD<"ld.shared", RRegs32, load_shared>;
 
 defm STg : PTX_ST<"st.global", RRegs32, store_global>;
-defm STc : PTX_ST<"st.const",  RRegs32, store_constant>;
 defm STl : PTX_ST<"st.local",  RRegs32, store_local>;
 defm STp : PTX_ST<"st.param",  RRegs32, store_parameter>;
 defm STs : PTX_ST<"st.shared", RRegs32, store_shared>;

diff --git a/test/CodeGen/PTX/ld.ll b/test/CodeGen/PTX/ld.ll
index baafbc2d3d220b263a2e092472771abe01ab2547..836c4d41045a9092c8c852b550efd4f5aad4ce57 100644 (file)
--- a/test/CodeGen/PTX/ld.ll
+++ b/test/CodeGen/PTX/ld.ll
@@ -30,7 +30,8 @@ entry:
 define ptx_device i32 @t3(i32* %p, i32 %q) {
 entry:
 ;CHECK: shl.b32 r0, r2, 2;
-;CHECK: ld.global.s32 r0, [r1+r0];
+;CHECK: add.s32 r0, r1, r0;
+;CHECK: ld.global.s32 r0, [r0];
   %i = getelementptr i32* %p, i32 %q
   %x = load i32* %i
   ret i32 %x

diff --git a/test/CodeGen/PTX/st.ll b/test/CodeGen/PTX/st.ll
index ed482b24a402ef20523acfef57f37ef83b83fa26..2cbacb9ee59cee3fb522e43d79aea84a6fcad30f 100644 (file)
--- a/test/CodeGen/PTX/st.ll
+++ b/test/CodeGen/PTX/st.ll
@@ -31,7 +31,8 @@ define ptx_device void @t3(i32* %p, i32 %q, i32 %x) {
 ;CHECK: .reg .s32 r0;
 entry:
 ;CHECK: shl.b32 r0, r2, 2;
-;CHECK: st.global.s32 [r1+r0], r3;
+;CHECK: add.s32 r0, r1, r0;
+;CHECK: st.global.s32 [r0], r3;
   %i = getelementptr i32* %p, i32 %q
   store i32 %x, i32* %i
   ret void
@@ -45,14 +46,6 @@ entry:
   ret void
 }
 
-define ptx_device void @t4_const(i32 %x) {
-entry:
-;CHECK: st.const.s32 [array_constant], r1;
-  %i = getelementptr [10 x i32] addrspace(1)* @array_constant, i32 0, i32 0
-  store i32 %x, i32 addrspace(1)* %i
-  ret void
-}
-
 define ptx_device void @t4_local(i32 %x) {
 entry:
 ;CHECK: st.local.s32 [array_local], r1;