6502/llvm-6502
1
0
Fork 0
mirror of https://github.com/c64scene-ar/llvm-6502.git synced 2024-07-22 09:29:31 +00:00
Code Issues Projects Releases Wiki Activity
bb418bc23b
llvm-6502/lib/Target/NVPTX
History
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 Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
MCTargetDesc MC: Remove MCSubtargetInfo() default constructor 2015-07-10 22:43:42 +00:00
TargetInfo
cl_common_defines.h Update a couple of header inclusion guards 2014-05-31 21:26:09 +00:00
CMakeLists.txt Add NVPTXPeephole pass to reduce unnecessary address cast 2015-06-24 20:20:16 +00:00
LLVMBuild.txt Update libdeps in NVPTXCodeGen, since r225944. 2015-01-14 23:01:36 +00:00
Makefile
ManagedStringPool.h Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTX.h Add NVPTXPeephole pass to reduce unnecessary address cast 2015-06-24 20:20:16 +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 Move most user of TargetMachine::getDataLayout to the Module one 2015-07-16 06:11:10 +00:00
NVPTXAsmPrinter.h Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXAssignValidGlobalNames.cpp Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXFavorNonGenericAddrSpaces.cpp Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXFrameLowering.cpp [NVPTX] cleanups and refacotring in NVPTXFrameLowering.cpp 2015-06-30 21:28:31 +00:00
NVPTXFrameLowering.h Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXGenericToNVVM.cpp [opaque pointer type] More GEP IRBuilder API migrations... 2015-04-03 21:33:42 +00:00
NVPTXImageOptimizer.cpp Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXInstrFormats.td [NVPTX] Add more surface/texture intrinsics, including CUDA unified texture fetch 2014-07-17 11:59:04 +00:00
NVPTXInstrInfo.cpp [CodeGen] ArrayRef'ize cond/pred in various TII APIs. NFC. 2015-06-11 19:30:37 +00:00
NVPTXInstrInfo.h [CodeGen] ArrayRef'ize cond/pred in various TII APIs. NFC. 2015-06-11 19:30:37 +00:00
NVPTXInstrInfo.td Reapply r235977 "[DebugInfo] Add debug locations to constant SD nodes" 2015-04-28 14:05:47 +00:00
NVPTXIntrinsics.td NVPTX: support direct f16 <-> f64 conversions via intrinsics. 2014-07-18 08:30:10 +00:00
NVPTXISelDAGToDAG.cpp [NVPTX] make load on global readonly memory to use ldg 2015-07-20 21:28:54 +00:00
NVPTXISelDAGToDAG.h Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXISelLowering.cpp Remove getDataLayout() from TargetLowering 2015-07-09 02:09:52 +00:00
NVPTXISelLowering.h Re-instate the EVT parameter to getScalarShiftAmountTy() for OOT user 2015-07-09 15:12:23 +00:00
NVPTXLowerAggrCopies.cpp Use inbounds GEPs for memcpy and memset lowering 2015-07-17 16:42:33 +00:00
NVPTXLowerAggrCopies.h Move unreferenced passes into the cpp file 2015-03-09 15:50:58 +00:00
NVPTXLowerAlloca.cpp Add NVPTXLowerAlloca pass to convert alloca'ed memory to local address 2015-06-17 22:31:02 +00:00
NVPTXLowerKernelArgs.cpp [NVPTX] noop when kernel pointers are already global 2015-06-26 22:35:43 +00:00
NVPTXMachineFunctionInfo.h Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXMCExpr.cpp MC: Add target hook to control symbol quoting 2015-06-09 00:31:39 +00:00
NVPTXMCExpr.h MC: Add target hook to control symbol quoting 2015-06-09 00:31:39 +00:00
NVPTXPeephole.cpp [NVPTX] Move NVPTXPeephole after NVPTXPrologEpilogPass 2015-07-01 20:08:06 +00:00
NVPTXPrologEpilogPass.cpp Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXRegisterInfo.cpp Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXRegisterInfo.h Have getCalleeSavedRegs take a non-null MachineFunction all the 2015-03-11 21:41:28 +00:00
NVPTXRegisterInfo.td Add NVPTXPeephole pass to reduce unnecessary address cast 2015-06-24 20:20:16 +00:00
NVPTXReplaceImageHandles.cpp Cosmetic cleanups - NFC 2015-07-08 16:33:21 +00:00
NVPTXSection.h Implement unique sections with an unique ID. 2015-04-04 18:02:01 +00:00
NVPTXSubtarget.cpp Remove getDataLayout() from TargetSelectionDAGInfo (had no users) 2015-07-09 02:10:08 +00:00
NVPTXSubtarget.h Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXTargetMachine.cpp [NVPTX] enable SpeculativeExecution in NVPTX 2015-07-16 20:13:48 +00:00
NVPTXTargetMachine.h Replace string GNU Triples with llvm::Triple in TargetMachine. NFC. 2015-06-11 19:41:26 +00:00
NVPTXTargetObjectFile.h Remove DataLayout from TargetLoweringObjectFile, redirect to Module 2015-07-16 06:04:17 +00:00
NVPTXTargetTransformInfo.cpp Enable partial and runtime loop unrolling for NVPTX. 2015-07-13 18:33:21 +00:00
NVPTXTargetTransformInfo.h Enable partial and runtime loop unrolling for NVPTX. 2015-07-13 18:33:21 +00:00
NVPTXUtilities.cpp Simplify boolean expressions with true and false using clang-tidy 2015-03-23 16:26:23 +00:00
NVPTXUtilities.h Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00
NVPTXVector.td Reapply r235977 "[DebugInfo] Add debug locations to constant SD nodes" 2015-04-28 14:05:47 +00:00
NVVMReflect.cpp Revert r240137 (Fixed/added namespace ending comments using clang-tidy. NFC) 2015-06-23 09:49:53 +00:00