From 8902ecb6826d0677c32f1fb9aac41070fee1ac7d Mon Sep 17 00:00:00 2001
From: Che-Liang Chiou <clchiou@gmail.com>
Date: Fri, 18 Mar 2011 11:23:56 +0000
Subject: [PATCH] ptx: fix parameter order that is reversed
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@127874 91177308-0d34-0410-b5e6-96231b3b80d8
---
lib/Target/PTX/PTXAsmPrinter.cpp | 9 ++++-----
test/CodeGen/PTX/parameter-order.ll | 8 ++++++++
2 files changed, 12 insertions(+), 5 deletions(-)
create mode 100644 test/CodeGen/PTX/parameter-order.ll
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
+}