6502/llvm-6502
1
0
Fork 0
mirror of https://github.com/c64scene-ar/llvm-6502.git synced 2025-02-20 14:29:27 +00:00
Code Issues Projects Releases Wiki Activity

[NVPTX] run SROA after NVPTXFavorNonGenericAddrSpaces

Browse Source 
Summary:
This cleans up most allocas NVPTXLowerKernelArgs emits for byval
parameters.

Test Plan: makes bug21465.ll more stronger to verify no redundant local load/store.

Reviewers: eliben, jholewinski

Reviewed By: eliben, jholewinski

Subscribers: jholewinski, llvm-commits

Differential Revision: http://reviews.llvm.org/D10322

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@239368 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Jingyue Wu 2015-06-09 00:05:56 +00:00
parent 04ed4447ee
commit 4e04297ac3
2 changed files with 8 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
lib/Target/NVPTX/NVPTXTargetMachine.cpp
View File

@ -165,6 +165,11 @@ void NVPTXPassConfig::addIRPasses() {
addPass(createGenericToNVVMPass());
addPass(createNVPTXLowerKernelArgsPass(&getNVPTXTargetMachine()));
addPass(createNVPTXFavorNonGenericAddrSpacesPass());
// NVPTXLowerKernelArgs emits alloca for byval parameters which can often
// be eliminated by SROA. We do not run SROA right after NVPTXLowerKernelArgs
// because we plan to merge NVPTXLowerKernelArgs and
// NVPTXFavorNonGenericAddrSpaces into one pass.
addPass(createSROAPass());
// FavorNonGenericAddrSpaces shortcuts unnecessary addrspacecasts, and leave
// them unused. We could remove dead code in an ad-hoc manner, but that
// requires manual work and might be error-prone.

5
test/CodeGen/NVPTX/bug21465.ll
View File

@ -14,9 +14,10 @@ entry:
; CHECK: addrspacecast %struct.S* %input to %struct.S addrspace(101)*
%b = getelementptr inbounds %struct.S, %struct.S* %input, i64 0, i32 1
%0 = load i32, i32* %b, align 4
; PTX: ld.param.u32 %r{{[0-9]+}}, {{\[}}[[BASE:%rd[0-9]+]]{{\]}}
; PTX-NEXT: ld.param.u32 %r{{[0-9]+}}, {{\[}}[[BASE]]+4{{\]}}
; PTX-NOT: ld.param.u32 {{%r[0-9]+}}, [{{%rd[0-9]+}}]
; PTX: ld.param.u32 [[value:%r[0-9]+]], [{{%rd[0-9]+}}+4]
store i32 %0, i32* %output, align 4
; PTX-NEXT: st.global.u32 [{{%rd[0-9]+}}], [[value]]
ret void
}