Jingyue Wu c9f86c1260 [NVPTX] make load on global readonly memory to use ldg ... Summary: [NVPTX] make load on global readonly memory to use ldg Summary: As describe in [1], ld.global.nc may be used to load memory by nvcc when __restrict__ is used and compiler can detect whether read-only data cache is safe to use. This patch will try to check whether ldg is safe to use and use them to replace ld.global when possible. This change can improve the performance by 18~29% on affected kernels (ratt*_kernel and rwdot*_kernel) in S3D benchmark of shoc [2]. Patched by Xuetian Weng. [1] http://docs.nvidia.com/cuda/kepler-tuning-guide/#read-only-data-cache [2] https://github.com/vetter/shoc Test Plan: test/CodeGen/NVPTX/load-with-non-coherent-cache.ll Reviewers: jholewinski, jingyue Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D11314 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@242713 91177308-0d34-0410-b5e6-96231b3b80d8		2015-07-20 21:28:54 +00:00
..
InstPrinter
MCTargetDesc
TargetInfo
cl_common_defines.h
CMakeLists.txt
LLVMBuild.txt
Makefile
ManagedStringPool.h
NVPTX.h
NVPTX.td
NVPTXAllocaHoisting.cpp
NVPTXAllocaHoisting.h
NVPTXAsmPrinter.cpp
NVPTXAsmPrinter.h
NVPTXAssignValidGlobalNames.cpp
NVPTXFavorNonGenericAddrSpaces.cpp
NVPTXFrameLowering.cpp
NVPTXFrameLowering.h
NVPTXGenericToNVVM.cpp
NVPTXImageOptimizer.cpp
NVPTXInstrFormats.td
NVPTXInstrInfo.cpp
NVPTXInstrInfo.h
NVPTXInstrInfo.td
NVPTXIntrinsics.td
NVPTXISelDAGToDAG.cpp	[NVPTX] make load on global readonly memory to use ldg	2015-07-20 21:28:54 +00:00
NVPTXISelDAGToDAG.h
NVPTXISelLowering.cpp
NVPTXISelLowering.h
NVPTXLowerAggrCopies.cpp
NVPTXLowerAggrCopies.h
NVPTXLowerAlloca.cpp
NVPTXLowerKernelArgs.cpp
NVPTXMachineFunctionInfo.h
NVPTXMCExpr.cpp
NVPTXMCExpr.h
NVPTXPeephole.cpp
NVPTXPrologEpilogPass.cpp
NVPTXRegisterInfo.cpp
NVPTXRegisterInfo.h
NVPTXRegisterInfo.td
NVPTXReplaceImageHandles.cpp
NVPTXSection.h
NVPTXSubtarget.cpp
NVPTXSubtarget.h
NVPTXTargetMachine.cpp
NVPTXTargetMachine.h
NVPTXTargetObjectFile.h
NVPTXTargetTransformInfo.cpp
NVPTXTargetTransformInfo.h
NVPTXUtilities.cpp
NVPTXUtilities.h
NVPTXVector.td
NVVMReflect.cpp