#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/IR/Operation.h"
#include "mlir/Target/LLVMIR/ModuleTranslation.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/ADT/iterator_range.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/Support/FormatVariadic.h"
#include "mlir/Dialect/LLVMIR/NVVMConversions.inc"

Macros
#define	REDUX_F32_ID_IMPL(op, abs, hasNaN)
#define	GET_REDUX_F32_ID(op, hasAbs, hasNaN)
#define	TCGEN05LD(SHAPE, NUM)
#define	TCGEN05ST(SHAPE, NUM)

Functions
static llvm::Intrinsic::ID	getReduxIntrinsicId (llvm::Type *resultType, NVVM::ReduxKind kind, bool hasAbs, bool hasNaN)
static llvm::Intrinsic::ID	getShflIntrinsicId (llvm::Type *resultType, NVVM::ShflKind kind, bool withPredicate)
static llvm::Intrinsic::ID	getMatchSyncIntrinsicId (Type valType, NVVM::MatchSyncKind kind)
static llvm::Intrinsic::ID	getVoteSyncIntrinsicId (NVVM::VoteSyncKind kind)
static llvm::Intrinsic::ID	getLdMatrixIntrinsicId (NVVM::MMALayout layout, int32_t num, NVVM::LdStMatrixShapeAttr shape, NVVM::LdStMatrixEltType eltType)
static llvm::Intrinsic::ID	getStMatrixIntrinsicId (NVVM::MMALayout layout, int32_t num, NVVM::LdStMatrixShapeAttr shape, NVVM::LdStMatrixEltType eltType)
	Return the intrinsic ID associated with stmatrix for the given paramters.
static llvm::Intrinsic::ID	getStBulkIntrinsicId (LLVM::LLVMPointerType addrType)
	Return the intrinsic ID associated with st.bulk for the given address type.
static unsigned	getUnidirectionalFenceProxyID (NVVM::ProxyKind fromProxy, NVVM::ProxyKind toProxy, NVVM::MemScopeKind scope, bool isRelease)
static llvm::Intrinsic::ID	getTcgen05LdIntrinsicID (mlir::NVVM::Tcgen05LdStShape shape, uint32_t num)
static llvm::Intrinsic::ID	getTcgen05StIntrinsicID (mlir::NVVM::Tcgen05LdStShape shape, uint32_t num)
llvm::CallInst *	createIntrinsicCall (llvm::IRBuilderBase &builder, llvm::Intrinsic::ID intrinsic, ArrayRef< llvm::Value * > args={}, ArrayRef< llvm::Type * > tys={})
	Creates a call to an LLVM IR intrinsic function with the given arguments.

Macro Definition Documentation

◆ GET_REDUX_F32_ID

#define GET_REDUX_F32_ID	(	op,
		hasAbs,
		hasNaN )

Value:

hasAbs ? REDUX_F32_ID_IMPL(op, _abs, hasNaN) : REDUX_F32_ID_IMPL(op, , hasNaN)

REDUX_F32_ID_IMPL

#define REDUX_F32_ID_IMPL(op, abs, hasNaN)

Definition NVVMToLLVMIRTranslation.cpp:29

Definition at line 33 of file NVVMToLLVMIRTranslation.cpp.

Referenced by getReduxIntrinsicId().

◆ REDUX_F32_ID_IMPL

#define REDUX_F32_ID_IMPL	(	op,
		abs,
		hasNaN )

Value:

hasNaN ? llvm::Intrinsic::nvvm_redux_sync_f##op##abs##_NaN \

: llvm::Intrinsic::nvvm_redux_sync_f##op##abs

Definition at line 29 of file NVVMToLLVMIRTranslation.cpp.

◆ TCGEN05LD

#define TCGEN05LD	(		SHAPE,
			NUM )

Value:

llvm::Intrinsic::nvvm_tcgen05_ld_##SHAPE##_##NUM

Definition at line 294 of file NVVMToLLVMIRTranslation.cpp.

Referenced by getTcgen05LdIntrinsicID().

◆ TCGEN05ST

#define TCGEN05ST	(		SHAPE,
			NUM )

Value:

llvm::Intrinsic::nvvm_tcgen05_st_##SHAPE##_##NUM

Definition at line 347 of file NVVMToLLVMIRTranslation.cpp.

Referenced by getTcgen05StIntrinsicID().

Function Documentation

◆ createIntrinsicCall()

llvm::CallInst * mlir::LLVM::detail::createIntrinsicCall	(	llvm::IRBuilderBase &	builder,
		llvm::Intrinsic::ID	intrinsic,
		ArrayRef< llvm::Value * >	args = {},
		ArrayRef< llvm::Type * >	tys = {} )