LLVM 23.0.0git
Macros | Enumerations | Functions | Variables
NVPTXISelLowering.cpp File Reference
#include "NVPTXISelLowering.h"
#include "MCTargetDesc/NVPTXBaseInfo.h"
#include "NVPTX.h"
#include "NVPTXISelDAGToDAG.h"
#include "NVPTXSelectionDAGInfo.h"
#include "NVPTXSubtarget.h"
#include "NVPTXTargetMachine.h"
#include "NVPTXTargetObjectFile.h"
#include "NVPTXUtilities.h"
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/APInt.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/CodeGen/Analysis.h"
#include "llvm/CodeGen/ISDOpcodes.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/CodeGen/MachineJumpTableInfo.h"
#include "llvm/CodeGen/MachineMemOperand.h"
#include "llvm/CodeGen/SDPatternMatch.h"
#include "llvm/CodeGen/SelectionDAG.h"
#include "llvm/CodeGen/SelectionDAGNodes.h"
#include "llvm/CodeGen/TargetCallingConv.h"
#include "llvm/CodeGen/TargetLowering.h"
#include "llvm/CodeGen/ValueTypes.h"
#include "llvm/CodeGenTypes/MachineValueType.h"
#include "llvm/IR/Argument.h"
#include "llvm/IR/Attributes.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/DiagnosticInfo.h"
#include "llvm/IR/FPEnv.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/GlobalValue.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instruction.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/Module.h"
#include "llvm/IR/Type.h"
#include "llvm/IR/Value.h"
#include "llvm/Support/Alignment.h"
#include "llvm/Support/AtomicOrdering.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/CodeGen.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/KnownBits.h"
#include "llvm/Support/NVPTXAddrSpace.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/Target/TargetMachine.h"
#include "llvm/Target/TargetOptions.h"
#include <algorithm>
#include <cassert>
#include <cmath>
#include <cstdint>
#include <iterator>
#include <optional>
#include <string>
#include <tuple>
#include <utility>
#include <vector>

Go to the source code of this file.

Macros

#define DEBUG_TYPE   "nvptx-lower"
#define TCGEN05_LD_RED_INTR(SHAPE, NUM, TYPE)
#define TCGEN05_LD_RED_INST(SHAPE, NUM, TYPE)

Enumerations

enum  OperandSignedness { Signed = 0 , Unsigned , Unknown }

Functions

static bool IsPTXVectorType (MVT VT)
static std::optional< std::pair< unsigned int, MVT > > getVectorLoweringShape (EVT VectorEVT, const NVPTXSubtarget &STI, unsigned AddressSpace)
static void ComputePTXValueVTs (const TargetLowering &TLI, const DataLayout &DL, LLVMContext &Ctx, CallingConv::ID CallConv, Type *Ty, SmallVectorImpl< EVT > &ValueVTs, SmallVectorImpl< uint64_t > &Offsets, uint64_t StartingOffset=0)
 ComputePTXValueVTs - For the given Type Ty, returns the set of primitive legal-ish MVTs that compose it.
static EVT getVectorizedVT (EVT VT, unsigned N, LLVMContext &C)
static SDValue getExtractVectorizedValue (SDValue V, unsigned I, EVT VT, const SDLoc &dl, SelectionDAG &DAG)
template<typename T>
static SDValue getBuildVectorizedValue (unsigned N, const SDLoc &dl, SelectionDAG &DAG, T GetElement)
static EVT promoteScalarIntegerPTX (const EVT VT)
 PromoteScalarIntegerPTX Used to make sure the arguments/returns are suitable for passing and promote them to a larger size if they're not.
template<typename T>
static unsigned canMergeParamLoadStoresStartingAt (unsigned Idx, uint32_t AccessSize, const SmallVectorImpl< EVT > &ValueVTs, const SmallVectorImpl< T > &Offsets, Align ParamAlignment)
template<typename T>
static SmallVector< unsigned, 16 > VectorizePTXValueVTs (const SmallVectorImpl< EVT > &ValueVTs, const SmallVectorImpl< T > &Offsets, Align ParamAlignment, bool IsVAArg=false)
static bool shouldConvertToIndirectCall (const CallBase *CB, const GlobalAddressSDNode *Func)
static MachinePointerInfo refinePtrAS (SDValue &Ptr, SelectionDAG &DAG, const DataLayout &DL, const TargetLowering &TL)
static ISD::NodeType getExtOpcode (const ISD::ArgFlagsTy &Flags)
static SDValue correctParamType (SDValue V, EVT ExpectedVT, ISD::ArgFlagsTy Flags, SelectionDAG &DAG, SDLoc dl)
static SDValue getPRMT (SDValue A, SDValue B, SDValue Selector, SDLoc DL, SelectionDAG &DAG, unsigned Mode=NVPTX::PTXPrmtMode::NONE)
static SDValue getPRMT (SDValue A, SDValue B, uint64_t Selector, SDLoc DL, SelectionDAG &DAG, unsigned Mode=NVPTX::PTXPrmtMode::NONE)
static SDValue buildTreeReduction (const SmallVector< SDValue > &Elements, EVT EltTy, ArrayRef< std::pair< unsigned, unsigned > > Ops, const SDLoc &DL, const SDNodeFlags Flags, SelectionDAG &DAG)
 Reduces the elements using the scalar operations provided.
static ISD::NodeType getScalarOpcodeForReduction (unsigned ReductionOpcode)
static std::optional< unsigned > getScalar3OpcodeForReduction (unsigned ReductionOpcode)
 Get 3-input scalar reduction opcode.
static SDValue PromoteBinOpToF32 (SDNode *N, SelectionDAG &DAG)
static SDValue LowerVectorArith (SDValue Op, SelectionDAG &DAG)
static SDValue lowerTcgen05St (SDValue Op, SelectionDAG &DAG)
static SDValue lowerBSWAP (SDValue Op, SelectionDAG &DAG)
static unsigned getTcgen05MMADisableOutputLane (unsigned IID)
static SDValue LowerTcgen05MMADisableOutputLane (SDValue Op, SelectionDAG &DAG)
static std::optional< std::pair< SDValue, SDValue > > lowerTcgen05Ld (SDNode *N, SelectionDAG &DAG, bool HasOffset=false)
static SDValue reportInvalidTensormapReplaceUsage (SDValue Op, SelectionDAG &DAG, unsigned Val)
static SDValue lowerTensormapReplaceElemtype (SDValue Op, SelectionDAG &DAG)
static SDValue lowerTensormapReplaceSwizzleMode (SDValue Op, SelectionDAG &DAG)
static SDValue lowerIntrinsicVoid (SDValue Op, SelectionDAG &DAG)
static SDValue LowerClusterLaunchControlQueryCancel (SDValue Op, SelectionDAG &DAG)
static SDValue lowerCvtRSIntrinsics (SDValue Op, SelectionDAG &DAG)
static SDValue lowerPrmtIntrinsic (SDValue Op, SelectionDAG &DAG)
static unsigned getTcgen05LdRedID (Intrinsic::ID IID)
static std::optional< std::tuple< SDValue, SDValue, SDValue > > lowerTcgen05LdRed (SDNode *N, SelectionDAG &DAG)
static SDValue lowerIntrinsicWChain (SDValue Op, SelectionDAG &DAG)
static SDValue lowerIntrinsicWOChain (SDValue Op, SelectionDAG &DAG)
static SDValue lowerCTLZCTPOP (SDValue Op, SelectionDAG &DAG)
static SDValue expandFSH64 (SDValue A, SDValue B, SDValue ShiftAmount, SDLoc DL, unsigned Opcode, SelectionDAG &DAG)
static SDValue lowerFSH (SDValue Op, SelectionDAG &DAG)
static SDValue lowerROT (SDValue Op, SelectionDAG &DAG)
static SDValue lowerFREM (SDValue Op, SelectionDAG &DAG)
static SDValue lowerSELECT (SDValue Op, SelectionDAG &DAG)
static SDValue lowerMSTORE (SDValue Op, SelectionDAG &DAG)
static std::pair< MemSDNode *, uint32_t > convertMLOADToLoadWithUsedBytesMask (MemSDNode *N, SelectionDAG &DAG, const NVPTXSubtarget &STI)
static std::optional< std::pair< SDValue, SDValue > > replaceLoadVector (SDNode *N, SelectionDAG &DAG, const NVPTXSubtarget &STI)
 replaceLoadVector - Convert vector loads into multi-output scalar loads.
static void replaceLoadVector (SDNode *N, SelectionDAG &DAG, SmallVectorImpl< SDValue > &Results, const NVPTXSubtarget &STI)
static SDValue lowerLoadVector (SDNode *N, SelectionDAG &DAG, const NVPTXSubtarget &STI)
static SDValue lowerLOADi1 (LoadSDNode *LD, SelectionDAG &DAG)
static SDValue lowerSTOREVector (SDValue Op, SelectionDAG &DAG, const NVPTXSubtarget &STI)
static bool isConstZero (const SDValue &Operand)
static SDValue PerformADDCombineWithOperands (SDNode *N, SDValue N0, SDValue N1, TargetLowering::DAGCombinerInfo &DCI)
 PerformADDCombineWithOperands - Try DAG combinations for an ADD with operands N0 and N1.
