RPCS3/llvm
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm.git synced 2024-11-25 21:00:00 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

PTX: Continue to fix up the register mess.

Browse Source 
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@145947 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Justin Holewinski 2011-12-06 17:39:48 +00:00
parent e37a83f66b
commit 4c7ffb6a7e
9 changed files with 266 additions and 145 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

51
lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp
View File

@ -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;
}

59
lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h
View File

@ -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

73
lib/Target/PTX/PTXAsmPrinter.cpp
View File

@ -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<PTXSubtarget>();
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<PTXMachineFunctionInfo>();
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;
}

35
lib/Target/PTX/PTXISelLowering.cpp
View File

@ -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);
}
}

12
lib/Target/PTX/PTXInstrInfo.td
View File

@ -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", []>;

15
lib/Target/PTX/PTXMFInfoExtract.cpp
View File

@ -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;

154
lib/Target/PTX/PTXMachineFunctionInfo.h
View File

@ -35,15 +35,22 @@ private:
DenseSet<unsigned> RegArgs;
DenseSet<unsigned> RegRets;
typedef std::vector<unsigned> RegisterList;
typedef DenseMap<const TargetRegisterClass*, RegisterList> RegisterMap;
typedef DenseMap<unsigned, std::string> RegisterNameMap;
typedef DenseMap<int, std::string> FrameMap;
RegisterMap UsedRegs;
RegisterNameMap RegNames;
FrameMap FrameSymbols;
struct RegisterInfo {
unsigned Reg;
unsigned Type;
unsigned Space;
unsigned Offset;
unsigned Encoded;
};
typedef DenseMap<unsigned, RegisterInfo> 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.

10
test/CodeGen/PTX/mov.ll
View File

@ -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
}

2
test/CodeGen/PTX/parameter-order.ll
View File

@ -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