NVPTXISelLowering.cpp File Reference

#include "NVPTXISelLowering.h"
#include "MCTargetDesc/NVPTXBaseInfo.h"
#include "NVPTX.h"
#include "NVPTXMachineFunctionInfo.h"
#include "NVPTXSelectionDAGInfo.h"
#include "NVPTXSubtarget.h"
#include "NVPTXTargetMachine.h"
#include "NVPTXTargetObjectFile.h"
#include "NVPTXUtilities.h"
#include "NVVMProperties.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/MachineFrameInfo.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/LLVMContext.h"
#include "llvm/IR/Module.h"
#include "llvm/IR/NVVMIntrinsicUtils.h"
#include "llvm/IR/Type.h"
#include "llvm/IR/Value.h"
#include "llvm/MC/MCContext.h"
#include "llvm/MC/MCSymbol.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/Target/TargetMachine.h"
#include "llvm/Target/TargetOptions.h"
#include <algorithm>
#include <cassert>
#include <cmath>
#include <cstdint>
#include <iterator>
#include <optional>
#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)
#define	SPCOMPRESS_NODE_CASE(NumResults)
#define	SPDECOMPRESS_NODE_CASE(NumResults)

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 MachinePointerInfo	refinePtrAS (SDValue &Ptr, SelectionDAG &DAG)
static ISD::NodeType	getExtOpcode (const ISD::ArgFlagsTy &Flags)
static SDValue	correctParamType (SDValue V, EVT ExpectedVT, ISD::ArgFlagsTy Flags, SelectionDAG &DAG, SDLoc dl)
static SDValue	getSymbolNode (SelectionDAG &DAG, MCSymbol *Sym, EVT T)
static SDValue	getSymbolNode (SelectionDAG &DAG, const Twine &Name, EVT T)
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	lowerMUL_LOHI (SDValue Op, SelectionDAG &DAG)
static SDValue	lowerFPTRUNC_ROUND (SDValue Op, SelectionDAG &DAG, const NVPTXSubtarget &STI)
static SDValue	LowerVectorArith (SDValue Op, SelectionDAG &DAG)
static SDValue	lowerTcgen05St (SDValue Op, SelectionDAG &DAG, bool hasOffset=false)
static SDValue	lowerBSWAP (SDValue Op, SelectionDAG &DAG)
static SDValue	lowerStAsyncWithMbarrier (SDValue Op, SelectionDAG &DAG)
static SDValue	lowerStAsyncRelease (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 unsigned	getSPNumRegs (EVT VT)
static EVT	getSPRegType (EVT VT)
static void	appendSplitSPOperands (ArrayRef< SDUse > Input, SmallVectorImpl< SDValue > &Output, SelectionDAG &DAG)
static SDValue	joinSPRegs (EVT VT, ArrayRef< SDValue > Regs, const SDLoc &DL, SelectionDAG &DAG)
static unsigned	getSPCompressNodeOpcode (unsigned NumResults)
static unsigned	getSPDecompressNodeOpcode (unsigned NumResults)
static SDValue	lowerSPIntrinsic (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	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	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 bool	isNonCoalescableBuildVector (const SDValue &BV)
	Check if a v2f32 BUILD_VECTOR provably packs values from non-adjacent register pairs (non-coalescable).
static unsigned	getMinMax3Opcode (unsigned MinMax2Opcode)
	Get 3-input version of a 2-input min/max opcode.
static SDValue	PerformFMinMaxCombine (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const NVPTXSubtarget &STI)
	PerformFMinMaxCombine - Combine (fmaxnum (fmaxnum a, b), c) into (fmaxnum3 a, b, c).
static SDValue	combineSZExtToMulWide (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, CodeGenOptLevel OptLevel)
static bool	IsExtendedFromLte (SDValue Op, unsigned OptSize, bool &IsSigned)
	Check whether Op is explicitly extended from at most OptSize bits, setting IsSigned to the signedness of the extension.
static bool	AreMulWideOperandsDemotable (SDValue LHS, SDValue RHS, unsigned OptSize, const SelectionDAG &DAG, bool &IsSigned)
	AreMulWideOperandsDemotable - Checks if the given LHS and RHS operands can be demoted to OptSize bits without loss of information using the same signedness.
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	combineShiftOfLogicOp (SDNode *N, TargetLowering::DAGCombinerInfo &DCI)
	Commute SHL with a bitwise logic operation when doing so exposes a common shifted operand.
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, const NVPTXSubtarget &STI)
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 unsigned	getFAddWithNegOpcode (EVT VT, Intrinsic::ID IID, APFloat::roundingMode RoundingMode)
static SDValue	combineFAddWithNeg (SDNode *N, SelectionDAG &DAG, Intrinsic::ID AddIntrinsicID, APFloat::roundingMode RoundingMode)
static bool	isSupportedFPArith (SDNode *N, const NVPTXSubtarget &STI, unsigned ISDOpcode, Intrinsic::ID IID, APFloat::roundingMode RoundingMode)
static SDValue	diagnoseUnsupportedFPArith (SDNode *N, SelectionDAG &DAG, Intrinsic::ID IID, APFloat::roundingMode RoundingMode)
static SDValue	combineIntrinsicWOChain (SDNode *N, TargetLowering::DAGCombinerInfo &DCI, const NVPTXSubtarget &STI)
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 TargetLoweringBase::AtomicExpansionKind	getScalarAtomicRMWExpansion (AtomicRMWInst::BinOp Op, Type *Ty, unsigned AddressSpace, const Function &F, bool IgnoreDenormalMode, const NVPTXSubtarget &STI)
	Returns how NVPTX should expand a scalar atomicrmw.
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 >	AllowFTZAtomics ("nvptx-allow-ftz-atomics", cl::Hidden, cl::desc("NVPTX Specific: Lower atomicrmw fadd to atom.add even when its " "FTZ behavior does not match the function's denormal mode."), 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.

Macro Definition Documentation

DEBUG_TYPE

#define DEBUG_TYPE   "nvptx-lower"

Definition at line 82 of file NVPTXISelLowering.cpp.

SPCOMPRESS_NODE_CASE

#define SPCOMPRESS_NODE_CASE

(

NumResults

)

Value:

  case NumResults:                                                             \
    return NVPTXISD::SPCOMPRESS_R##NumResults

Referenced by getSPCompressNodeOpcode().

SPDECOMPRESS_NODE_CASE

#define SPDECOMPRESS_NODE_CASE

(

NumResults

)

Value:

  case NumResults:                                                             \
    return NVPTXISD::SPDECOMPRESS_R##NumResults

Referenced by getSPDecompressNodeOpcode().

TCGEN05_LD_RED_INST

#define TCGEN05_LD_RED_INST	(		SHAPE,
			NUM,
			TYPE )

Value:

  NVPTXISD::TCGEN05_LD_RED_##SHAPE##_X##NUM##_##TYPE

Definition at line 3131 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 3128 of file NVPTXISelLowering.cpp.

Referenced by getTcgen05LdRedID().

Function Documentation

appendSplitSPOperands()

void appendSplitSPOperands ( ArrayRef< SDUse > Input, SmallVectorImpl< SDValue > & Output, SelectionDAG & DAG )

static

Definition at line 3287 of file NVPTXISelLowering.cpp.

References DL, llvm::ISD::EXTRACT_SUBVECTOR, llvm::SelectionDAG::ExtractVectorElements(), llvm::SelectionDAG::getAnyExtOrTrunc(), llvm::SelectionDAG::getBitcast(), llvm::SelectionDAG::getContext(), llvm::EVT::getIntegerVT(), llvm::SelectionDAG::getNode(), llvm::EVT::getSizeInBits(), getSPNumRegs(), getSPRegType(), llvm::SDValue::getValueType(), llvm::SelectionDAG::getVectorIdxConstant(), llvm::EVT::getVectorNumElements(), I, llvm::EVT::isVector(), llvm::SmallVectorTemplateBase< T, bool >::push_back(), and llvm::seq().

Referenced by lowerSPIntrinsic().

AreMulWideOperandsDemotable()

bool AreMulWideOperandsDemotable ( SDValue LHS, SDValue RHS, unsigned OptSize, const SelectionDAG & DAG, bool & IsSigned )

static

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

The signedness of the operands is placed in IsSigned.

Definition at line 6677 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::computeKnownBits(), llvm::SelectionDAG::ComputeMaxSignificantBits(), llvm::KnownBits::countMaxActiveBits(), llvm::dyn_cast(), IsExtendedFromLte(), llvm::APInt::isIntN(), llvm::APInt::isSignedIntN(), LHS, and RHS.

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 1839 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, assert(), DL, E(), llvm::SmallVectorTemplateCommon< T, typename >::empty(), llvm::SelectionDAG::getNode(), I, 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 436 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 8497 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::getConstant().

Referenced by simplifyDemandedBitsForPRMT().

combineADDRSPACECAST()

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

static

Definition at line 7275 of file NVPTXISelLowering.cpp.

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

combineFAddWithNeg()

SDValue combineFAddWithNeg ( SDNode * N, SelectionDAG & DAG, Intrinsic::ID AddIntrinsicID, APFloat::roundingMode RoundingMode )

static

Definition at line 7457 of file NVPTXISelLowering.cpp.

References llvm::ISD::FNEG, getFAddWithNegOpcode(), llvm::SelectionDAG::getNode(), llvm::SDValue::getOpcode(), llvm::SDValue::getOperand(), N, and Opc.

Referenced by combineIntrinsicWOChain().

combineIntrinsicWOChain()

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

static

Definition at line 7523 of file NVPTXISelLowering.cpp.

References combineFAddWithNeg(), llvm::TargetLowering::DAGCombinerInfo::DAG, diagnoseUnsupportedFPArith(), llvm::ISD::FADD, llvm::ISD::FMUL, isSupportedFPArith(), lowerSPIntrinsic(), and N.

combineLOAD()

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

static

Definition at line 6337 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 6817 of file NVPTXISelLowering.cpp.

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

Referenced by PerformMULCombineWithOperands().

combineMulSelectConstOne()

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

static

Definition at line 6828 of file NVPTXISelLowering.cpp.

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

Referenced by PerformMULCombineWithOperands().

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 6242 of file NVPTXISelLowering.cpp.

References 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, Operands, 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 7346 of file NVPTXISelLowering.cpp.

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

combineProxyReg()

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

static

Definition at line 7559 of file NVPTXISelLowering.cpp.

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

combineShiftOfLogicOp()

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

static

Commute SHL with a bitwise logic operation when doing so exposes a common shifted operand.

For example:

Before: N = shl (zext (LogicOp X, C)), ShiftAmount OtherShift = shl (zext (OtherLogicOp X, OtherC)), ShiftAmount

After: ShiftedX = shl (zext X), ShiftAmount N = LogicOp ShiftedX, ShiftedC OtherShift = OtherLogicOp ShiftedX, ShiftedOtherC

ShiftedC = (zext C) << ShiftAmount and ShiftedOtherC = (zext OtherC) << ShiftAmount are folded constants. This replaces two variable shifts with the single shared ShiftedX. Requiring another matching shift avoids disrupting isolated address calculations where a shift may be folded into the addressing mode.

Definition at line 6917 of file NVPTXISelLowering.cpp.

References C(), llvm::TargetLowering::DAGCombinerInfo::CombineTo(), llvm::TargetLowering::DAGCombinerInfo::DAG, DL, llvm::SmallVectorTemplateCommon< T, typename >::empty(), llvm::SDNode::getFlags(), llvm::SelectionDAG::getNode(), llvm::isConstOrConstSplat(), llvm::ISD::isExtOpcode(), llvm::PatternMatch::m_BitwiseLogic(), llvm::MIPatternMatch::m_OneUse(), llvm::PatternMatch::m_Value(), N, llvm::SmallVectorTemplateBase< T, bool >::push_back(), llvm::ISD::SHL, llvm::SDNode::users(), and X.

Referenced by PerformSHLCombine().

combineSTORE()

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

static

Definition at line 6322 of file NVPTXISelLowering.cpp.

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

combineSZExtToMulWide()

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

static

Definition at line 6602 of file NVPTXISelLowering.cpp.

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

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 6128 of file NVPTXISelLowering.cpp.

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

Referenced by combineLOAD().

computeKnownBitsForLoadV()

void computeKnownBitsForLoadV ( const SDValue Op, KnownBits & Known )

static

Definition at line 8422 of file NVPTXISelLowering.cpp.

References assert(), llvm::cast(), llvm::getFromTypeWidthForLoad(), llvm::Known, and llvm::ISD::SEXTLOAD.

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

computeKnownBitsForPRMT()

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

static

Definition at line 8391 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::Known, llvm::KnownBits::makeConstant(), Mode, and llvm::seq().

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

computePRMT()

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

static

Definition at line 7327 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 312 of file NVPTXISelLowering.cpp.

References assert(), llvm::ComputeValueVTs(), DL, llvm::TargetLoweringBase::getNumRegistersForCallingConv(), llvm::TargetLoweringBase::getRegisterTypeForCallingConv(), llvm::MVT::getStoreSize(), I, Off, llvm::SmallVectorTemplateBase< T, bool >::push_back(), llvm::seq(), uint64_t, 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 3879 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 1279 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().

diagnoseUnsupportedFPArith()

SDValue diagnoseUnsupportedFPArith ( SDNode * N, SelectionDAG & DAG, Intrinsic::ID IID, APFloat::roundingMode RoundingMode )

static

Definition at line 7510 of file NVPTXISelLowering.cpp.

References llvm::LLVMContext::diagnose(), llvm::Intrinsic::getBaseName(), llvm::SelectionDAG::getContext(), getDebugLoc(), llvm::EVT::getEVTString(), llvm::MachineFunction::getFunction(), llvm::SelectionDAG::getMachineFunction(), llvm::SelectionDAG::getPOISON(), llvm::nvvm::GetRoundingModeName(), and N.

Referenced by combineIntrinsicWOChain().

expandFSH64()

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

static

Definition at line 3486 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, 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 378 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::seq(), T, and llvm::Values.

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

getExtOpcode()

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

static

Definition at line 1271 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 361 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().

getFAddWithNegOpcode()

unsigned getFAddWithNegOpcode ( EVT VT, Intrinsic::ID IID, APFloat::roundingMode RoundingMode )

static

Definition at line 7428 of file NVPTXISelLowering.cpp.

References llvm::nvvm::FPArithIsSaturating(), llvm::nvvm::FPArithShouldFTZ(), llvm::EVT::getScalarType(), llvm::EVT::getSimpleVT(), llvm::EVT::isVector(), and llvm::MVT::SimpleTy.

Referenced by combineFAddWithNeg().

getMinMax3Opcode()

unsigned getMinMax3Opcode ( unsigned MinMax2Opcode )

static

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

Definition at line 6548 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 1819 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 1828 of file NVPTXISelLowering.cpp.

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

getPRMTDemandedBits()

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

static

Definition at line 8471 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 7294 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, llvm::APInt::trunc(), and llvm::V1.

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

getScalar3OpcodeForReduction()

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

static

Get 3-input scalar reduction opcode.

Definition at line 1907 of file NVPTXISelLowering.cpp.

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

getScalarAtomicRMWExpansion()

TargetLoweringBase::AtomicExpansionKind getScalarAtomicRMWExpansion ( AtomicRMWInst::BinOp Op, Type * Ty, unsigned AddressSpace, const Function & F, bool IgnoreDenormalMode, const NVPTXSubtarget & STI )

static

Returns how NVPTX should expand a scalar atomicrmw.

Definition at line 7953 of file NVPTXISelLowering.cpp.

References llvm::AtomicRMWInst::Add, llvm::NVPTXAS::ADDRESS_SPACE_GLOBAL, llvm::NVPTXAS::ADDRESS_SPACE_SHARED, llvm::NVPTXAS::ADDRESS_SPACE_SHARED_CLUSTER, AllowFTZAtomics, llvm::AtomicRMWInst::And, assert(), llvm::BitWidth, llvm::cast(), F, llvm::AtomicRMWInst::FAdd, llvm::AtomicRMWInst::FSub, llvm::NVPTXSubtarget::hasAtomAddF64(), llvm::NVPTXSubtarget::hasAtomMinMaxAndOrXor(), llvm::APFloatBase::IEEEhalf(), llvm::APFloatBase::IEEEsingle(), llvm::AtomicRMWInst::isFPOperation(), llvm_unreachable, llvm::AtomicRMWInst::Max, llvm::AtomicRMWInst::Min, llvm::AtomicRMWInst::Or, llvm::DenormalMode::PreserveSign, llvm::AtomicRMWInst::Sub, llvm::AtomicRMWInst::UDecWrap, llvm::AtomicRMWInst::UIncWrap, llvm::AtomicRMWInst::UMax, llvm::AtomicRMWInst::UMin, UseNative, llvm::AtomicRMWInst::Xchg, and llvm::AtomicRMWInst::Xor.

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

getScalarOpcodeForReduction()

ISD::NodeType getScalarOpcodeForReduction ( unsigned ReductionOpcode )

static

Definition at line 1886 of file NVPTXISelLowering.cpp.

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

getSPCompressNodeOpcode()

unsigned getSPCompressNodeOpcode ( unsigned NumResults )

static

Definition at line 3340 of file NVPTXISelLowering.cpp.

References llvm_unreachable, and SPCOMPRESS_NODE_CASE.

Referenced by lowerSPIntrinsic().

getSPDecompressNodeOpcode()

unsigned getSPDecompressNodeOpcode ( unsigned NumResults )

static

Definition at line 3369 of file NVPTXISelLowering.cpp.

References llvm_unreachable, and SPDECOMPRESS_NODE_CASE.

Referenced by lowerSPIntrinsic().

getSPNumRegs()

unsigned getSPNumRegs ( EVT VT )

static

Definition at line 3273 of file NVPTXISelLowering.cpp.

References llvm::divideCeil(), and llvm::EVT::getSizeInBits().

Referenced by appendSplitSPOperands(), joinSPRegs(), and lowerSPIntrinsic().

getSPRegType()

EVT getSPRegType ( EVT VT )

static

Definition at line 3277 of file NVPTXISelLowering.cpp.

References assert(), llvm::EVT::getVectorElementType(), and llvm::EVT::isVector().

Referenced by appendSplitSPOperands(), and joinSPRegs().

getSymbolNode() [1/2]

SDValue getSymbolNode ( SelectionDAG & DAG, const Twine & Name, EVT T )

static

Definition at line 1298 of file NVPTXISelLowering.cpp.

References llvm::MachineFunction::getContext(), llvm::SelectionDAG::getMachineFunction(), getSymbolNode(), and T.

getSymbolNode() [2/2]

SDValue getSymbolNode ( SelectionDAG & DAG, MCSymbol * Sym, EVT T )

static

Definition at line 1294 of file NVPTXISelLowering.cpp.

References llvm::SelectionDAG::getMCSymbol(), llvm::SelectionDAG::getNode(), and T.

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

getTcgen05LdRedID()

unsigned getTcgen05LdRedID ( Intrinsic::ID IID )

static

Definition at line 3134 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 2707 of file NVPTXISelLowering.cpp.

References llvm_unreachable.

Referenced by LowerTcgen05MMADisableOutputLane().

getVectorizedVT()

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

static

Definition at line 352 of file NVPTXISelLowering.cpp.

References 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 210 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 6799 of file NVPTXISelLowering.cpp.

References llvm::dyn_cast().

Referenced by combineMulSelectConstOne(), and matchMADConstOnePattern().

isConstZero()

bool isConstZero ( const SDValue & Operand )

static

Definition at line 5990 of file NVPTXISelLowering.cpp.

References llvm::dyn_cast().

Referenced by PerformADDCombineWithOperands().

IsExtendedFromLte()

bool IsExtendedFromLte ( SDValue Op, unsigned OptSize, bool & IsSigned )

static

Check whether Op is explicitly extended from at most OptSize bits, setting IsSigned to the signedness of the extension.

Definition at line 6655 of file NVPTXISelLowering.cpp.

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

Referenced by AreMulWideOperandsDemotable(), and TryMULWIDECombine().

isNonCoalescableBuildVector()

bool isNonCoalescableBuildVector ( const SDValue & BV )

static

Check if a v2f32 BUILD_VECTOR provably packs values from non-adjacent register pairs (non-coalescable).

Definition at line 6376 of file NVPTXISelLowering.cpp.

References llvm::ISD::BUILD_VECTOR, llvm::dyn_cast(), llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SDValue::getOpcode(), llvm::SDValue::getOperand(), and llvm::SDValue::getValueType().

IsPTXVectorType()

bool IsPTXVectorType ( MVT VT )

static

Definition at line 165 of file NVPTXISelLowering.cpp.

References llvm::MVT::SimpleTy.

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

isSupportedFPArith()

bool isSupportedFPArith ( SDNode * N, const NVPTXSubtarget & STI, unsigned ISDOpcode, Intrinsic::ID IID, APFloat::roundingMode RoundingMode )

static

Definition at line 7486 of file NVPTXISelLowering.cpp.

References llvm::nvvm::FPArithIsSaturating(), llvm::nvvm::FPArithShouldFTZ(), llvm::ElementCount::getFixed(), llvm::EVT::getScalarType(), llvm::EVT::getSimpleVT(), llvm::EVT::getVectorElementCount(), llvm::NVPTXSubtarget::hasF32x2Instructions(), llvm::NVPTXSubtarget::hasNativeBF16Support(), llvm::EVT::isVector(), N, llvm::APFloatBase::rmNearestTiesToEven, and llvm::MVT::SimpleTy.

Referenced by combineIntrinsicWOChain().

joinSPRegs()

SDValue joinSPRegs ( EVT VT, ArrayRef< SDValue > Regs, const SDLoc & DL, SelectionDAG & DAG )

static

Definition at line 3321 of file NVPTXISelLowering.cpp.

References assert(), llvm::ISD::CONCAT_VECTORS, DL, llvm::ArrayRef< T >::front(), llvm::SmallVectorTemplateCommon< T, typename >::front(), llvm::SelectionDAG::getBitcast(), llvm::SelectionDAG::getBuildVector(), llvm::SelectionDAG::getNode(), llvm::EVT::getSizeInBits(), getSPNumRegs(), getSPRegType(), llvm::EVT::isVector(), llvm::SmallVectorTemplateBase< T, bool >::push_back(), Reg, llvm::ArrayRef< T >::size(), and llvm::SmallVectorTemplateCommon< T, typename >::size().

Referenced by lowerSPIntrinsic().

lowerBSWAP()

SDValue lowerBSWAP ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 2585 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 3004 of file NVPTXISelLowering.cpp.

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

Referenced by lowerIntrinsicWOChain().

lowerCTLZCTPOP()

SDValue lowerCTLZCTPOP ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 3476 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 3044 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().

lowerFPTRUNC_ROUND()

SDValue lowerFPTRUNC_ROUND ( SDValue Op, SelectionDAG & DAG, const NVPTXSubtarget & STI )

static

Definition at line 2473 of file NVPTXISelLowering.cpp.

References llvm::LLVMContext::diagnose(), DL, llvm::ISD::FP_ROUND, llvm::SelectionDAG::getContext(), getDebugLoc(), llvm::MachineFunction::getFunction(), llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMachineFunction(), llvm::SelectionDAG::getNode(), llvm::SelectionDAG::getPOISON(), llvm::NearestTiesToEven, llvm::TowardNegative, llvm::TowardPositive, and llvm::TowardZero.

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

lowerFSH()

SDValue lowerFSH ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 3530 of file NVPTXISelLowering.cpp.

References expandFSH64().

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

lowerIntrinsicVoid()

SDValue lowerIntrinsicVoid ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 2904 of file NVPTXISelLowering.cpp.

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

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

lowerIntrinsicWChain()

SDValue lowerIntrinsicWChain ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 3243 of file NVPTXISelLowering.cpp.

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

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

lowerIntrinsicWOChain()

SDValue lowerIntrinsicWOChain ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 3438 of file NVPTXISelLowering.cpp.

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

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

lowerLOADi1()

SDValue lowerLOADi1 ( LoadSDNode * LD, SelectionDAG & DAG )

static

Definition at line 4053 of file NVPTXISelLowering.cpp.

References assert(), llvm::SelectionDAG::getExtLoad(), llvm::SelectionDAG::getMergeValues(), llvm::SelectionDAG::getNode(), llvm::SDValue::getValue(), 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 4042 of file NVPTXISelLowering.cpp.

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

Referenced by combineLOAD().

lowerMSTORE()

SDValue lowerMSTORE ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 3575 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, and llvm::NVPTXISD::StoreV8.

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

lowerMUL_LOHI()

SDValue lowerMUL_LOHI ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 2459 of file NVPTXISelLowering.cpp.

References DL, llvm::SelectionDAG::getContext(), llvm::SelectionDAG::getMergeValues(), llvm::SelectionDAG::getNode(), llvm::Hi, llvm::Lo, Opc, llvm::ISD::SMUL_LOHI, llvm::SelectionDAG::SplitScalar(), and llvm::EVT::widenIntegerElementType().

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

lowerPrmtIntrinsic()

SDValue lowerPrmtIntrinsic ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 3099 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 3535 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 3541 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().

lowerSPIntrinsic()

SDValue lowerSPIntrinsic ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 3390 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::SmallVectorImpl< T >::append(), appendSplitSPOperands(), assert(), DL, llvm::SelectionDAG::getMergeValues(), llvm::SelectionDAG::getNode(), llvm::EVT::getScalarSizeInBits(), getSPCompressNodeOpcode(), getSPDecompressNodeOpcode(), getSPNumRegs(), llvm::SelectionDAG::getTargetConstant(), llvm::SDValue::getValue(), llvm::EVT::getVectorNumElements(), joinSPRegs(), N, llvm::SmallVectorTemplateBase< T, bool >::push_back(), Reg, llvm::seq(), and llvm::SmallVectorTemplateCommon< T, typename >::size().

Referenced by combineIntrinsicWOChain(), and lowerIntrinsicWOChain().

lowerStAsyncRelease()

SDValue lowerStAsyncRelease ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 2655 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::LLVMContext::diagnose(), DL, llvm::SelectionDAG::getContext(), llvm::MachineFunction::getFunction(), llvm::SelectionDAG::getMachineFunction(), llvm::Intrinsic::getName(), llvm::SelectionDAG::getNode(), llvm::SelectionDAG::getTargetConstant(), llvm_unreachable, N, and llvm::ISD::ZERO_EXTEND.

Referenced by lowerIntrinsicVoid().

lowerStAsyncWithMbarrier()

SDValue lowerStAsyncWithMbarrier ( SDValue Op, SelectionDAG & DAG )

static

Definition at line 2623 of file NVPTXISelLowering.cpp.

References AbstractManglingParser< Derived, Alloc >::Ops, llvm::ISD::BITCAST, llvm::LLVMContext::diagnose(), DL, llvm::ISD::EXTRACT_VECTOR_ELT, llvm::SelectionDAG::getContext(), llvm::MachineFunction::getFunction(), llvm::SelectionDAG::getIntPtrConstant(), llvm::SelectionDAG::getMachineFunction(), llvm::Intrinsic::getName(), llvm::SelectionDAG::getNode(), and N.

Referenced by lowerIntrinsicVoid().

lowerSTOREVector()

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

static

Definition at line 4124 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, 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 2810 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 3199 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 2779 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, bool hasOffset = false )

