From 4c7ffb6a7e2881e8de85ad47c0d679074de420fa Mon Sep 17 00:00:00 2001 From: Justin Holewinski Date: Tue, 6 Dec 2011 17:39:48 +0000 Subject: [PATCH] PTX: Continue to fix up the register mess. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@145947 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp | 51 +++--- lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h | 59 +++++++ lib/Target/PTX/PTXAsmPrinter.cpp | 73 +++------ lib/Target/PTX/PTXISelLowering.cpp | 35 +++- lib/Target/PTX/PTXInstrInfo.td | 12 +- lib/Target/PTX/PTXMFInfoExtract.cpp | 15 +- lib/Target/PTX/PTXMachineFunctionInfo.h | 154 ++++++++++-------- test/CodeGen/PTX/mov.ll | 10 +- test/CodeGen/PTX/parameter-order.ll | 2 +- 9 files changed, 266 insertions(+), 145 deletions(-) diff --git a/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp b/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp index a2743487888..5fecb852481 100644 --- a/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp +++ b/lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp @@ -39,32 +39,45 @@ StringRef PTXInstPrinter::getOpcodeName(unsigned Opcode) const { void PTXInstPrinter::printRegName(raw_ostream &OS, unsigned RegNo) const { // Decode the register number into type and offset - unsigned RegType = RegNo & 0xF; - unsigned RegOffset = RegNo >> 4; + unsigned RegSpace = RegNo & 0x7; + unsigned RegType = (RegNo >> 3) & 0x7; + unsigned RegOffset = RegNo >> 6; // Print the register OS << "%"; - switch (RegType) { + switch (RegSpace) { default: - llvm_unreachable("Unknown register type!"); - case PTXRegisterType::Pred: - OS << "p"; + llvm_unreachable("Unknown register space!"); + case PTXRegisterSpace::Reg: + switch (RegType) { + default: + llvm_unreachable("Unknown register type!"); + case PTXRegisterType::Pred: + OS << "p"; + break; + case PTXRegisterType::B16: + OS << "rh"; + break; + case PTXRegisterType::B32: + OS << "r"; + break; + case PTXRegisterType::B64: + OS << "rd"; + break; + case PTXRegisterType::F32: + OS << "f"; + break; + case PTXRegisterType::F64: + OS << "fd"; + break; + } break; - case PTXRegisterType::B16: - OS << "rh"; + case PTXRegisterSpace::Return: + OS << "ret"; break; - case PTXRegisterType::B32: - OS << "r"; - break; - case PTXRegisterType::B64: - OS << "rd"; - break; - case PTXRegisterType::F32: - OS << "f"; - break; - case PTXRegisterType::F64: - OS << "fd"; + case PTXRegisterSpace::Argument: + OS << "arg"; break; } diff --git a/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h b/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h index d376d367ab6..77a298d2406 100644 --- a/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h +++ b/lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h @@ -17,6 +17,8 @@ #ifndef PTXBASEINFO_H #define PTXBASEINFO_H +#include "llvm/Support/ErrorHandling.h" +#include "llvm/Support/raw_ostream.h" #include "PTXMCTargetDesc.h" namespace llvm { @@ -69,6 +71,63 @@ namespace llvm { F64 }; } // namespace PTXRegisterType + + namespace PTXRegisterSpace { + // Register space encoded in MCOperands + enum { + Reg = 0, + Local, + Param, + Argument, + Return + }; + } + + inline static void decodeRegisterName(raw_ostream &OS, + unsigned EncodedReg) { + OS << "%"; + + unsigned RegSpace = EncodedReg & 0x7; + unsigned RegType = (EncodedReg >> 3) & 0x7; + unsigned RegOffset = EncodedReg >> 6; + + switch (RegSpace) { + default: + llvm_unreachable("Unknown register space!"); + case PTXRegisterSpace::Reg: + switch (RegType) { + default: + llvm_unreachable("Unknown register type!"); + case PTXRegisterType::Pred: + OS << "p"; + break; + case PTXRegisterType::B16: + OS << "rh"; + break; + case PTXRegisterType::B32: + OS << "r"; + break; + case PTXRegisterType::B64: + OS << "rd"; + break; + case PTXRegisterType::F32: + OS << "f"; + break; + case PTXRegisterType::F64: + OS << "fd"; + break; + } + break; + case PTXRegisterSpace::Return: + OS << "ret"; + break; + case PTXRegisterSpace::Argument: + OS << "arg"; + break; + } + + OS << RegOffset; + } } // namespace llvm #endif diff --git a/lib/Target/PTX/PTXAsmPrinter.cpp b/lib/Target/PTX/PTXAsmPrinter.cpp index db614ad36a3..77ed71d0a64 100644 --- a/lib/Target/PTX/PTXAsmPrinter.cpp +++ b/lib/Target/PTX/PTXAsmPrinter.cpp @@ -51,23 +51,23 @@ using namespace llvm; static const char PARAM_PREFIX[] = "__param_"; static const char RETURN_PREFIX[] = "__ret_"; -static const char *getRegisterTypeName(unsigned RegNo, - const MachineRegisterInfo& MRI) { - const TargetRegisterClass *TRC = MRI.getRegClass(RegNo); - -#define TEST_REGCLS(cls, clsstr) \ - if (PTX::cls ## RegisterClass == TRC) return # clsstr; - - TEST_REGCLS(RegPred, pred); - TEST_REGCLS(RegI16, b16); - TEST_REGCLS(RegI32, b32); - TEST_REGCLS(RegI64, b64); - TEST_REGCLS(RegF32, b32); - TEST_REGCLS(RegF64, b64); -#undef TEST_REGCLS - - llvm_unreachable("Not in any register class!"); - return NULL; +static const char *getRegisterTypeName(unsigned RegType) { + switch (RegType) { + default: + llvm_unreachable("Unknown register type"); + case PTXRegisterType::Pred: + return ".pred"; + case PTXRegisterType::B16: + return ".b16"; + case PTXRegisterType::B32: + return ".b32"; + case PTXRegisterType::B64: + return ".b64"; + case PTXRegisterType::F32: + return ".f32"; + case PTXRegisterType::F64: + return ".f64"; + } } static const char *getStateSpaceName(unsigned addressSpace) { @@ -188,32 +188,32 @@ void PTXAsmPrinter::EmitFunctionBodyStart() { unsigned numRegs; // pred - numRegs = MFI->getNumRegistersForClass(PTX::RegPredRegisterClass); + numRegs = MFI->countRegisters(PTXRegisterType::Pred, PTXRegisterSpace::Reg); if(numRegs > 0) os << "\t.reg .pred %p<" << numRegs << ">;\n"; // i16 - numRegs = MFI->getNumRegistersForClass(PTX::RegI16RegisterClass); + numRegs = MFI->countRegisters(PTXRegisterType::B16, PTXRegisterSpace::Reg); if(numRegs > 0) os << "\t.reg .b16 %rh<" << numRegs << ">;\n"; // i32 - numRegs = MFI->getNumRegistersForClass(PTX::RegI32RegisterClass); + numRegs = MFI->countRegisters(PTXRegisterType::B32, PTXRegisterSpace::Reg); if(numRegs > 0) os << "\t.reg .b32 %r<" << numRegs << ">;\n"; // i64 - numRegs = MFI->getNumRegistersForClass(PTX::RegI64RegisterClass); + numRegs = MFI->countRegisters(PTXRegisterType::B64, PTXRegisterSpace::Reg); if(numRegs > 0) os << "\t.reg .b64 %rd<" << numRegs << ">;\n"; // f32 - numRegs = MFI->getNumRegistersForClass(PTX::RegF32RegisterClass); + numRegs = MFI->countRegisters(PTXRegisterType::F32, PTXRegisterSpace::Reg); if(numRegs > 0) os << "\t.reg .f32 %f<" << numRegs << ">;\n"; // f64 - numRegs = MFI->getNumRegistersForClass(PTX::RegF64RegisterClass); + numRegs = MFI->countRegisters(PTXRegisterType::F64, PTXRegisterSpace::Reg); if(numRegs > 0) os << "\t.reg .f64 %fd<" << numRegs << ">;\n"; @@ -368,7 +368,6 @@ void PTXAsmPrinter::EmitFunctionEntryLabel() { const PTXParamManager &PM = MFI->getParamManager(); const bool isKernel = MFI->isKernel(); const PTXSubtarget& ST = TM.getSubtarget(); - const MachineRegisterInfo& MRI = MF->getRegInfo(); SmallString<128> decl; raw_svector_ostream os(decl); @@ -391,7 +390,7 @@ void PTXAsmPrinter::EmitFunctionEntryLabel() { if (i != b) os << ", "; - os << ".reg ." << getRegisterTypeName(*i, MRI) << ' ' + os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' ' << MFI->getRegisterName(*i); } } @@ -450,7 +449,7 @@ void PTXAsmPrinter::EmitFunctionEntryLabel() { if (i != b) os << ", "; - os << ".reg ." << getRegisterTypeName(*i, MRI) << ' ' + os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' ' << MFI->getRegisterName(*i); } } @@ -521,34 +520,14 @@ MCOperand PTXAsmPrinter::GetSymbolRef(const MachineOperand &MO, MCOperand PTXAsmPrinter::lowerOperand(const MachineOperand &MO) { MCOperand MCOp; const PTXMachineFunctionInfo *MFI = MF->getInfo(); - const MachineRegisterInfo& MRI = MF->getRegInfo(); - const TargetRegisterClass* TRC; - unsigned RegType; - unsigned RegOffset; unsigned EncodedReg; switch (MO.getType()) { default: llvm_unreachable("Unknown operand type"); case MachineOperand::MO_Register: if (MO.getReg() > 0) { - TRC = MRI.getRegClass(MO.getReg()); - // Determine which PTX register type to use - if (TRC == PTX::RegPredRegisterClass) - RegType = PTXRegisterType::Pred; - else if (TRC == PTX::RegI16RegisterClass) - RegType = PTXRegisterType::B16; - else if (TRC == PTX::RegI32RegisterClass) - RegType = PTXRegisterType::B32; - else if (TRC == PTX::RegI64RegisterClass) - RegType = PTXRegisterType::B64; - else if (TRC == PTX::RegF32RegisterClass) - RegType = PTXRegisterType::F32; - else if (TRC == PTX::RegF64RegisterClass) - RegType = PTXRegisterType::F64; - // Determine our virtual register offset - RegOffset = MFI->getOffsetForRegister(TRC, MO.getReg()); // Encode the register - EncodedReg = (RegOffset << 4) | RegType; + EncodedReg = MFI->getEncodedRegister(MO.getReg()); } else { EncodedReg = 0; } diff --git a/lib/Target/PTX/PTXISelLowering.cpp b/lib/Target/PTX/PTXISelLowering.cpp index 17191fbc3e0..a012297e835 100644 --- a/lib/Target/PTX/PTXISelLowering.cpp +++ b/lib/Target/PTX/PTXISelLowering.cpp @@ -243,6 +243,30 @@ SDValue PTXTargetLowering:: for (unsigned i = 0, e = Ins.size(); i != e; ++i) { EVT RegVT = Ins[i].VT; TargetRegisterClass* TRC = getRegClassFor(RegVT); + unsigned RegType; + + // Determine which register class we need + if (RegVT == MVT::i1) { + RegType = PTXRegisterType::Pred; + } + else if (RegVT == MVT::i16) { + RegType = PTXRegisterType::B16; + } + else if (RegVT == MVT::i32) { + RegType = PTXRegisterType::B32; + } + else if (RegVT == MVT::i64) { + RegType = PTXRegisterType::B64; + } + else if (RegVT == MVT::f32) { + RegType = PTXRegisterType::F32; + } + else if (RegVT == MVT::f64) { + RegType = PTXRegisterType::F64; + } + else { + llvm_unreachable("Unknown parameter type"); + } // Use a unique index in the instruction to prevent instruction folding. // Yes, this is a hack. @@ -253,7 +277,7 @@ SDValue PTXTargetLowering:: InVals.push_back(ArgValue); - MFI->addArgReg(Reg); + MFI->addRegister(Reg, RegType, PTXRegisterSpace::Argument); } } @@ -304,25 +328,32 @@ SDValue PTXTargetLowering:: for (unsigned i = 0, e = Outs.size(); i != e; ++i) { EVT RegVT = Outs[i].VT; TargetRegisterClass* TRC = 0; + unsigned RegType; // Determine which register class we need if (RegVT == MVT::i1) { TRC = PTX::RegPredRegisterClass; + RegType = PTXRegisterType::Pred; } else if (RegVT == MVT::i16) { TRC = PTX::RegI16RegisterClass; + RegType = PTXRegisterType::B16; } else if (RegVT == MVT::i32) { TRC = PTX::RegI32RegisterClass; + RegType = PTXRegisterType::B32; } else if (RegVT == MVT::i64) { TRC = PTX::RegI64RegisterClass; + RegType = PTXRegisterType::B64; } else if (RegVT == MVT::f32) { TRC = PTX::RegF32RegisterClass; + RegType = PTXRegisterType::F32; } else if (RegVT == MVT::f64) { TRC = PTX::RegF64RegisterClass; + RegType = PTXRegisterType::F64; } else { llvm_unreachable("Unknown parameter type"); @@ -335,7 +366,7 @@ SDValue PTXTargetLowering:: Chain = DAG.getNode(PTXISD::WRITE_PARAM, dl, MVT::Other, Copy, OutReg); - MFI->addRetReg(Reg); + MFI->addRegister(Reg, RegType, PTXRegisterSpace::Return); } } diff --git a/lib/Target/PTX/PTXInstrInfo.td b/lib/Target/PTX/PTXInstrInfo.td index bcd5bcf734e..19a862f4ec8 100644 --- a/lib/Target/PTX/PTXInstrInfo.td +++ b/lib/Target/PTX/PTXInstrInfo.td @@ -825,17 +825,17 @@ let hasSideEffects = 1 in { ///===- Parameter Passing Pseudo-Instructions -----------------------------===// def READPARAMPRED : InstPTX<(outs RegPred:$a), (ins i32imm:$b), - "mov.pred\t$a, %param$b", []>; + "mov.pred\t$a, %arg$b", []>; def READPARAMI16 : InstPTX<(outs RegI16:$a), (ins i32imm:$b), - "mov.b16\t$a, %param$b", []>; + "mov.b16\t$a, %arg$b", []>; def READPARAMI32 : InstPTX<(outs RegI32:$a), (ins i32imm:$b), - "mov.b32\t$a, %param$b", []>; + "mov.b32\t$a, %arg$b", []>; def READPARAMI64 : InstPTX<(outs RegI64:$a), (ins i32imm:$b), - "mov.b64\t$a, %param$b", []>; + "mov.b64\t$a, %arg$b", []>; def READPARAMF32 : InstPTX<(outs RegF32:$a), (ins i32imm:$b), - "mov.f32\t$a, %param$b", []>; + "mov.f32\t$a, %arg$b", []>; def READPARAMF64 : InstPTX<(outs RegF64:$a), (ins i32imm:$b), - "mov.f64\t$a, %param$b", []>; + "mov.f64\t$a, %arg$b", []>; def WRITEPARAMPRED : InstPTX<(outs), (ins RegPred:$a), "//w", []>; def WRITEPARAMI16 : InstPTX<(outs), (ins RegI16:$a), "//w", []>; diff --git a/lib/Target/PTX/PTXMFInfoExtract.cpp b/lib/Target/PTX/PTXMFInfoExtract.cpp index 81585928aea..26ec6239f49 100644 --- a/lib/Target/PTX/PTXMFInfoExtract.cpp +++ b/lib/Target/PTX/PTXMFInfoExtract.cpp @@ -58,7 +58,20 @@ bool PTXMFInfoExtract::runOnMachineFunction(MachineFunction &MF) { for (unsigned i = 0; i < MRI.getNumVirtRegs(); ++i) { unsigned Reg = TargetRegisterInfo::index2VirtReg(i); const TargetRegisterClass *TRC = MRI.getRegClass(Reg); - MFI->addVirtualRegister(TRC, Reg); + unsigned RegType; + if (TRC == PTX::RegPredRegisterClass) + RegType = PTXRegisterType::Pred; + else if (TRC == PTX::RegI16RegisterClass) + RegType = PTXRegisterType::B16; + else if (TRC == PTX::RegI32RegisterClass) + RegType = PTXRegisterType::B32; + else if (TRC == PTX::RegI64RegisterClass) + RegType = PTXRegisterType::B64; + else if (TRC == PTX::RegF32RegisterClass) + RegType = PTXRegisterType::F32; + else if (TRC == PTX::RegF64RegisterClass) + RegType = PTXRegisterType::F64; + MFI->addRegister(Reg, RegType, PTXRegisterSpace::Reg); } return false; diff --git a/lib/Target/PTX/PTXMachineFunctionInfo.h b/lib/Target/PTX/PTXMachineFunctionInfo.h index 6f3b826db96..1a2878cbec3 100644 --- a/lib/Target/PTX/PTXMachineFunctionInfo.h +++ b/lib/Target/PTX/PTXMachineFunctionInfo.h @@ -35,15 +35,22 @@ private: DenseSet RegArgs; DenseSet RegRets; - typedef std::vector RegisterList; - typedef DenseMap RegisterMap; - typedef DenseMap RegisterNameMap; typedef DenseMap FrameMap; - RegisterMap UsedRegs; - RegisterNameMap RegNames; FrameMap FrameSymbols; + struct RegisterInfo { + unsigned Reg; + unsigned Type; + unsigned Space; + unsigned Offset; + unsigned Encoded; + }; + + typedef DenseMap RegisterInfoMap; + + RegisterInfoMap RegInfo; + PTXParamManager ParamManager; public: @@ -51,13 +58,7 @@ public: PTXMachineFunctionInfo(MachineFunction &MF) : IsKernel(false) { - UsedRegs[PTX::RegPredRegisterClass] = RegisterList(); - UsedRegs[PTX::RegI16RegisterClass] = RegisterList(); - UsedRegs[PTX::RegI32RegisterClass] = RegisterList(); - UsedRegs[PTX::RegI64RegisterClass] = RegisterList(); - UsedRegs[PTX::RegF32RegisterClass] = RegisterList(); - UsedRegs[PTX::RegF64RegisterClass] = RegisterList(); - } + } /// getParamManager - Returns the PTXParamManager instance for this function. PTXParamManager& getParamManager() { return ParamManager; } @@ -78,81 +79,106 @@ public: reg_iterator retreg_begin() const { return RegRets.begin(); } reg_iterator retreg_end() const { return RegRets.end(); } + /// addRegister - Adds a virtual register to the set of all used registers + void addRegister(unsigned Reg, unsigned RegType, unsigned RegSpace) { + if (!RegInfo.count(Reg)) { + RegisterInfo Info; + Info.Reg = Reg; + Info.Type = RegType; + Info.Space = RegSpace; + + // Determine register offset + Info.Offset = 0; + for(RegisterInfoMap::const_iterator i = RegInfo.begin(), + e = RegInfo.end(); i != e; ++i) { + const RegisterInfo& RI = i->second; + if (RI.Space == RegSpace) + if (RI.Space != PTXRegisterSpace::Reg || RI.Type == Info.Type) + Info.Offset++; + } + + // Encode the register data into a single register number + Info.Encoded = (Info.Offset << 6) | (Info.Type << 3) | Info.Space; + + RegInfo[Reg] = Info; + + if (RegSpace == PTXRegisterSpace::Argument) + RegArgs.insert(Reg); + else if (RegSpace == PTXRegisterSpace::Return) + RegRets.insert(Reg); + } + } + + /// countRegisters - Returns the number of registers of the given type and + /// space. + unsigned countRegisters(unsigned RegType, unsigned RegSpace) const { + unsigned Count = 0; + for(RegisterInfoMap::const_iterator i = RegInfo.begin(), e = RegInfo.end(); + i != e; ++i) { + const RegisterInfo& RI = i->second; + if (RI.Type == RegType && RI.Space == RegSpace) + Count++; + } + return Count; + } + + /// getEncodedRegister - Returns the encoded value of the register. + unsigned getEncodedRegister(unsigned Reg) const { + return RegInfo.lookup(Reg).Encoded; + } + /// addRetReg - Adds a register to the set of return-value registers. void addRetReg(unsigned Reg) { if (!RegRets.count(Reg)) { RegRets.insert(Reg); - std::string name; - name = "%ret"; - name += utostr(RegRets.size() - 1); - RegNames[Reg] = name; } } /// addArgReg - Adds a register to the set of function argument registers. void addArgReg(unsigned Reg) { RegArgs.insert(Reg); - std::string name; - name = "%param"; - name += utostr(RegArgs.size() - 1); - RegNames[Reg] = name; - } - - /// addVirtualRegister - Adds a virtual register to the set of all used - /// registers in the function. - void addVirtualRegister(const TargetRegisterClass *TRC, unsigned Reg) { - std::string name; - - // Do not count registers that are argument/return registers. - if (!RegRets.count(Reg) && !RegArgs.count(Reg)) { - UsedRegs[TRC].push_back(Reg); - if (TRC == PTX::RegPredRegisterClass) - name = "%p"; - else if (TRC == PTX::RegI16RegisterClass) - name = "%rh"; - else if (TRC == PTX::RegI32RegisterClass) - name = "%r"; - else if (TRC == PTX::RegI64RegisterClass) - name = "%rd"; - else if (TRC == PTX::RegF32RegisterClass) - name = "%f"; - else if (TRC == PTX::RegF64RegisterClass) - name = "%fd"; - else - llvm_unreachable("Invalid register class"); - - name += utostr(UsedRegs[TRC].size() - 1); - RegNames[Reg] = name; - } } /// getRegisterName - Returns the name of the specified virtual register. This /// name is used during PTX emission. - const char *getRegisterName(unsigned Reg) const { - if (RegNames.count(Reg)) - return RegNames.find(Reg)->second.c_str(); + std::string getRegisterName(unsigned Reg) const { + if (RegInfo.count(Reg)) { + const RegisterInfo& RI = RegInfo.lookup(Reg); + std::string Name; + raw_string_ostream NameStr(Name); + decodeRegisterName(NameStr, RI.Encoded); + NameStr.flush(); + return Name; + } else if (Reg == PTX::NoRegister) return "%noreg"; else llvm_unreachable("Register not in register name map"); } - /// getNumRegistersForClass - Returns the number of virtual registers that are - /// used for the specified register class. - unsigned getNumRegistersForClass(const TargetRegisterClass *TRC) const { - return UsedRegs.lookup(TRC).size(); + /// getEncodedRegisterName - Returns the name of the encoded register. + std::string getEncodedRegisterName(unsigned EncodedReg) const { + std::string Name; + raw_string_ostream NameStr(Name); + decodeRegisterName(NameStr, EncodedReg); + NameStr.flush(); + return Name; + } + + /// getRegisterType - Returns the type of the specified virtual register. + unsigned getRegisterType(unsigned Reg) const { + if (RegInfo.count(Reg)) + return RegInfo.lookup(Reg).Type; + else + llvm_unreachable("Unknown register"); } /// getOffsetForRegister - Returns the offset of the virtual register - unsigned getOffsetForRegister(const TargetRegisterClass *TRC, - unsigned Reg) const { - const RegisterList &RegList = UsedRegs.lookup(TRC); - for (unsigned i = 0, e = RegList.size(); i != e; ++i) { - if (RegList[i] == Reg) - return i; - } - //llvm_unreachable("Unknown virtual register"); - return 0; + unsigned getOffsetForRegister(unsigned Reg) const { + if (RegInfo.count(Reg)) + return RegInfo.lookup(Reg).Offset; + else + return 0; } /// getFrameSymbol - Returns the symbol name for the given FrameIndex. diff --git a/test/CodeGen/PTX/mov.ll b/test/CodeGen/PTX/mov.ll index 75555a7c5e9..9e501be03ee 100644 --- a/test/CodeGen/PTX/mov.ll +++ b/test/CodeGen/PTX/mov.ll @@ -31,31 +31,31 @@ define ptx_device double @t1_f64() { } define ptx_device i16 @t2_u16(i16 %x) { -; CHECK: mov.b16 %ret{{[0-9]+}}, %param{{[0-9]+}}; +; CHECK: mov.b16 %ret{{[0-9]+}}, %arg{{[0-9]+}}; ; CHECK: ret; ret i16 %x } define ptx_device i32 @t2_u32(i32 %x) { -; CHECK: mov.b32 %ret{{[0-9]+}}, %param{{[0-9]+}}; +; CHECK: mov.b32 %ret{{[0-9]+}}, %arg{{[0-9]+}}; ; CHECK: ret; ret i32 %x } define ptx_device i64 @t2_u64(i64 %x) { -; CHECK: mov.b64 %ret{{[0-9]+}}, %param{{[0-9]+}}; +; CHECK: mov.b64 %ret{{[0-9]+}}, %arg{{[0-9]+}}; ; CHECK: ret; ret i64 %x } define ptx_device float @t3_f32(float %x) { -; CHECK: mov.f32 %ret{{[0-9]+}}, %param{{[0-9]+}}; +; CHECK: mov.f32 %ret{{[0-9]+}}, %arg{{[0-9]+}}; ; CHECK: ret; ret float %x } define ptx_device double @t3_f64(double %x) { -; CHECK: mov.f64 %ret{{[0-9]+}}, %param{{[0-9]+}}; +; CHECK: mov.f64 %ret{{[0-9]+}}, %arg{{[0-9]+}}; ; CHECK: ret; ret double %x } diff --git a/test/CodeGen/PTX/parameter-order.ll b/test/CodeGen/PTX/parameter-order.ll index 09015dae1db..377f17379fe 100644 --- a/test/CodeGen/PTX/parameter-order.ll +++ b/test/CodeGen/PTX/parameter-order.ll @@ -1,6 +1,6 @@ ; RUN: llc < %s -march=ptx32 | FileCheck %s -; CHECK: .func (.reg .b32 %ret{{[0-9]+}}) test_parameter_order (.reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}}) +; CHECK: .func (.reg .b32 %ret{{[0-9]+}}) test_parameter_order (.reg .f32 %arg{{[0-9]+}}, .reg .b32 %arg{{[0-9]+}}, .reg .b32 %arg{{[0-9]+}}, .reg .f32 %arg{{[0-9]+}}) define ptx_device i32 @test_parameter_order(float %a, i32 %b, i32 %c, float %d) { ; CHECK: sub.u32 %ret{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}} %result = sub i32 %b, %c