[llvm-commits] [llvm] r127874 - in /llvm/trunk: lib/Target/PTX/PTXAsmPrinter.cpp test/CodeGen/PTX/parameter-order.ll

[Che-Liang Chiou]

Author: clchiou
Date: Fri Mar 18 06:23:56 2011
New Revision: 127874

URL: http://llvm.org/viewvc/llvm-project?rev=127874&view=rev
Log:
ptx: fix parameter order that is reversed

Added:
    llvm/trunk/test/CodeGen/PTX/parameter-order.ll
Modified:
    llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp

Modified: llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp?rev=127874&r1=127873&r2=127874&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp Fri Mar 18 06:23:56 2011
@@ -381,9 +381,8 @@
     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 @@
         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!");

Added: llvm/trunk/test/CodeGen/PTX/parameter-order.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/parameter-order.ll?rev=127874&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/parameter-order.ll (added)
+++ llvm/trunk/test/CodeGen/PTX/parameter-order.ll Fri Mar 18 06:23:56 2011
@@ -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
+}