diff options
| author | Che-Liang Chiou <clchiou@gmail.com> | 2011-03-18 11:23:56 +0000 |
|---|---|---|
| committer | Che-Liang Chiou <clchiou@gmail.com> | 2011-03-18 11:23:56 +0000 |
| commit | 8902ecb6826d0677c32f1fb9aac41070fee1ac7d (patch) | |
| tree | a090b42e2664203411cd7769f02b7639beb93bef | |
| parent | 88d3367baa066b4924a9303291aee084c154fff1 (diff) | |
| download | external_llvm-8902ecb6826d0677c32f1fb9aac41070fee1ac7d.zip external_llvm-8902ecb6826d0677c32f1fb9aac41070fee1ac7d.tar.gz external_llvm-8902ecb6826d0677c32f1fb9aac41070fee1ac7d.tar.bz2 | |
ptx: fix parameter order that is reversed
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@127874 91177308-0d34-0410-b5e6-96231b3b80d8
| -rw-r--r-- | lib/Target/PTX/PTXAsmPrinter.cpp | 9 | ||||
| -rw-r--r-- | test/CodeGen/PTX/parameter-order.ll | 8 |
2 files changed, 12 insertions, 5 deletions
diff --git a/lib/Target/PTX/PTXAsmPrinter.cpp b/lib/Target/PTX/PTXAsmPrinter.cpp index 19a699e..dd3e895 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 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 +} |