static SDValue PerformFADDCombineWithOperands (SDNode *N, SDValue N0, SDValue N1, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
static SDValue combineUnpackingMovIntoLoad (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
 Fold unpacking movs into a load by increasing the number of return values.
static SDValue combinePackingMovIntoStore (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, unsigned Front, unsigned Back)
 Fold packing movs into a store.
static SDValue combineSTORE (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const NVPTXSubtarget &STI)
static SDValue combineLOAD (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const NVPTXSubtarget &STI)
static SDValue PerformADDCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
 PerformADDCombine - Target-specific dag combine xforms for ISD::ADD.
static SDValue PerformFADDCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
 PerformFADDCombine - Target-specific dag combine xforms for ISD::FADD.
static unsigned getMinMax3Opcode (unsigned MinMax2Opcode)
 Get 3-input version of a 2-input min/max opcode.
static SDValue PerformFMinMaxCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, unsigned PTXVersion, unsigned SmVersion)
 PerformFMinMaxCombine - Combine (fmaxnum (fmaxnum a, b), c) into (fmaxnum3 a, b, c).
static SDValue PerformREMCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
static SDValue combineMulWide (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
static bool IsMulWideOperandDemotable (SDValue Op, unsigned OptSize, OperandSignedness &S)
 IsMulWideOperandDemotable - Checks if the provided DAG node is an operand that can be demoted to OptSize bits without loss of information.
static bool AreMulWideOperandsDemotable (SDValue LHS, SDValue RHS, unsigned OptSize, bool &IsSigned)
 AreMulWideOperandsDemotable - Checks if the given LHS and RHS operands can be demoted to OptSize bits without loss of information.
static SDValue TryMULWIDECombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
 TryMULWIDECombine - Attempt to replace a multiply of M bits with a multiply of M/2 bits that produces an M-bit result (i.e.
static bool isConstOne (const SDValue &Operand)
static SDValue matchMADConstOnePattern (SDValue Add)
static SDValue combineMADConstOne (SDValue X, SDValue Add, EVT VT, SDLoc DL, TargetLowering::DAGCombinerInfo &DCI)
static SDValue combineMulSelectConstOne (SDValue X, SDValue Select, EVT VT, SDLoc DL, TargetLowering::DAGCombinerInfo &DCI)
static SDValue PerformMULCombineWithOperands (SDNode *N, SDValue N0, SDValue N1, TargetLowering::DAGCombinerInfo &DCI)
static SDValue PerformMULCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
 PerformMULCombine - Runs PTX-specific DAG combine patterns on MUL nodes.
static SDValue PerformSHLCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
 PerformSHLCombine - Runs PTX-specific DAG combine patterns on SHL nodes.
static SDValue PerformSETCCCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, unsigned int SmVersion)
static SDValue PerformEXTRACTCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
static SDValue PerformSELECTShiftCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
 Transform patterns like: (select (ugt shift_amt, BitWidth-1), 0, (srl/shl x, shift_amt)) (select (ult shift_amt, BitWidth), (srl/shl x, shift_amt), 0) Into: (NVPTXISD::SRL_CLAMP x, shift_amt) or (NVPTXISD::SHL_CLAMP x, shift_amt)
static SDValue PerformVSELECTCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
static SDValue PerformBUILD_VECTORCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
static SDValue combineADDRSPACECAST (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
static APInt getPRMTSelector (const APInt &Selector, unsigned Mode)
static APInt computePRMT (APInt A, APInt B, APInt Selector, unsigned Mode)
static SDValue combinePRMT (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
static SDValue sinkProxyReg (SDValue R, SDValue Chain, TargetLowering::DAGCombinerInfo &DCI)
static SDValue combineProxyReg (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
static void ReplaceBITCAST (SDNode *Node, SelectionDAG &DAG, SmallVectorImpl< SDValue > &Results)
static void ReplaceINTRINSIC_W_CHAIN (SDNode *N, SelectionDAG &DAG, SmallVectorImpl< SDValue > &Results)
static void ReplaceCopyFromReg_128 (SDNode *N, SelectionDAG &DAG, SmallVectorImpl< SDValue > &Results)
static void replaceProxyReg (SDNode *N, SelectionDAG &DAG, const TargetLowering &TLI, SmallVectorImpl< SDValue > &Results)
static void replaceAtomicSwap128 (SDNode *N, SelectionDAG &DAG, const NVPTXSubtarget &STI, SmallVectorImpl< SDValue > &Results)
static void computeKnownBitsForPRMT (const SDValue Op, KnownBits &Known, const SelectionDAG &DAG, unsigned Depth)
static void computeKnownBitsForLoadV (const SDValue Op, KnownBits &Known)
static std::pair< APInt, APInt > getPRMTDemandedBits (const APInt &SelectorVal, const APInt &DemandedBits)
static SDValue canonicalizePRMTInput (SDValue Op, SelectionDAG &DAG)
static SDValue simplifyDemandedBitsForPRMT (SDValue PRMT, const APInt &DemandedBits, SelectionDAG &DAG, const TargetLowering &TLI, unsigned Depth)

Variables

static cl::opt< bool > sched4reg ("nvptx-sched4reg", cl::desc("NVPTX Specific: schedule for register pressue"), cl::init(false))
static cl::opt< unsigned > FMAContractLevelOpt ("nvptx-fma-level", cl::Hidden, cl::desc("NVPTX Specific: FMA contraction (0: don't do it" " 1: do it 2: do it aggressively"), cl::init(2))
static cl::opt< NVPTX::DivPrecisionLevel > UsePrecDivF32 ("nvptx-prec-divf32", cl::Hidden, cl::desc("NVPTX Specific: Override the precision of the lowering for f32 fdiv"), cl::values(clEnumValN(NVPTX::DivPrecisionLevel::Approx, "0", "Use div.approx"), clEnumValN(NVPTX::DivPrecisionLevel::Full, "1", "Use div.full"), clEnumValN(NVPTX::DivPrecisionLevel::IEEE754, "2", "Use IEEE Compliant F32 div.rnd if available (default)"), clEnumValN(NVPTX::DivPrecisionLevel::IEEE754_NoFTZ, "3", "Use IEEE Compliant F32 div.rnd if available, no FTZ")), cl::init(NVPTX::DivPrecisionLevel::IEEE754))
static cl::opt< bool > UsePrecSqrtF32 ("nvptx-prec-sqrtf32", cl::Hidden, cl::desc("NVPTX Specific: 0 use sqrt.approx, 1 use sqrt.rn."), cl::init(true))
static cl::opt< bool > UseApproxLog2F32 ("nvptx-approx-log2f32", cl::desc("NVPTX Specific: whether to use lg2.approx for log2"), cl::init(false))
 Whereas CUDA's implementation (see libdevice) uses ex2.approx for exp2(), it does NOT use lg2.approx for log2, so this is disabled by default.
static cl::opt< bool > ForceMinByValParamAlign ("nvptx-force-min-byval-param-align", cl::Hidden, cl::desc("NVPTX Specific: force 4-byte minimal alignment for byval" " params of device functions."), cl::init(false))

Macro Definition Documentation

◆ DEBUG_TYPE

#define DEBUG_TYPE   "nvptx-lower"

Definition at line 78 of file NVPTXISelLowering.cpp.

◆ TCGEN05_LD_RED_INST

#define TCGEN05_LD_RED_INST ( SHAPE,
NUM,
TYPE )
Value:
NVPTXISD::TCGEN05_LD_RED_##SHAPE##_X##NUM##_##TYPE
_
#define _
Definition HexagonMCCodeEmitter.cpp:46

Definition at line 3013 of file NVPTXISelLowering.cpp.

Referenced by getTcgen05LdRedID().

◆ TCGEN05_LD_RED_INTR

#define TCGEN05_LD_RED_INTR ( SHAPE,
NUM,
TYPE )
Value:
Intrinsic::nvvm_tcgen05_ld_red_##SHAPE##_x##NUM##_##TYPE

Definition at line 3010 of file NVPTXISelLowering.cpp.

Referenced by getTcgen05LdRedID().

Enumeration Type Documentation

◆ OperandSignedness

enum OperandSignedness
Enumerator
Signed 
Unsigned 
Unknown 

Definition at line 6206 of file NVPTXISelLowering.cpp.

Function Documentation

◆ AreMulWideOperandsDemotable()

bool AreMulWideOperandsDemotable ( SDValue LHS,
SDValue RHS,
unsigned OptSize,
bool & IsSigned )
static

AreMulWideOperandsDemotable - Checks if the given LHS and RHS operands can be demoted to OptSize bits without loss of information.

If the operands contain a constant, it should appear as the RHS operand. The signedness of the operands is placed in IsSigned.

Definition at line 6242 of file NVPTXISelLowering.cpp.

References llvm::dyn_cast(), llvm::APInt::isIntN(), IsMulWideOperandDemotable(), llvm::APInt::isSignedIntN(), LHS, RHS, Signed, Unknown, and Unsigned.

Referenced by TryMULWIDECombine().

◆ buildTreeReduction()

SDValue buildTreeReduction ( const SmallVector< SDValue > & Elements,
EVT EltTy,
ArrayRef< std::pair< unsigned, unsigned > > Ops,
const SDLoc & DL,
const SDNodeFlags Flags,
SelectionDAG & DAG )
static

Reduces the elements using the scalar operations provided.

The operations are sorted descending in number of inputs they take. The flags on the original reduction operation will be propagated to each scalar operation. Nearby elements are grouped in tree reduction, unlike the shuffle reduction used in ExpandReductions and SelectionDAG.

Definition at line 1890 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, assert(), DL, E(), llvm::SmallVectorTemplateCommon< T, typename >::empty(), llvm::SelectionDAG::getNode(), I, OpIdx, llvm::SmallVectorTemplateBase< T, bool >::push_back(), and llvm::ArrayRef< T >::slice().

◆ canMergeParamLoadStoresStartingAt()

template<typename T>
unsigned canMergeParamLoadStoresStartingAt ( unsigned Idx,
uint32_t AccessSize,
const SmallVectorImpl< EVT > & ValueVTs,
const SmallVectorImpl< T > & Offsets,
Align ParamAlignment )
static

Definition at line 427 of file NVPTXISelLowering.cpp.

References llvm::EVT::getStoreSize(), and llvm::SmallVectorTemplateCommon< T, typename >::size().

Referenced by VectorizePTXValueVTs().

◆ canonicalizePRMTInput()

SDValue canonicalizePRMTInput ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 7509 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::getConstant(), and SDValue().

Referenced by simplifyDemandedBitsForPRMT().

◆ combineADDRSPACECAST()

SDValue combineADDRSPACECAST ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6686 of file NVPTXISelLowering.cpp.

References assert(), llvm::cast(), llvm::dyn_cast(), N, and SDValue().

◆ combineLOAD()

SDValue combineLOAD ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
const NVPTXSubtarget & STI )
static

Definition at line 6032 of file NVPTXISelLowering.cpp.

References combineUnpackingMovIntoLoad(), llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::TargetLowering::DAGCombinerInfo::isBeforeLegalize(), llvm::ISD::LOAD, lowerLoadVector(), and N.

◆ combineMADConstOne()

SDValue combineMADConstOne ( SDValue X,
SDValue Add,
EVT VT,
SDLoc DL,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6362 of file NVPTXISelLowering.cpp.

References llvm::ISD::ADD, llvm::Add, llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::SelectionDAG::getNode(), matchMADConstOnePattern(), llvm::ISD::MUL, Mul, SDValue(), X, and Y.

Referenced by PerformMULCombineWithOperands().

◆ combineMulSelectConstOne()

SDValue combineMulSelectConstOne ( SDValue X,
SDValue Select,
EVT VT,
SDLoc DL,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6373 of file NVPTXISelLowering.cpp.

References Cond, llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::SelectionDAG::getNode(), isConstOne(), matchMADConstOnePattern(), llvm::ISD::MUL, SDValue(), llvm::ISD::SELECT, Select, X, and Y.

Referenced by PerformMULCombineWithOperands().

◆ combineMulWide()

SDValue combineMulWide ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

Definition at line 6171 of file NVPTXISelLowering.cpp.

References llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getNode(), llvm::EVT::getSizeInBits(), llvm::isa(), llvm::ISD::MUL, N, llvm::None, RHS, SDValue(), llvm::ISD::SHL, llvm::ISD::SIGN_EXTEND, and llvm::ISD::ZERO_EXTEND.

◆ combinePackingMovIntoStore()

SDValue combinePackingMovIntoStore ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
unsigned Front,
unsigned Back )
static

Fold packing movs into a store.

ex: v1: v2f16 = BUILD_VECTOR a:f16, b:f16 v2: v2f16 = BUILD_VECTOR c:f16, d:f16 StoreV2 v1, v2

...is turned into...

StoreV4 a, b, c, d

Definition at line 5937 of file NVPTXISelLowering.cpp.

References llvm::SmallVectorImpl< T >::append(), llvm::ISD::BITCAST, llvm::ISD::BUILD_VECTOR, llvm::cast(), llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::ISD::FP_ROUND, llvm::SelectionDAG::getMemIntrinsicNode(), llvm::TargetLowering::DAGCombinerInfo::isAfterLegalizeDAG(), llvm::NVPTX::isPackedVectorTy(), llvm_unreachable, N, SDValue(), llvm::ISD::STORE, llvm::NVPTXISD::StoreV2, llvm::NVPTXISD::StoreV4, llvm::NVPTXISD::StoreV8, and llvm::ISD::TRUNCATE.

Referenced by combineSTORE().

◆ combinePRMT()

SDValue combinePRMT ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

Definition at line 6757 of file NVPTXISelLowering.cpp.

References computePRMT(), llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::SelectionDAG::getConstant(), llvm::isa(), N, llvm::None, and SDValue().

◆ combineProxyReg()

SDValue combineProxyReg ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6839 of file NVPTXISelLowering.cpp.

References llvm::ISD::LOAD, N, Reg, SDValue(), and sinkProxyReg().

◆ combineSTORE()

SDValue combineSTORE ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
const NVPTXSubtarget & STI )
static

Definition at line 6017 of file NVPTXISelLowering.cpp.

References llvm::cast(), combinePackingMovIntoStore(), llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::TargetLowering::DAGCombinerInfo::isBeforeLegalize(), lowerSTOREVector(), N, SDValue(), and llvm::ISD::STORE.

◆ combineUnpackingMovIntoLoad()

SDValue combineUnpackingMovIntoLoad ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

Fold unpacking movs into a load by increasing the number of return values.

ex: L: v2f16,ch = load

a: f16 = extractelt L:0, 0 b: f16 = extractelt L:0, 1 use(a, b)

...is turned into...

L: f16,f16,ch = LoadV2

use(L:0, L:1)

Definition at line 5823 of file NVPTXISelLowering.cpp.

References llvm::all_of(), llvm::TargetLowering::DAGCombinerInfo::isAfterLegalizeDAG(), llvm::NVPTX::isPackedVectorTy(), N, and SDValue().

Referenced by combineLOAD().

◆ computeKnownBitsForLoadV()

void computeKnownBitsForLoadV ( const SDValue Op,
KnownBits & Known )
static

Definition at line 7446 of file NVPTXISelLowering.cpp.

References assert(), llvm::cast(), llvm::KnownBits::getBitWidth(), llvm::NVPTXDAGToDAGISel::getFromTypeWidthForLoad(), llvm::APInt::setHighBits(), llvm::ISD::SEXTLOAD, and llvm::KnownBits::Zero.

Referenced by llvm::NVPTXTargetLowering::computeKnownBitsForTargetNode().

◆ computeKnownBitsForPRMT()

void computeKnownBitsForPRMT ( const SDValue Op,
KnownBits & Known,
const SelectionDAG & DAG,
unsigned Depth )
static

Definition at line 7415 of file NVPTXISelLowering.cpp.

References A(), llvm::KnownBits::ashr(), assert(), B(), llvm::SelectionDAG::computeKnownBits(), llvm::KnownBits::concat(), llvm::Depth, llvm::dyn_cast(), llvm::APInt::extractBits(), llvm::ConstantSDNode::getAPIntValue(), llvm::KnownBits::getBitWidth(), llvm::APInt::getHiBits(), llvm::APInt::getLoBits(), getPRMTSelector(), llvm::APInt::getZExtValue(), I, llvm::KnownBits::insertBits(), Mode, and llvm::seq().

Referenced by llvm::NVPTXTargetLowering::computeKnownBitsForTargetNode().

◆ computePRMT()

APInt computePRMT ( APInt A,
APInt B,
APInt Selector,
unsigned Mode )
static

Definition at line 6738 of file NVPTXISelLowering.cpp.

References A(), assert(), B(), llvm::APInt::extractBits(), llvm::APInt::getBitWidth(), llvm::APInt::getHiBits(), llvm::APInt::getLoBits(), getPRMTSelector(), llvm::APInt::getZExtValue(), I, Mode, and llvm::seq().

Referenced by combinePRMT().

◆ ComputePTXValueVTs()

void ComputePTXValueVTs ( const TargetLowering & TLI,
const DataLayout & DL,
LLVMContext & Ctx,
CallingConv::ID CallConv,
Type * Ty,
SmallVectorImpl< EVT > & ValueVTs,
SmallVectorImpl< uint64_t > & Offsets,
uint64_t StartingOffset = 0 )
static

ComputePTXValueVTs - For the given Type Ty, returns the set of primitive legal-ish MVTs that compose it.

Unlike ComputeValueVTs, this will legalize the types as required by the calling convention (with special handling for i8s). NOTE: This is a band-aid for code that expects ComputeValueVTs to return the same number of types as the Ins/Outs arrays in LowerFormalArguments, LowerCall, and LowerReturn.

Definition at line 303 of file NVPTXISelLowering.cpp.

References assert(), llvm::ComputeValueVTs(), DL, llvm::TargetLoweringBase::getNumRegistersForCallingConv(), llvm::TargetLoweringBase::getRegisterTypeForCallingConv(), llvm::MVT::getStoreSize(), I, llvm::SmallVectorTemplateBase< T, bool >::push_back(), llvm::seq(), and llvm::zip().

Referenced by llvm::NVPTXTargetLowering::LowerCall(), llvm::NVPTXTargetLowering::LowerFormalArguments(), and llvm::NVPTXTargetLowering::LowerReturn().

◆ convertMLOADToLoadWithUsedBytesMask()

std::pair< MemSDNode *, uint32_t > convertMLOADToLoadWithUsedBytesMask ( MemSDNode * N,
SelectionDAG & DAG,
const NVPTXSubtarget & STI )
static

Definition at line 3605 of file NVPTXISelLowering.cpp.

References assert(), llvm::cast(), DL, llvm::SelectionDAG::getLoad(), llvm::SDValue::getNode(), llvm::SDValue::getOpcode(), llvm::EVT::getSizeInBits(), llvm::EVT::getVectorElementType(), llvm::NVPTXSubtarget::hasUsedBytesMaskPragma(), llvm::EVT::isVector(), N, llvm::ISD::POISON, llvm::reverse(), and llvm::ISD::UNDEF.

Referenced by replaceLoadVector().

◆ correctParamType()

SDValue correctParamType ( SDValue V,
EVT ExpectedVT,
ISD::ArgFlagsTy Flags,
SelectionDAG & DAG,
SDLoc dl )
static

Definition at line 1358 of file NVPTXISelLowering.cpp.

References assert(), llvm::EVT::bitsGT(), llvm::EVT::bitsLT(), getExtOpcode(), llvm::SelectionDAG::getNode(), llvm::EVT::isInteger(), and llvm::ISD::TRUNCATE.

Referenced by llvm::NVPTXTargetLowering::LowerCall(), llvm::NVPTXTargetLowering::LowerFormalArguments(), and llvm::NVPTXTargetLowering::LowerReturn().

◆ expandFSH64()

SDValue expandFSH64 ( SDValue A,
SDValue B,
SDValue ShiftAmount,
SDLoc DL,
unsigned Opcode,
SelectionDAG & DAG )
static

Definition at line 3201 of file NVPTXISelLowering.cpp.

References A(), assert(), B(), DL, llvm::dyn_cast(), llvm::ISD::FSHL, llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getNode(), llvm::SDValue::getValue(), High, llvm::Low, SDValue(), and llvm::NVPTXISD::UNPACK_VECTOR.

Referenced by lowerFSH(), and lowerROT().

◆ getBuildVectorizedValue()

template<typename T>
SDValue getBuildVectorizedValue ( unsigned N,
const SDLoc & dl,
SelectionDAG & DAG,
T GetElement )
inlinestatic

Definition at line 369 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::ExtractVectorElements(), llvm::SelectionDAG::getBuildVector(), llvm::SelectionDAG::getContext(), llvm::SDValue::getValueType(), llvm::EVT::getVectorVT(), I, llvm::EVT::isVector(), N, llvm::SmallVectorTemplateBase< T, bool >::push_back(), llvm::seq(), llvm::SmallVectorTemplateCommon< T, typename >::size(), and T.

Referenced by llvm::NVPTXTargetLowering::LowerCall(), and llvm::NVPTXTargetLowering::LowerReturn().

◆ getExtOpcode()

ISD::NodeType getExtOpcode ( const ISD::ArgFlagsTy & Flags)
static

Definition at line 1350 of file NVPTXISelLowering.cpp.

References llvm::ISD::ANY_EXTEND, llvm::ISD::SIGN_EXTEND, and llvm::ISD::ZERO_EXTEND.

Referenced by correctParamType().

◆ getExtractVectorizedValue()

SDValue getExtractVectorizedValue ( SDValue V,
unsigned I,
EVT VT,
const SDLoc & dl,
SelectionDAG & DAG )
static

Definition at line 352 of file NVPTXISelLowering.cpp.

References assert(), llvm::ISD::EXTRACT_SUBVECTOR, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getNode(), llvm::SelectionDAG::getVectorIdxConstant(), llvm::EVT::getVectorNumElements(), I, and llvm::EVT::isVector().

Referenced by llvm::NVPTXTargetLowering::LowerCall(), and llvm::NVPTXTargetLowering::LowerFormalArguments().

◆ getMinMax3Opcode()

unsigned getMinMax3Opcode ( unsigned MinMax2Opcode)
static

Get 3-input version of a 2-input min/max opcode.

Definition at line 6090 of file NVPTXISelLowering.cpp.

References llvm::ISD::FMAXIMUM, llvm::ISD::FMAXIMUMNUM, llvm::ISD::FMAXNUM, llvm::ISD::FMINIMUM, llvm::ISD::FMINIMUMNUM, llvm::ISD::FMINNUM, and llvm_unreachable.

Referenced by PerformFMinMaxCombine().

◆ getPRMT() [1/2]

SDValue getPRMT ( SDValue A,
SDValue B,
SDValue Selector,
SDLoc DL,
SelectionDAG & DAG,
unsigned Mode = NVPTX::PTXPrmtMode::NONE )
static

Definition at line 1870 of file NVPTXISelLowering.cpp.

References A(), assert(), B(), DL, llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getNode(), llvm::SDValue::getValueType(), Mode, and llvm::NVPTX::PTXPrmtMode::NONE.

Referenced by getPRMT(), lowerBSWAP(), lowerPrmtIntrinsic(), PerformBUILD_VECTORCombine(), and simplifyDemandedBitsForPRMT().

◆ getPRMT() [2/2]

SDValue getPRMT ( SDValue A,
SDValue B,
uint64_t Selector,
SDLoc DL,
SelectionDAG & DAG,
unsigned Mode = NVPTX::PTXPrmtMode::NONE )
static

Definition at line 1879 of file NVPTXISelLowering.cpp.

References A(), B(), DL, llvm::SelectionDAG::getConstant(), getPRMT(), Mode, and llvm::NVPTX::PTXPrmtMode::NONE.

◆ getPRMTDemandedBits()

std::pair< APInt, APInt > getPRMTDemandedBits ( const APInt & SelectorVal,
const APInt & DemandedBits )
static

Definition at line 7483 of file NVPTXISelLowering.cpp.

References llvm::APInt::extractBits(), llvm::APInt::getHiBits(), llvm::APInt::getLoBits(), llvm::APInt::getZExtValue(), I, llvm::seq(), and llvm::APInt::setBit().

Referenced by simplifyDemandedBitsForPRMT().

◆ getPRMTSelector()

APInt getPRMTSelector ( const APInt & Selector,
unsigned Mode )
static

Definition at line 6705 of file NVPTXISelLowering.cpp.

References assert(), llvm::NVPTX::PTXPrmtMode::B4E, llvm::NVPTX::PTXPrmtMode::ECL, llvm::NVPTX::PTXPrmtMode::ECR, llvm::NVPTX::PTXPrmtMode::F4E, llvm::APInt::getBitWidth(), llvm::APInt::getZExtValue(), llvm_unreachable, Mode, llvm::NVPTX::PTXPrmtMode::NONE, llvm::NVPTX::PTXPrmtMode::RC16, llvm::NVPTX::PTXPrmtMode::RC8, S1, and llvm::APInt::trunc().

Referenced by computeKnownBitsForPRMT(), computePRMT(), and simplifyDemandedBitsForPRMT().

◆ getScalar3OpcodeForReduction()

std::optional< unsigned > getScalar3OpcodeForReduction ( unsigned ReductionOpcode)
static

Get 3-input scalar reduction opcode.

Definition at line 1954 of file NVPTXISelLowering.cpp.

References llvm::ISD::VECREDUCE_FMAX, llvm::ISD::VECREDUCE_FMAXIMUM, llvm::ISD::VECREDUCE_FMIN, and llvm::ISD::VECREDUCE_FMINIMUM.

◆ getScalarOpcodeForReduction()

ISD::NodeType getScalarOpcodeForReduction ( unsigned ReductionOpcode)
static

Definition at line 1937 of file NVPTXISelLowering.cpp.

References llvm::ISD::FMAXIMUM, llvm::ISD::FMAXNUM, llvm::ISD::FMINIMUM, llvm::ISD::FMINNUM, llvm_unreachable, llvm::ISD::VECREDUCE_FMAX, llvm::ISD::VECREDUCE_FMAXIMUM, llvm::ISD::VECREDUCE_FMIN, and llvm::ISD::VECREDUCE_FMINIMUM.

◆ getTcgen05LdRedID()

unsigned getTcgen05LdRedID ( Intrinsic::ID IID)
static

Definition at line 3016 of file NVPTXISelLowering.cpp.

References F32, llvm_unreachable, TCGEN05_LD_RED_INST, and TCGEN05_LD_RED_INTR.

Referenced by lowerTcgen05LdRed().

◆ getTcgen05MMADisableOutputLane()

unsigned getTcgen05MMADisableOutputLane ( unsigned IID)
static

Definition at line 2617 of file NVPTXISelLowering.cpp.

References llvm_unreachable.

Referenced by LowerTcgen05MMADisableOutputLane().

◆ getVectorizedVT()

EVT getVectorizedVT ( EVT VT,
unsigned N,
LLVMContext & C )
static

Definition at line 343 of file NVPTXISelLowering.cpp.

References llvm::CallingConv::C, llvm::EVT::getScalarType(), llvm::EVT::getVectorNumElements(), llvm::EVT::getVectorVT(), llvm::EVT::isVector(), and N.

Referenced by llvm::NVPTXTargetLowering::LowerCall(), and llvm::NVPTXTargetLowering::LowerFormalArguments().

◆ getVectorLoweringShape()

std::optional< std::pair< unsigned int, MVT > > getVectorLoweringShape ( EVT VectorEVT,
const NVPTXSubtarget & STI,
unsigned AddressSpace )
static

Definition at line 201 of file NVPTXISelLowering.cpp.

References llvm::EVT::getSimpleVT(), llvm::EVT::getSizeInBits(), llvm::MVT::getSizeInBits(), llvm::MVT::getVectorElementType(), llvm::MVT::getVectorNumElements(), llvm::MVT::getVectorVT(), llvm::NVPTXSubtarget::has256BitVectorLoadStore(), llvm::NVPTXSubtarget::hasF32x2Instructions(), llvm::EVT::isScalarInteger(), llvm::EVT::isSimple(), llvm::MVT::isVector(), and llvm::MVT::SimpleTy.

Referenced by lowerSTOREVector(), and replaceLoadVector().

◆ isConstOne()

bool isConstOne ( const SDValue & Operand)
static

Definition at line 6344 of file NVPTXISelLowering.cpp.

References llvm::dyn_cast().

Referenced by combineMulSelectConstOne(), and matchMADConstOnePattern().

◆ isConstZero()

bool isConstZero ( const SDValue & Operand)
static

Definition at line 5682 of file NVPTXISelLowering.cpp.

References llvm::dyn_cast().

Referenced by PerformADDCombineWithOperands().

◆ IsMulWideOperandDemotable()

bool IsMulWideOperandDemotable ( SDValue Op,
unsigned OptSize,
OperandSignedness & S )
static

IsMulWideOperandDemotable - Checks if the provided DAG node is an operand that can be demoted to OptSize bits without loss of information.

The signedness of the operand, if determinable, is placed in S.

Definition at line 6215 of file NVPTXISelLowering.cpp.

References llvm::EVT::getFixedSizeInBits(), llvm::ISD::SIGN_EXTEND, llvm::ISD::SIGN_EXTEND_INREG, Signed, Unknown, Unsigned, and llvm::ISD::ZERO_EXTEND.

Referenced by AreMulWideOperandsDemotable().

◆ IsPTXVectorType()

bool IsPTXVectorType ( MVT VT)
static

Definition at line 156 of file NVPTXISelLowering.cpp.

References llvm::MVT::SimpleTy.

Referenced by llvm::NVPTXTargetLowering::NVPTXTargetLowering(), and PerformEXTRACTCombine().

◆ lowerBSWAP()

SDValue lowerBSWAP ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2579 of file NVPTXISelLowering.cpp.

References llvm::ISD::ANY_EXTEND, llvm::ISD::BITCAST, DL, llvm::SelectionDAG::getBitcast(), llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getNode(), getPRMT(), llvm::EVT::getSimpleVT(), llvm::SDValue::getValue(), llvm_unreachable, llvm::MVT::SimpleTy, llvm::ISD::TRUNCATE, and llvm::NVPTXISD::UNPACK_VECTOR.

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ LowerClusterLaunchControlQueryCancel()

SDValue LowerClusterLaunchControlQueryCancel ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2886 of file NVPTXISelLowering.cpp.

References llvm::ISD::BITCAST, llvm::cast(), DL, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getNode(), llvm_unreachable, N, and SDValue().

Referenced by lowerIntrinsicWOChain().

◆ lowerCTLZCTPOP()

SDValue lowerCTLZCTPOP ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3191 of file NVPTXISelLowering.cpp.

References assert(), DL, llvm::SelectionDAG::getNode(), llvm::SDNodeFlags::NonNeg, and llvm::ISD::ZERO_EXTEND.

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerCvtRSIntrinsics()

SDValue lowerCvtRSIntrinsics ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2926 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, DL, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getNode(), llvm_unreachable, and N.

Referenced by lowerIntrinsicWOChain().

◆ lowerFREM()

SDValue lowerFREM ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3256 of file NVPTXISelLowering.cpp.

References llvm::SDNodeFlags::AllowContract, DL, llvm::ISD::FABS, llvm::ISD::FDIV, llvm::ISD::FMUL, llvm::ISD::FSUB, llvm::ISD::FTRUNC, llvm::SelectionDAG::getConstantFP(), llvm::APFloat::getInf(), llvm::SelectionDAG::getNode(), llvm::SelectionDAG::getSelect(), llvm::SelectionDAG::getSetCC(), Mul, llvm::ISD::SETEQ, llvm::Sub, X, and Y.

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerFSH()

SDValue lowerFSH ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3245 of file NVPTXISelLowering.cpp.

References expandFSH64().

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerIntrinsicVoid()

SDValue lowerIntrinsicVoid ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2802 of file NVPTXISelLowering.cpp.

References llvm::cast(), llvm::SDValue::getNode(), LowerTcgen05MMADisableOutputLane(), lowerTcgen05St(), lowerTensormapReplaceElemtype(), lowerTensormapReplaceSwizzleMode(), and N.

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerIntrinsicWChain()

SDValue lowerIntrinsicWChain ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3125 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::getMergeValues(), lowerTcgen05Ld(), lowerTcgen05LdRed(), and SDValue().

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerIntrinsicWOChain()

SDValue lowerIntrinsicWOChain ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3155 of file NVPTXISelLowering.cpp.

References LowerClusterLaunchControlQueryCancel(), lowerCvtRSIntrinsics(), and lowerPrmtIntrinsic().

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerLOADi1()

SDValue lowerLOADi1 ( LoadSDNode * LD,
SelectionDAG & DAG )
static

Definition at line 3779 of file NVPTXISelLowering.cpp.

References assert(), llvm::SelectionDAG::getExtLoad(), llvm::SelectionDAG::getMergeValues(), llvm::SelectionDAG::getNode(), llvm::ISD::NON_EXTLOAD, llvm::ISD::TRUNCATE, and llvm::ISD::ZEXTLOAD.

◆ lowerLoadVector()

SDValue lowerLoadVector ( SDNode * N,
SelectionDAG & DAG,
const NVPTXSubtarget & STI )
static

Definition at line 3768 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::getMergeValues(), N, replaceLoadVector(), and SDValue().

Referenced by combineLOAD().

◆ lowerMSTORE()

SDValue lowerMSTORE ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3318 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, assert(), llvm::ISD::BUILD_VECTOR, llvm::cast(), DL, llvm::enumerate(), llvm::ISD::EXTRACT_VECTOR_ELT, llvm::MemSDNode::getAlign(), llvm::SelectionDAG::getEVTAlign(), llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMemIntrinsicNode(), llvm::MemSDNode::getMemOperand(), llvm::MemSDNode::getMemoryVT(), llvm::SelectionDAG::getNode(), llvm::SelectionDAG::getRegister(), llvm::EVT::getSimpleVT(), llvm::SDValue::getValueType(), llvm::EVT::getVectorElementType(), llvm::EVT::getVectorNumElements(), llvm::SelectionDAG::getVTList(), I, llvm::EVT::isVector(), llvm_unreachable, N, llvm::MCRegister::NoRegister, llvm::Offset, llvm::MVT::SimpleTy, llvm::NVPTXISD::StoreV4, llvm::NVPTXISD::StoreV8, and llvm::ISD::UNDEF.

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerPrmtIntrinsic()

SDValue lowerPrmtIntrinsic ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2981 of file NVPTXISelLowering.cpp.

References A(), B(), llvm::NVPTX::PTXPrmtMode::B4E, DL, llvm::NVPTX::PTXPrmtMode::ECL, llvm::NVPTX::PTXPrmtMode::ECR, llvm::NVPTX::PTXPrmtMode::F4E, llvm::SelectionDAG::getConstant(), getPRMT(), llvm_unreachable, Mode, llvm::NVPTX::PTXPrmtMode::NONE, llvm::NVPTX::PTXPrmtMode::RC16, and llvm::NVPTX::PTXPrmtMode::RC8.

Referenced by lowerIntrinsicWOChain().

◆ lowerROT()

SDValue lowerROT ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3250 of file NVPTXISelLowering.cpp.

References expandFSH64(), llvm::ISD::FSHL, llvm::ISD::FSHR, and llvm::ISD::ROTL.

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerSELECT()

SDValue lowerSELECT ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 3284 of file NVPTXISelLowering.cpp.

References llvm::ISD::AND, assert(), Cond, DL, llvm::SelectionDAG::getAnyExtOrTrunc(), llvm::SelectionDAG::getFreeze(), llvm::SelectionDAG::getNode(), llvm::SelectionDAG::getNOT(), llvm::SelectionDAG::getSelect(), llvm::ISD::OR, Select, and llvm::ISD::TRUNCATE.

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ lowerSTOREVector()

SDValue lowerSTOREVector ( SDValue Op,
SelectionDAG & DAG,
const NVPTXSubtarget & STI )
static

Definition at line 3850 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::ISD::ANY_EXTEND, assert(), llvm::cast(), DL, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::ExtractVectorElements(), llvm::SelectionDAG::getBitcast(), llvm::SelectionDAG::getBuildVector(), llvm::SelectionDAG::getContext(), llvm::SelectionDAG::getDataLayout(), llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMemIntrinsicNode(), llvm::SelectionDAG::getNode(), llvm::DataLayout::getPrefTypeAlign(), llvm::EVT::getSizeInBits(), llvm::EVT::getTypeForEVT(), llvm::SDValue::getValueType(), llvm::EVT::getVectorElementType(), getVectorLoweringShape(), llvm::EVT::getVectorNumElements(), llvm::MVT::getVectorVT(), llvm::SelectionDAG::getVTList(), I, llvm::EVT::isVector(), N, SDValue(), llvm::seq(), llvm::NVPTXISD::StoreV2, llvm::NVPTXISD::StoreV4, and llvm::NVPTXISD::StoreV8.

Referenced by combineSTORE().

◆ lowerTcgen05Ld()

std::optional< std::pair< SDValue, SDValue > > lowerTcgen05Ld ( SDNode * N,
SelectionDAG & DAG,
bool HasOffset = false )
static

Definition at line 2708 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::ISD::BUILD_VECTOR, llvm::cast(), DL, llvm::SelectionDAG::getMemIntrinsicNode(), llvm::MemSDNode::getMemOperand(), llvm::MemSDNode::getMemoryVT(), llvm::SelectionDAG::getNode(), llvm::SDValue::getValue(), llvm::EVT::getVectorNumElements(), llvm::SelectionDAG::getVTList(), llvm::ISD::INTRINSIC_W_CHAIN, llvm::EVT::isVector(), N, and llvm::SmallVectorTemplateBase< T, bool >::push_back().

Referenced by lowerIntrinsicWChain(), and ReplaceINTRINSIC_W_CHAIN().

◆ lowerTcgen05LdRed()

std::optional< std::tuple< SDValue, SDValue, SDValue > > lowerTcgen05LdRed ( SDNode * N,
SelectionDAG & DAG )
static

Definition at line 3081 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::ISD::BUILD_VECTOR, llvm::cast(), DL, llvm::SelectionDAG::getMemIntrinsicNode(), llvm::MemSDNode::getMemOperand(), llvm::MemSDNode::getMemoryVT(), llvm::SelectionDAG::getNode(), getTcgen05LdRedID(), llvm::SDValue::getValue(), llvm::EVT::getVectorElementType(), llvm::EVT::getVectorNumElements(), llvm::SelectionDAG::getVTList(), llvm::EVT::isFloatingPoint(), llvm::EVT::isVector(), N, and llvm::SmallVectorTemplateBase< T, bool >::push_back().

Referenced by lowerIntrinsicWChain(), and ReplaceINTRINSIC_W_CHAIN().

◆ LowerTcgen05MMADisableOutputLane()

SDValue LowerTcgen05MMADisableOutputLane ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2677 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::cast(), DL, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMemIntrinsicNode(), llvm::MemSDNode::getMemOperand(), llvm::MemSDNode::getMemoryVT(), llvm::SelectionDAG::getNode(), getTcgen05MMADisableOutputLane(), llvm::SDValue::getValueType(), llvm::EVT::getVectorElementType(), llvm::EVT::getVectorNumElements(), I, llvm::EVT::isVector(), and N.

Referenced by lowerIntrinsicVoid().

◆ lowerTcgen05St()

SDValue lowerTcgen05St ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2553 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::cast(), DL, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMemIntrinsicNode(), llvm::MemSDNode::getMemOperand(), llvm::MemSDNode::getMemoryVT(), llvm::SelectionDAG::getNode(), llvm::SDValue::getValueType(), llvm::EVT::getVectorElementType(), llvm::EVT::getVectorNumElements(), I, llvm::ISD::INTRINSIC_VOID, llvm::EVT::isVector(), and N.

Referenced by lowerIntrinsicVoid().

◆ lowerTensormapReplaceElemtype()

SDValue lowerTensormapReplaceElemtype ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2774 of file NVPTXISelLowering.cpp.

References DL, llvm::SelectionDAG::getSubtarget(), llvm::NVPTXSubtarget::hasTensormapReplaceElemtypeSupport(), N, and reportInvalidTensormapReplaceUsage().

Referenced by lowerIntrinsicVoid().

◆ lowerTensormapReplaceSwizzleMode()

SDValue lowerTensormapReplaceSwizzleMode ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2788 of file NVPTXISelLowering.cpp.

References DL, llvm::SelectionDAG::getSubtarget(), llvm::NVPTXSubtarget::hasTensormapReplaceSwizzleModeSupport(), N, and reportInvalidTensormapReplaceUsage().

Referenced by lowerIntrinsicVoid().

◆ LowerVectorArith()

SDValue LowerVectorArith ( SDValue Op,
SelectionDAG & DAG )
static

Definition at line 2533 of file NVPTXISelLowering.cpp.

References llvm::ISD::BUILD_VECTOR, DL, E(), llvm::SelectionDAG::getNode(), I, llvm::SmallVectorTemplateBase< T, bool >::push_back(), and llvm::transform().

Referenced by llvm::NVPTXTargetLowering::LowerOperation().

◆ matchMADConstOnePattern()

SDValue matchMADConstOnePattern ( SDValue Add)
static

Definition at line 6349 of file NVPTXISelLowering.cpp.

References llvm::ISD::ADD, llvm::Add, isConstOne(), and SDValue().

Referenced by combineMADConstOne(), and combineMulSelectConstOne().

◆ PerformADDCombine()

SDValue PerformADDCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

PerformADDCombine - Target-specific dag combine xforms for ISD::ADD.

Definition at line 6047 of file NVPTXISelLowering.cpp.

References llvm::SDValue::getValueType(), llvm::EVT::isVector(), N, llvm::None, PerformADDCombineWithOperands(), and SDValue().

◆ PerformADDCombineWithOperands()

SDValue PerformADDCombineWithOperands ( SDNode * N,
SDValue N0,
SDValue N1,
TargetLowering::DAGCombinerInfo & DCI )
static

PerformADDCombineWithOperands - Try DAG combinations for an ADD with operands N0 and N1.

This is a helper for PerformADDCombine that is called with the default operands, and if that fails, with commuted operands.

Definition at line 5692 of file NVPTXISelLowering.cpp.

References llvm::ISD::ADD, llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::SDValue::getNode(), llvm::SelectionDAG::getNode(), llvm::SDValue::getOpcode(), llvm::SDNode::getOperand(), llvm::SelectionDAG::getSelect(), llvm::SDValue::getValueType(), llvm::SDNode::hasOneUse(), isConstZero(), llvm::ISD::MUL, Mul, N, SDValue(), and llvm::ISD::SELECT.

◆ PerformBUILD_VECTORCombine()

SDValue PerformBUILD_VECTORCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6625 of file NVPTXISelLowering.cpp.

References assert(), llvm::ISD::BITCAST, llvm::cast(), llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::SelectionDAG::getBitcast(), getPRMT(), llvm::EVT::getVectorNumElements(), llvm::EVT::is32BitVector(), llvm::isa(), llvm::TargetLowering::DAGCombinerInfo::isAfterLegalizeDAG(), llvm::NVPTX::isPackedVectorTy(), N, SDValue(), llvm::ISD::SRL, and llvm::ISD::TRUNCATE.

◆ PerformEXTRACTCombine()

SDValue PerformEXTRACTCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6485 of file NVPTXISelLowering.cpp.

References llvm::ISD::allOperandsUndef(), llvm::ISD::ANY_EXTEND, llvm::ISD::BITCAST, llvm::EVT::changeTypeToInteger(), llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::dyn_cast(), llvm::ISD::FREEZE, llvm::SelectionDAG::getConstant(), llvm::MVT::getIntegerVT(), llvm::SelectionDAG::getNode(), llvm::EVT::getScalarSizeInBits(), llvm::EVT::getSimpleVT(), llvm::EVT::getSizeInBits(), llvm::EVT::getVectorElementType(), llvm::EVT::getVectorNumElements(), llvm::NVPTX::isPackedVectorTy(), IsPTXVectorType(), llvm::EVT::isSimple(), llvm::ISD::LOAD, N, SDValue(), llvm::ISD::SRA, llvm::ISD::TRUNCATE, and llvm::Vector.

◆ PerformFADDCombine()

SDValue PerformFADDCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

PerformFADDCombine - Target-specific dag combine xforms for ISD::FADD.

Definition at line 6071 of file NVPTXISelLowering.cpp.

References llvm::SDValue::getValueType(), llvm::EVT::isVector(), N, PerformFADDCombineWithOperands(), and SDValue().

◆ PerformFADDCombineWithOperands()

SDValue PerformFADDCombineWithOperands ( SDNode * N,
SDValue N0,
SDValue N1,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

Definition at line 5732 of file NVPTXISelLowering.cpp.

References llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::ISD::FADD, llvm::ISD::FMA, llvm::ISD::FMUL, llvm::SDNode::getFlags(), llvm::SDNode::getIROrder(), llvm::SelectionDAG::getMachineFunction(), llvm::SDValue::getNode(), llvm::SelectionDAG::getNode(), llvm::SDValue::getOpcode(), llvm::SDValue::getOperand(), llvm::SelectionDAG::getTargetLoweringInfo(), llvm::SDValue::getValueType(), llvm::SDNodeFlags::hasAllowContract(), llvm::isa(), N, SDValue(), and llvm::SDNode::users().

Referenced by PerformFADDCombine().

◆ PerformFMinMaxCombine()

SDValue PerformFMinMaxCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
unsigned PTXVersion,
unsigned SmVersion )
static

PerformFMinMaxCombine - Combine (fmaxnum (fmaxnum a, b), c) into (fmaxnum3 a, b, c).

Also covers other llvm min/max intrinsics.

Definition at line 6109 of file NVPTXISelLowering.cpp.

References A(), B(), llvm::CallingConv::C, llvm::TargetLowering::DAGCombinerInfo::DAG, getMinMax3Opcode(), llvm::SelectionDAG::getNode(), llvm::SDValue::getOpcode(), llvm::SDValue::getOperand(), llvm::SDValue::hasOneUse(), N, and SDValue().

◆ PerformMULCombine()

SDValue PerformMULCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

PerformMULCombine - Runs PTX-specific DAG combine patterns on MUL nodes.

Definition at line 6431 of file NVPTXISelLowering.cpp.

References N, llvm::None, PerformMULCombineWithOperands(), SDValue(), and TryMULWIDECombine().

◆ PerformMULCombineWithOperands()

SDValue PerformMULCombineWithOperands ( SDNode * N,
SDValue N0,
SDValue N1,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6403 of file NVPTXISelLowering.cpp.

References combineMADConstOne(), combineMulSelectConstOne(), DL, llvm::SDValue::getValueType(), llvm::EVT::isVector(), N, and SDValue().

Referenced by PerformMULCombine().

◆ PerformREMCombine()

SDValue PerformREMCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

Definition at line 6139 of file NVPTXISelLowering.cpp.

References assert(), llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::Default, DL, llvm::SelectionDAG::getNode(), llvm::SDNode::getOperand(), llvm::ISD::MUL, N, llvm::ISD::SDIV, SDValue(), llvm::ISD::SREM, llvm::ISD::SUB, llvm::ISD::UDIV, llvm::ISD::UREM, and llvm::SDNode::users().

◆ PerformSELECTShiftCombine()

SDValue PerformSELECTShiftCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

Transform patterns like: (select (ugt shift_amt, BitWidth-1), 0, (srl/shl x, shift_amt)) (select (ult shift_amt, BitWidth), (srl/shl x, shift_amt), 0) Into: (NVPTXISD::SRL_CLAMP x, shift_amt) or (NVPTXISD::SHL_CLAMP x, shift_amt)

These patterns arise from C/C++ code like shift >= 32 ? 0 : x >> shift which guards against undefined behavior. PTX shr/shl instructions clamp shift amounts >= BitWidth to produce 0 for logical shifts, making the guard redundant.

Note: We only handle SRL and SHL, not SRA, because arithmetic right shifts could produce 0 or -1 when shift >= BitWidth. Note: We don't handle uge or ule. These don't appear because of canonicalization.

Definition at line 6552 of file NVPTXISelLowering.cpp.

References llvm::BitWidth, llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::SelectionDAG::getNode(), llvm::SDValue::getOpcode(), llvm::SDValue::getOperand(), llvm::SDValue::getValueType(), llvm::TargetLowering::DAGCombinerInfo::isAfterLegalizeDAG(), llvm::PatternMatch::m_Deferred(), llvm::PatternMatch::m_Select(), llvm::PatternMatch::m_Shl(), llvm::PatternMatch::m_SpecificInt(), llvm::PatternMatch::m_TruncOrSelf(), llvm::PatternMatch::m_Value(), llvm::PatternMatch::m_Zero(), N, SDValue(), llvm::ISD::SETUGT, llvm::ISD::SETULT, and llvm::ISD::SRL.

◆ PerformSETCCCombine()

SDValue PerformSETCCCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
unsigned int SmVersion )
static

Definition at line 6458 of file NVPTXISelLowering.cpp.

References A(), B(), llvm::ISD::BUILD_VECTOR, llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::SelectionDAG::getNode(), llvm::SDValue::getValue(), llvm::SelectionDAG::getVTList(), N, SDValue(), llvm::NVPTXISD::SETP_BF16X2, and llvm::NVPTXISD::SETP_F16X2.

◆ PerformSHLCombine()

SDValue PerformSHLCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI,
CodeGenOptLevel OptLevel )
static

PerformSHLCombine - Runs PTX-specific DAG combine patterns on SHL nodes.

Definition at line 6446 of file NVPTXISelLowering.cpp.

References N, llvm::None, SDValue(), and TryMULWIDECombine().

◆ PerformVSELECTCombine()

SDValue PerformVSELECTCombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6592 of file NVPTXISelLowering.cpp.

References llvm::ISD::BUILD_VECTOR, llvm::CallingConv::C, llvm::TargetLowering::DAGCombinerInfo::DAG, DL, E(), llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getAnyExtOrTrunc(), llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getNode(), llvm::SDValue::getValueType(), I, N, SDValue(), and llvm::ISD::SELECT.

◆ PromoteBinOpToF32()

SDValue PromoteBinOpToF32 ( SDNode * N,
SelectionDAG & DAG )
static

Definition at line 2419 of file NVPTXISelLowering.cpp.

References DL, llvm::SelectionDAG::getContext(), llvm::SelectionDAG::getFPExtendOrRound(), llvm::SelectionDAG::getNode(), llvm::EVT::getVectorElementCount(), llvm::EVT::getVectorVT(), llvm::EVT::isVector(), and N.

◆ promoteScalarIntegerPTX()

EVT promoteScalarIntegerPTX ( const EVT VT)
static

PromoteScalarIntegerPTX Used to make sure the arguments/returns are suitable for passing and promote them to a larger size if they're not.

The promoted type is placed in PromoteVT if the function returns true.

Definition at line 393 of file NVPTXISelLowering.cpp.

References llvm::EVT::getFixedSizeInBits(), llvm::EVT::isScalarInteger(), llvm_unreachable, and llvm::PowerOf2Ceil().

Referenced by llvm::NVPTXTargetLowering::LowerCall(), and llvm::NVPTXTargetLowering::LowerReturn().

◆ refinePtrAS()

MachinePointerInfo refinePtrAS ( SDValue & Ptr,
SelectionDAG & DAG,
const DataLayout & DL,
const TargetLowering & TL )
static

Definition at line 1326 of file NVPTXISelLowering.cpp.

References llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, llvm::NVPTXAS::ADDRESS_SPACE_LOCAL, llvm::ISD::ADDRSPACECAST, llvm::cast(), DL, llvm::ISD::FrameIndex, llvm::SelectionDAG::getAddrSpaceCast(), llvm::SDNode::getOpcode(), llvm::SDNode::getOperand(), and llvm::TargetLoweringBase::getPointerTy().

Referenced by llvm::NVPTXTargetLowering::LowerCall().

◆ replaceAtomicSwap128()

void replaceAtomicSwap128 ( SDNode * N,
SelectionDAG & DAG,
const NVPTXSubtarget & STI,
SmallVectorImpl< SDValue > & Results )
static

Definition at line 7156 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, assert(), llvm::NVPTXISD::ATOMIC_CMP_SWAP_B128, llvm::ISD::ATOMIC_SWAP, llvm::NVPTXISD::ATOMIC_SWAP_B128, llvm::ISD::BUILD_PAIR, llvm::cast(), llvm::LLVMContext::diagnose(), llvm::ISD::EXTRACT_ELEMENT, llvm::SelectionDAG::getContext(), llvm::SDLoc::getDebugLoc(), llvm::MachineFunction::getFunction(), llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMachineFunction(), llvm::SelectionDAG::getMemIntrinsicNode(), llvm::MemSDNode::getMemOperand(), llvm::SelectionDAG::getNode(), llvm::SDNode::getOperand(), llvm::SelectionDAG::getUNDEF(), llvm::SelectionDAG::getVTList(), llvm::NVPTXSubtarget::hasAtomSwap128(), N, llvm::SDNode::ops(), and Results.

◆ ReplaceBITCAST()

void ReplaceBITCAST ( SDNode * Node,
SelectionDAG & DAG,
SmallVectorImpl< SDValue > & Results )
static

Definition at line 6911 of file NVPTXISelLowering.cpp.

References llvm::ISD::BUILD_VECTOR, DL, llvm::SelectionDAG::getBitcast(), llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getNode(), Results, llvm::ISD::SRL, and llvm::ISD::TRUNCATE.

◆ ReplaceCopyFromReg_128()

void ReplaceCopyFromReg_128 ( SDNode * N,
SelectionDAG & DAG,
SmallVectorImpl< SDValue > & Results )
static

Definition at line 7116 of file NVPTXISelLowering.cpp.

References assert(), llvm::ISD::BUILD_PAIR, llvm::ISD::CopyFromReg, DL, llvm::SelectionDAG::getNode(), llvm::SDValue::getValue(), N, Reg, and Results.

◆ ReplaceINTRINSIC_W_CHAIN()

void ReplaceINTRINSIC_W_CHAIN ( SDNode * N,
SelectionDAG & DAG,
SmallVectorImpl< SDValue > & Results )
static

Definition at line 6933 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::SmallVectorImpl< T >::append(), assert(), llvm::cast(), DL, llvm::SDNode::getAsZExtVal(), llvm::SelectionDAG::getBuildVector(), llvm::SelectionDAG::getMemIntrinsicNode(), llvm::MemSDNode::getMemOperand(), llvm::MemSDNode::getMemoryVT(), llvm::SDValue::getNode(), llvm::SelectionDAG::getNode(), llvm::EVT::getSimpleVT(), llvm::EVT::getSizeInBits(), llvm::SDValue::getValue(), llvm::EVT::getVectorElementType(), llvm::EVT::getVectorNumElements(), llvm::SelectionDAG::getVTList(), llvm::ISD::INTRINSIC_W_CHAIN, llvm::EVT::isSimple(), llvm::EVT::isVector(), llvm::NVPTXISD::LDUV2, llvm::NVPTXISD::LDUV4, lowerTcgen05Ld(), lowerTcgen05LdRed(), N, llvm::SmallVectorTemplateBase< T, bool >::push_back(), Results, llvm::MVT::SimpleTy, and llvm::ISD::TRUNCATE.

◆ replaceLoadVector() [1/2]

std::optional< std::pair< SDValue, SDValue > > replaceLoadVector ( SDNode * N,
SelectionDAG & DAG,
const NVPTXSubtarget & STI )
static

replaceLoadVector - Convert vector loads into multi-output scalar loads.

Definition at line 3658 of file NVPTXISelLowering.cpp.

References assert(), llvm::cast(), convertMLOADToLoadWithUsedBytesMask(), DL, llvm::SelectionDAG::ExtractVectorElements(), llvm::SelectionDAG::getBitcast(), llvm::SelectionDAG::getBuildVector(), llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getContext(), llvm::SelectionDAG::getDataLayout(), llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMemIntrinsicNode(), llvm::SelectionDAG::getNode(), llvm::EVT::getScalarType(), llvm::EVT::getSizeInBits(), llvm::EVT::getTypeForEVT(), llvm::SDValue::getValue(), llvm::EVT::getVectorElementType(), getVectorLoweringShape(), llvm::EVT::getVectorNumElements(), llvm::MVT::getVectorVT(), llvm::SelectionDAG::getVTList(), I, llvm::EVT::isVector(), llvm::NVPTXISD::LoadV2, llvm::NVPTXISD::LoadV4, llvm::NVPTXISD::LoadV8, llvm::ISD::MLOAD, N, llvm::SmallVectorTemplateBase< T, bool >::push_back(), llvm::seq(), llvm::SmallVectorTemplateCommon< T, typename >::size(), llvm::SmallVector, and llvm::ISD::TRUNCATE.

Referenced by lowerLoadVector(), and replaceLoadVector().

◆ replaceLoadVector() [2/2]

void replaceLoadVector ( SDNode * N,
SelectionDAG & DAG,
SmallVectorImpl< SDValue > & Results,
const NVPTXSubtarget & STI )
static

Definition at line 3761 of file NVPTXISelLowering.cpp.

References N, replaceLoadVector(), and Results.

◆ replaceProxyReg()

void replaceProxyReg ( SDNode * N,
SelectionDAG & DAG,
const TargetLowering & TLI,
SmallVectorImpl< SDValue > & Results )
static

Definition at line 7140 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::getAnyExtOrTrunc(), llvm::SelectionDAG::getContext(), llvm::SelectionDAG::getNode(), llvm::TargetLoweringBase::getRegisterType(), N, Reg, and Results.

◆ reportInvalidTensormapReplaceUsage()

SDValue reportInvalidTensormapReplaceUsage ( SDValue Op,
SelectionDAG & DAG,
unsigned Val )
static

Definition at line 2751 of file NVPTXISelLowering.cpp.

References llvm::LLVMContext::diagnose(), DL, llvm::dyn_cast(), llvm::PointerType::get(), llvm::SelectionDAG::getContext(), llvm::MachineFunction::getFunction(), llvm::SelectionDAG::getMachineFunction(), llvm::Intrinsic::getName(), llvm::GlobalValue::getParent(), and N.

Referenced by lowerTensormapReplaceElemtype(), and lowerTensormapReplaceSwizzleMode().

◆ shouldConvertToIndirectCall()

bool shouldConvertToIndirectCall ( const CallBase * CB,
const GlobalAddressSDNode * Func )
static

Definition at line 1317 of file NVPTXISelLowering.cpp.

References llvm::dyn_cast(), and llvm::CallBase::getFunctionType().

Referenced by llvm::NVPTXTargetLowering::LowerCall().

◆ simplifyDemandedBitsForPRMT()

SDValue simplifyDemandedBitsForPRMT ( SDValue PRMT,
const APInt & DemandedBits,
SelectionDAG & DAG,
const TargetLowering & TLI,
unsigned Depth )
static

Definition at line 7517 of file NVPTXISelLowering.cpp.

References assert(), canonicalizePRMTInput(), llvm::Depth, llvm::dyn_cast(), llvm::SDValue::getConstantOperandVal(), llvm::APInt::getLoBits(), llvm::SDValue::getOpcode(), llvm::SDValue::getOperand(), getPRMT(), getPRMTDemandedBits(), getPRMTSelector(), llvm::APInt::getZExtValue(), Mode, SDValue(), and llvm::TargetLowering::SimplifyMultipleUseDemandedBits().

Referenced by llvm::NVPTXTargetLowering::SimplifyDemandedBitsForTargetNode().

◆ sinkProxyReg()

SDValue sinkProxyReg ( SDValue R,
SDValue Chain,
TargetLowering::DAGCombinerInfo & DCI )
static

Definition at line 6783 of file NVPTXISelLowering.cpp.

References A(), AbstractManglingParser< Derived, Alloc >::Ops, llvm::ISD::ANY_EXTEND, B(), llvm::ISD::BITCAST, llvm::ISD::BUILD_VECTOR, llvm::ISD::Constant, llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getNode(), llvm::TargetLowering::DAGCombinerInfo::isBeforeLegalize(), llvm::ISD::LOAD, llvm::NVPTXISD::LoadV2, llvm::NVPTXISD::LoadV4, llvm::ISD::OR, SDValue(), llvm::ISD::SHL, llvm::ISD::SIGN_EXTEND, sinkProxyReg(), llvm::ISD::SRA, llvm::ISD::SRL, llvm::ISD::TRUNCATE, and llvm::ISD::ZERO_EXTEND.

Referenced by combineProxyReg(), and sinkProxyReg().

◆ TryMULWIDECombine()

SDValue TryMULWIDECombine ( SDNode * N,
TargetLowering::DAGCombinerInfo & DCI )
static

TryMULWIDECombine - Attempt to replace a multiply of M bits with a multiply of M/2 bits that produces an M-bit result (i.e.

mul.wide). This transform works on both multiply DAG nodes and SHL DAG nodes with a constant shift amount.

Definition at line 6278 of file NVPTXISelLowering.cpp.

References AreMulWideOperandsDemotable(), llvm::BitWidth, llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::dyn_cast(), llvm::ConstantSDNode::getAPIntValue(), llvm::SelectionDAG::getConstant(), llvm::SelectionDAG::getNode(), llvm::EVT::getSizeInBits(), llvm::isa(), LHS, llvm::ISD::MUL, N, Opc, RHS, SDValue(), llvm::APInt::sge(), llvm::ISD::SHL, Signed, llvm::APInt::slt(), std::swap(), and llvm::ISD::TRUNCATE.

Referenced by PerformMULCombine(), and PerformSHLCombine().

◆ VectorizePTXValueVTs()

template<typename T>
SmallVector< unsigned, 16 > VectorizePTXValueVTs ( const SmallVectorImpl< EVT > & ValueVTs,
const SmallVectorImpl< T > & Offsets,
Align ParamAlignment,
bool IsVAArg = false )
static

Definition at line 481 of file NVPTXISelLowering.cpp.

References assert(), canMergeParamLoadStoresStartingAt(), E(), I, and llvm::SmallVectorTemplateCommon< T, typename >::size().

Referenced by llvm::NVPTXTargetLowering::LowerCall(), llvm::NVPTXTargetLowering::LowerFormalArguments(), and llvm::NVPTXTargetLowering::LowerReturn().

Variable Documentation

◆ FMAContractLevelOpt

cl::opt< unsigned > FMAContractLevelOpt("nvptx-fma-level", cl::Hidden, cl::desc("NVPTX Specific: FMA contraction (0: don't do it" " 1: do it 2: do it aggressively"), cl::init(2)) ( "nvptx-fma-level" ,
cl::Hidden  )
static

Referenced by llvm::NVPTXTargetLowering::allowFMA().

◆ ForceMinByValParamAlign

cl::opt< bool > ForceMinByValParamAlign("nvptx-force-min-byval-param-align", cl::Hidden, cl::desc("NVPTX Specific: force 4-byte minimal alignment for byval" " params of device functions."), cl::init(false)) ( "nvptx-force-min-byval-param-align" ,
cl::Hidden ,
cl::desc("NVPTX Specific: force 4-byte minimal alignment for byval" " params of device functions.") ,
cl::init(false)  )
static

Referenced by llvm::NVPTXTargetLowering::getFunctionByValParamAlign().

◆ sched4reg

cl::opt< bool > sched4reg("nvptx-sched4reg", cl::desc("NVPTX Specific: schedule for register pressue"), cl::init(false)) ( "nvptx-sched4reg" ,
cl::desc("NVPTX Specific: schedule for register pressue") ,
cl::init(false)  )
static

Referenced by llvm::NVPTXTargetLowering::NVPTXTargetLowering().

◆ UseApproxLog2F32

cl::opt< bool > UseApproxLog2F32("nvptx-approx-log2f32", cl::desc("NVPTX Specific: whether to use lg2.approx for log2"), cl::init(false)) ( "nvptx-approx-log2f32" ,
cl::desc("NVPTX Specific: whether to use lg2.approx for log2") ,
cl::init(false)  )
static

Whereas CUDA's implementation (see libdevice) uses ex2.approx for exp2(), it does NOT use lg2.approx for log2, so this is disabled by default.

Referenced by llvm::NVPTXTargetLowering::NVPTXTargetLowering().

◆ UsePrecDivF32

cl::opt< NVPTX::DivPrecisionLevel > UsePrecDivF32("nvptx-prec-divf32", cl::Hidden, cl::desc( "NVPTX Specific: Override the precision of the lowering for f32 fdiv"), cl::values( clEnumValN(NVPTX::DivPrecisionLevel::Approx, "0", "Use div.approx"), clEnumValN(NVPTX::DivPrecisionLevel::Full, "1", "Use div.full"), clEnumValN(NVPTX::DivPrecisionLevel::IEEE754, "2", "Use IEEE Compliant F32 div.rnd if available (default)"), clEnumValN(NVPTX::DivPrecisionLevel::IEEE754_NoFTZ, "3", "Use IEEE Compliant F32 div.rnd if available, no FTZ")), cl::init(NVPTX::DivPrecisionLevel::IEEE754)) ( "nvptx-prec-divf32" ,
cl::Hidden ,
cl::desc( "NVPTX Specific: Override the precision of the lowering for f32 fdiv") ,
cl::values( clEnumValN(NVPTX::DivPrecisionLevel::Approx, "0", "Use div.approx"), clEnumValN(NVPTX::DivPrecisionLevel::Full, "1", "Use div.full"), clEnumValN(NVPTX::DivPrecisionLevel::IEEE754, "2", "Use IEEE Compliant F32 div.rnd if available (default)"), clEnumValN(NVPTX::DivPrecisionLevel::IEEE754_NoFTZ, "3", "Use IEEE Compliant F32 div.rnd if available, no FTZ")) ,
cl::init(NVPTX::DivPrecisionLevel::IEEE754)  )
static

Referenced by llvm::NVPTXTargetLowering::getDivF32Level().

◆ UsePrecSqrtF32

cl::opt< bool > UsePrecSqrtF32("nvptx-prec-sqrtf32", cl::Hidden, cl::desc("NVPTX Specific: 0 use sqrt.approx, 1 use sqrt.rn."), cl::init(true)) ( "nvptx-prec-sqrtf32" ,
cl::Hidden ,
cl::desc("NVPTX Specific: 0 use sqrt.approx, 1 use sqrt.rn.") ,
cl::init(true)  )
static

Referenced by llvm::NVPTXTargetLowering::usePrecSqrtF32().

Generated on for LLVM by doxygen 1.14.0