Jingyue Wu ef056a9111 [NVPTXFavorNonGenericAddrSpaces] recursively trace into GEP and BitCast ...

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

Test Plan: @ld_int_from_global_float in access-non-generic.ll

Reviewers: broune, eliben, jholewinski, meheff

Subscribers: jholewinski, llvm-commits

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

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@238574 91177308-0d34-0410-b5e6-96231b3b80d8

2015-05-29 17:00:27 +00:00

..

InstPrinter

Remove unused MCSubtargetInfo argument from the NVPTX MCInstPrinter ctors.

2015-03-30 22:03:16 +00:00

MCTargetDesc

MC: MCCodeGenInfo naming update. NFC.

2015-05-15 19:13:31 +00:00

TargetInfo

…

cl_common_defines.h

…

CMakeLists.txt

NVPTX: Remove dead code.

2015-03-02 13:16:28 +00:00

LLVMBuild.txt

Update libdeps in NVPTXCodeGen, since r225944.

2015-01-14 23:01:36 +00:00

Makefile

…

ManagedStringPool.h

…

NVPTX.h

[PM] Remove a bunch of stale TTI creation method declarations. I nuked

2015-02-01 00:22:15 +00:00

NVPTX.td

[NVPTX] Associate a minimum PTX version for each SM architecture

2015-03-30 19:30:55 +00:00

NVPTXAllocaHoisting.cpp

NVPTX: move NVPTXAllocaHoisting into the cpp file

2015-03-10 19:20:52 +00:00

NVPTXAllocaHoisting.h

NVPTX: move NVPTXAllocaHoisting into the cpp file

2015-03-10 19:20:52 +00:00

NVPTXAsmPrinter.cpp

Don't call utostr in Twine/raw_ostream contexts.

2015-05-28 11:24:24 +00:00

NVPTXAsmPrinter.h

Remove 3 includes from MCInstrDesc.h and explicitly include them where needed

2015-05-15 21:58:42 +00:00

NVPTXAssignValidGlobalNames.cpp

[PM] Remove the old 'PassManager.h' header file at the top level of

2015-02-13 10:01:29 +00:00

NVPTXFavorNonGenericAddrSpaces.cpp

[NVPTXFavorNonGenericAddrSpaces] recursively trace into GEP and BitCast

2015-05-29 17:00:27 +00:00

NVPTXFrameLowering.cpp

[ShrinkWrap] Add (a simplified version) of shrink-wrapping.

2015-05-05 17:38:16 +00:00

NVPTXFrameLowering.h

[ShrinkWrap] Add (a simplified version) of shrink-wrapping.

2015-05-05 17:38:16 +00:00

NVPTXGenericToNVVM.cpp

[opaque pointer type] More GEP IRBuilder API migrations...

2015-04-03 21:33:42 +00:00

NVPTXImageOptimizer.cpp

[cleanup] Re-sort all the #include lines in LLVM using

2015-01-14 11:23:27 +00:00

NVPTXInstrFormats.td

…

NVPTXInstrInfo.cpp

Remove all use of is64bit off of NVPTXSubtarget and clean up code

2015-02-19 00:08:27 +00:00

NVPTXInstrInfo.h

Remove all use of is64bit off of NVPTXSubtarget and clean up code

2015-02-19 00:08:27 +00:00

NVPTXInstrInfo.td

Reapply r235977 "[DebugInfo] Add debug locations to constant SD nodes"

2015-04-28 14:05:47 +00:00

NVPTXIntrinsics.td

…

NVPTXISelDAGToDAG.cpp

Reapply r235977 "[DebugInfo] Add debug locations to constant SD nodes"

2015-04-28 14:05:47 +00:00

NVPTXISelDAGToDAG.h

Reapply r235977 "[DebugInfo] Add debug locations to constant SD nodes"

2015-04-28 14:05:47 +00:00

NVPTXISelLowering.cpp

Move alignment from MCSectionData to MCSection.

2015-05-21 19:20:38 +00:00

NVPTXISelLowering.h

Change getTargetNodeName() to produce compiler warnings for missing cases, fix them

2015-05-07 21:33:59 +00:00

NVPTXLowerAggrCopies.cpp

Reduce dyn_cast<> to isa<> or cast<> where possible.

2015-04-10 11:24:51 +00:00

NVPTXLowerAggrCopies.h

Move unreferenced passes into the cpp file

2015-03-09 15:50:58 +00:00

NVPTXLowerStructArgs.cpp

Don't use LLVM_LIBRARY_VISIBILITY in cpp files.

2015-03-10 20:07:44 +00:00

NVPTXMachineFunctionInfo.h

…

NVPTXMCExpr.cpp

[NVPTX] Handle addrspacecast constant expressions in aggregate initializers

2015-04-28 17:18:30 +00:00

NVPTXMCExpr.h

Move alignment from MCSectionData to MCSection.

2015-05-21 19:20:38 +00:00

NVPTXPrologEpilogPass.cpp

[ShrinkWrap] Add (a simplified version) of shrink-wrapping.

2015-05-05 17:38:16 +00:00

NVPTXRegisterInfo.cpp

Have getCalleeSavedRegs take a non-null MachineFunction all the

2015-03-11 21:41:28 +00:00

NVPTXRegisterInfo.h

Have getCalleeSavedRegs take a non-null MachineFunction all the

2015-03-11 21:41:28 +00:00

NVPTXRegisterInfo.td

…

NVPTXReplaceImageHandles.cpp

Remove all use of getDrvInterface off of NVPTXSubtarget and clean

2015-02-19 00:08:23 +00:00

NVPTXSection.h

Implement unique sections with an unique ID.

2015-04-04 18:02:01 +00:00

NVPTXSubtarget.cpp

Remove all use of is64bit off of NVPTXSubtarget and clean up code

2015-02-19 00:08:27 +00:00

NVPTXSubtarget.h

Remove all use of is64bit off of NVPTXSubtarget and clean up code

2015-02-19 00:08:27 +00:00

NVPTXTargetMachine.cpp

[NaryReassociate] Run EarlyCSE after NaryReassociate

2015-05-28 04:56:52 +00:00

NVPTXTargetMachine.h

Use raw_pwrite_stream in the object writer/streamer.

2015-04-14 22:14:34 +00:00

NVPTXTargetObjectFile.h

Move alignment from MCSectionData to MCSection.

2015-05-21 19:20:38 +00:00

NVPTXTargetTransformInfo.cpp

Divergence analysis for GPU programs

2015-04-10 05:03:50 +00:00

NVPTXTargetTransformInfo.h

Divergence analysis for GPU programs

2015-04-10 05:03:50 +00:00

NVPTXUtilities.cpp

Simplify boolean expressions with true and false using clang-tidy

2015-03-23 16:26:23 +00:00

NVPTXUtilities.h

…

NVPTXVector.td

Reapply r235977 "[DebugInfo] Add debug locations to constant SD nodes"

2015-04-28 14:05:47 +00:00

NVVMReflect.cpp

Re-sort includes with sort-includes.py and insert raw_ostream.h where it's used.

2015-03-23 19:32:43 +00:00