static

Definition at line 2554 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 2876 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 2890 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 2534 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 6804 of file NVPTXISelLowering.cpp.

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

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 6352 of file NVPTXISelLowering.cpp.

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

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 6000 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, and llvm::ISD::SELECT.

PerformBUILD_VECTORCombine()

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

static

Definition at line 7214 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, llvm::ISD::SRL, llvm::ISD::TRUNCATE, and uint64_t.

PerformEXTRACTCombine()

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

static

Definition at line 7069 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::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, llvm::peekThroughFreeze(), llvm::ISD::SRA, llvm::ISD::TRUNCATE, uint64_t, and llvm::Vector.

PerformFMinMaxCombine()

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

static

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

Also covers other llvm min/max intrinsics.

Definition at line 6567 of file NVPTXISelLowering.cpp.

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

PerformMULCombine()

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

static

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

Definition at line 6886 of file NVPTXISelLowering.cpp.

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

PerformMULCombineWithOperands()

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

static

Definition at line 6858 of file NVPTXISelLowering.cpp.

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

Referenced by PerformMULCombine().

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 code like s >= 32 ? 0 : x >> s. In LLVM, over-shifting a value results in poison, but PTX shr/shl instructions clamp the shift amount to BitWidth, making the guard redundant.

Note: We only handle SRL and SHL, not SRA, because arithmetic right shifts can 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 7133 of file NVPTXISelLowering.cpp.

