RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-24 03:33:20 +01:00
Code Issues Projects Releases Wiki Activity

ptx: fix parameter order that is reversed

Browse Source 
llvm-svn: 127874
This commit is contained in:
Che-Liang Chiou 2011-03-18 11:23:56 +00:00
parent f4a2c17cf5
commit 2b173c0443
2 changed files with 12 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -381,9 +381,8 @@ void PTXAsmPrinter::EmitFunctionDeclaration() {
decl += " (";
if (isKernel) {
unsigned cnt = 0;
//for (int i = 0, e = MFI->getNumArg(); i != e; ++i) {
for(PTXMachineFunctionInfo::reg_reverse_iterator
i = MFI->argRegReverseBegin(), e = MFI->argRegReverseEnd(), b = i;
for(PTXMachineFunctionInfo::reg_iterator
i = MFI->argRegBegin(), e = MFI->argRegEnd(), b = i;
i != e; ++i) {
reg = *i;
assert(reg != PTX::NoRegister && "Not a valid register!");
@ -396,8 +395,8 @@ void PTXAsmPrinter::EmitFunctionDeclaration() {
decl += utostr(++cnt);
}
} else {
for (PTXMachineFunctionInfo::reg_reverse_iterator
i = MFI->argRegReverseBegin(), e = MFI->argRegReverseEnd(), b = i;
for (PTXMachineFunctionInfo::reg_iterator
i = MFI->argRegBegin(), e = MFI->argRegEnd(), b = i;
i != e; ++i) {
reg = *i;
assert(reg != PTX::NoRegister && "Not a valid register!");

8
test/CodeGen/PTX/parameter-order.ll Normal file
View File

@ -0,0 +1,8 @@
; RUN: llc < %s -march=ptx | FileCheck %s
; CHECK: .func (.reg .u32 r0) test_parameter_order (.reg .u32 r1, .reg .u32 r2)
define ptx_device i32 @test_parameter_order(i32 %x, i32 %y) {
; CHECK: sub.u32 r0, r1, r2
%z = sub i32 %x, %y
ret i32 %z
}