Page MenuHomePhabricator

[NVPTXFavorNonGenericAddrSpaces] recursively trace into GEP and BitCast
ClosedPublic

Authored by jingyue on May 27 2015, 3:01 PM.

Details

Summary

This patch allows NVPTXFavorNonGenericAddrSpaces to remove addrspacecast
from longer chains consisting of GEPs and BitCasts. For example, it can
now optimize

%0 = addrspacecast [10 x float] addrspace(3)* @a to [10 x float]*
%1 = gep [10 x float]* %0, i64 0, i64 %i
%2 = bitcast float* %1 to i32*
%3 = load i32* %2 ; emits ld.u32

to

%0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
%1 = bitcast float addrspace(3)* %0 to i32 addrspace(3)*
%3 = load i32 addrspace(3)* %1 ; emits ld.shared.f32

Diff Detail

Repository
rL LLVM

Event Timeline

jingyue updated this revision to Diff 26638.May 27 2015, 3:01 PM
jingyue retitled this revision from to [NVPTXFavorNonGenericAddrSpaces] recursively trace into GEP and BitCast.
jingyue updated this object.
jingyue edited the test plan for this revision. (Show Details)
jingyue added reviewers: broune, eliben, meheff, jholewinski.
jingyue added a subscriber: Unknown Object (MLST).
meheff edited edge metadata.May 28 2015, 9:50 AM

LGTM

lib/Target/NVPTX/NVPTXFavorNonGenericAddrSpaces.cpp
32 ↗(On Diff #26638)

still emits ld.u32?

187 ↗(On Diff #26638)

BC's source operand could have started out as an addrspace cast, right? That is, it wasn't necessarily replaced.

222 ↗(On Diff #26638)

This doesn't need to be static.

jingyue updated this revision to Diff 26789.May 29 2015, 9:57 AM
jingyue edited edge metadata.

addressed Mark's comments

This revision was automatically updated to reflect the committed changes.
jingyue added inline comments.May 29 2015, 10:30 AM
lib/Target/NVPTX/NVPTXFavorNonGenericAddrSpaces.cpp
32 ↗(On Diff #26638)

Done. Carelessly copied the old comments :)

187 ↗(On Diff #26638)

Right. Updated this comment and another one below.

222 ↗(On Diff #26638)

Done.