diff --git a/lib/Target/PTX/PTXAsmPrinter.cpp b/lib/Target/PTX/PTXAsmPrinter.cpp
index 19a699e44fe..dd3e895eaa2 100644
--- 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
index 00000000000..dbbbb67a140
--- /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
+}