ptx: fix parameter order that is reversed

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

diff --git a/lib/Target/PTX/PTXAsmPrinter.cpp b/lib/Target/PTX/PTXAsmPrinter.cpp
index 19a699e44fe300b857a2f36f918457d25a66b8c6..dd3e895eaa20cf94aec4a60cf2c8f0d2e25ff950 100644 (file)
--- a/lib/Target/PTX/PTXAsmPrinter.cpp
+++ b/lib/Target/PTX/PTXAsmPrinter.cpp
@@ -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!");

diff --git a/test/CodeGen/PTX/parameter-order.ll b/test/CodeGen/PTX/parameter-order.ll
new file mode 100644 (file)
index 0000000..dbbbb67
--- /dev/null
+++ b/test/CodeGen/PTX/parameter-order.ll
@@ -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
+}