References llvm::BitWidth, llvm::SelectionDAG::computeKnownBits(), llvm::KnownBits::countMaxActiveBits(), llvm::TargetLowering::DAGCombinerInfo::DAG, llvm::SelectionDAG::getNode(), llvm::SDValue::getOpcode(), llvm::SDValue::getOperand(), llvm::SDValue::getValueSizeInBits(), 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, llvm::ISD::SETUGT, llvm::ISD::SETULT, and llvm::ISD::SRL.

PerformSETCCCombine()

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

static

Definition at line 7042 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, 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 7025 of file NVPTXISelLowering.cpp.

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

PerformVSELECTCombine()

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

static

Definition at line 7181 of file NVPTXISelLowering.cpp.

References llvm::ISD::BUILD_VECTOR, 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, and llvm::ISD::SELECT.

PromoteBinOpToF32()

SDValue PromoteBinOpToF32 ( SDNode * N, SelectionDAG & DAG )

static

Definition at line 2375 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 402 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 )

static

Definition at line 1250 of file NVPTXISelLowering.cpp.

References llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, llvm::ISD::ADDRSPACECAST, llvm::cast(), llvm::dyn_cast(), llvm::MachineFunction::getFrameInfo(), llvm::SelectionDAG::getMachineFunction(), llvm::MachineFrameInfo::getObjectAllocation(), llvm::SDNode::getOpcode(), and llvm::SDNode::getOperand().

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

