6502/llvm-6502
1
0
Fork 0
mirror of https://github.com/c64scene-ar/llvm-6502.git synced 2025-11-02 22:23:10 +00:00
Code Issues Projects Releases Wiki Activity

ptx: fix parameter order that is reversed

Browse Source 
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@127874 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Che-Liang Chiou
2011-03-18 11:23:56 +00:00
parent 88d3367baa
commit 8902ecb682
2 changed files with 12 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
lib/Target/PTX/PTXAsmPrinter.cpp
View File

@@ -381,9 +381,8 @@ void PTXAsmPrinter::EmitFunctionDeclaration() {
decl += " ("; decl += " (";
if (isKernel) { if (isKernel) {
unsigned cnt = 0; unsigned cnt = 0;
//for (int i = 0, e = MFI->getNumArg(); i != e; ++i) { for(PTXMachineFunctionInfo::reg_iterator
for(PTXMachineFunctionInfo::reg_reverse_iterator i = MFI->argRegBegin(), e = MFI->argRegEnd(), b = i;
i = MFI->argRegReverseBegin(), e = MFI->argRegReverseEnd(), b = i;
i != e; ++i) { i != e; ++i) {
reg = *i; reg = *i;
assert(reg != PTX::NoRegister && "Not a valid register!"); assert(reg != PTX::NoRegister && "Not a valid register!");
@@ -396,8 +395,8 @@ void PTXAsmPrinter::EmitFunctionDeclaration() {
decl += utostr(++cnt); decl += utostr(++cnt);
} }
} else { } else {
for (PTXMachineFunctionInfo::reg_reverse_iterator for (PTXMachineFunctionInfo::reg_iterator
i = MFI->argRegReverseBegin(), e = MFI->argRegReverseEnd(), b = i; i = MFI->argRegBegin(), e = MFI->argRegEnd(), b = i;
i != e; ++i) { i != e; ++i) {
reg = *i; reg = *i;
assert(reg != PTX::NoRegister && "Not a valid register!"); assert(reg != PTX::NoRegister && "Not a valid register!");

8
test/CodeGen/PTX/parameter-order.ll Normal file
View File

@@ -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
}