llvm-6502/lib/Target/NVPTX
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