replaceAtomicSwap128()

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

static

Definition at line 7881 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 7633 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 7841 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 7655 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 3932 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 4035 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 7865 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 2853 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().

simplifyDemandedBitsForPRMT()

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

static

Definition at line 8505 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, and llvm::TargetLowering::SimplifyMultipleUseDemandedBits().

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

sinkProxyReg()

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

static

Definition at line 7372 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, 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 6716 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::SelectionDAG::getNodeIfExists(), llvm::EVT::getSizeInBits(), llvm::Hi, llvm::isa(), IsExtendedFromLte(), LHS, llvm::ISD::MUL, llvm::ISD::MULHS, llvm::ISD::MULHU, N, Opc, RHS, llvm::APInt::sge(), llvm::ISD::SHL, llvm::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 490 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

AllowFTZAtomics

cl::opt< bool > AllowFTZAtomics("nvptx-allow-ftz-atomics", cl::Hidden, cl::desc("NVPTX Specific: Lower atomicrmw fadd to atom.add even when its " "FTZ behavior does not match the function's denormal mode."), cl::init(true)) ( "nvptx-allow-ftz-atomics" , cl::Hidden , cl::desc("NVPTX Specific: Lower atomicrmw fadd to atom.add even when its " "FTZ behavior does not match the function's denormal mode.") , cl::init(true)  )

static

Referenced by getScalarAtomicRMWExpansion().

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().

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().