1
0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-26 04:32:44 +01:00
llvm-mirror/lib/Target/NVPTX
Steffen Larsen 1e7a7bb573 [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions
Adds NVPTX builtins and intrinsics for the CUDA PTX `redux.sync` instructions
for `sm_80` architecture or newer.

PTX ISA description of `redux.sync`:
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-redux-sync

Authored-by: Steffen Larsen <steffen.larsen@codeplay.com>

Differential Revision: https://reviews.llvm.org/D100124
2021-05-17 09:46:59 -07:00
..
MCTargetDesc [MC] Untangle MCContext and MCObjectFileInfo 2021-05-05 10:03:02 -07:00
TargetInfo llvmbuildectomy - replace llvm-build by plain cmake 2020-11-13 10:35:24 +01:00
cl_common_defines.h
CMakeLists.txt [NVPTX] Enable lowering of atomics on local memory 2021-04-26 20:12:12 -04:00
ManagedStringPool.h
NVPTX.h [NewPM][NVPTX] Port NVPTX opt passes 2021-01-07 15:12:35 -08:00
NVPTX.td [CUDA, NVPTX] Allow targeting sm_86 GPUs. 2021-02-09 11:01:10 -08:00
NVPTXAllocaHoisting.cpp
NVPTXAllocaHoisting.h
NVPTXAsmPrinter.cpp Do not construct std::string from nullptr 2020-11-05 15:23:26 -08:00
NVPTXAsmPrinter.h
NVPTXAssignValidGlobalNames.cpp
NVPTXAtomicLower.cpp [NVPTX] Enable lowering of atomics on local memory 2021-04-26 20:12:12 -04:00
NVPTXAtomicLower.h [NVPTX] Enable lowering of atomics on local memory 2021-04-26 20:12:12 -04:00
NVPTXFrameLowering.cpp [SVE] Return StackOffset for TargetFrameLowering::getFrameIndexReference. 2020-11-05 11:02:18 +00:00
NVPTXFrameLowering.h [SVE] Return StackOffset for TargetFrameLowering::getFrameIndexReference. 2020-11-05 11:02:18 +00:00
NVPTXGenericToNVVM.cpp
NVPTXImageOptimizer.cpp
NVPTXInstrFormats.td [NVPTX] [TableGen] Use new features of TableGen to simplify and clarify. 2020-11-06 09:20:19 -05:00
NVPTXInstrInfo.cpp
NVPTXInstrInfo.h
NVPTXInstrInfo.td [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async instructions 2021-05-17 09:46:59 -07:00
NVPTXIntrinsics.td [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions 2021-05-17 09:46:59 -07:00
NVPTXISelDAGToDAG.cpp [llvm] Use the default value of drop_begin (NFC) 2021-01-18 10:16:36 -08:00
NVPTXISelDAGToDAG.h
NVPTXISelLowering.cpp [ValueTypes] Rename MVT::getVectorNumElements() to MVT::getVectorMinNumElements(). Fix some misuses of getVectorNumElements() 2021-05-12 07:46:45 -07:00
NVPTXISelLowering.h [llvm][nvptx] add atomicity to counter in ISelLowering 2021-01-19 10:20:20 +01:00
NVPTXLowerAggrCopies.cpp
NVPTXLowerAggrCopies.h
NVPTXLowerAlloca.cpp
NVPTXLowerArgs.cpp [NVPTX] Fix unused var warning with asserts disabled 2021-04-29 09:54:03 +01:00
NVPTXMachineFunctionInfo.h
NVPTXMCExpr.cpp
NVPTXMCExpr.h
NVPTXPeephole.cpp
NVPTXPrologEpilogPass.cpp [NFC][CodeGen] Tidy up TargetRegisterInfo stack realignment functions 2021-03-30 17:31:39 +01:00
NVPTXProxyRegErasure.cpp
NVPTXRegisterInfo.cpp
NVPTXRegisterInfo.h
NVPTXRegisterInfo.td [NVPTX] [TableGen] Use new features of TableGen to simplify and clarify. 2020-11-06 09:20:19 -05:00
NVPTXReplaceImageHandles.cpp
NVPTXSubtarget.cpp [X86][MC][Target] Initial backend support a tune CPU to support -mtune 2020-08-14 15:31:50 -07:00
NVPTXSubtarget.h [X86][MC][Target] Initial backend support a tune CPU to support -mtune 2020-08-14 15:31:50 -07:00
NVPTXTargetMachine.cpp [NewPM] Hide pass manager debug logging behind -debug-pass-manager-verbose 2021-05-07 21:51:47 -07:00
NVPTXTargetMachine.h [NewPM] Hide pass manager debug logging behind -debug-pass-manager-verbose 2021-05-07 21:51:47 -07:00
NVPTXTargetObjectFile.h
NVPTXTargetTransformInfo.cpp [TTI] NFC: Change getTypeLegalizationCost to return InstructionCost. 2021-04-30 22:51:51 +03:00
NVPTXTargetTransformInfo.h [TTI] NFC: Change getArithmeticInstrCost to return InstructionCost 2021-04-14 17:20:36 +01:00
NVPTXUtilities.cpp
NVPTXUtilities.h
NVVMIntrRange.cpp [NewPM][NVPTX] Port NVPTX opt passes 2021-01-07 15:12:35 -08:00
NVVMReflect.cpp [NewPM][NVPTX] Port NVPTX opt passes 2021-01-07 15:12:35 -08:00