Page MenuHomePhabricator

[NVPTXFavorNonGenericAddrSpaces] recursively trace into GEP and BitCast

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



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


%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


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


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
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)