MLIR: lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp File Reference

#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"

Go to the source code of this file.

Macros
#define	REDUX_F32_ID_IMPL(op, abs, hasNaN)

#define	GET_REDUX_F32_ID(op, hasAbs, hasNaN) hasAbs ? REDUX_F32_ID_IMPL(op, _abs, hasNaN) : REDUX_F32_ID_IMPL(op, , hasNaN)

#define	TCGEN05LD(SHAPE, NUM) llvm::Intrinsic::nvvm_tcgen05_ld_##SHAPE##_##NUM

#define	TCGEN05ST(SHAPE, NUM) llvm::Intrinsic::nvvm_tcgen05_st_##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)
	Return the intrinsic ID associated with ldmatrix for the given paramters. More...

static llvm::Intrinsic::ID	getStBulkIntrinsicId (LLVM::LLVMPointerType addrType)
	Return the intrinsic ID associated with st.bulk for the given address type. More...

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)

Macro Definition Documentation

◆ GET_REDUX_F32_ID

#define GET_REDUX_F32_ID	(		op,
			hasAbs,
			hasNaN
	)		hasAbs ? REDUX_F32_ID_IMPL(op, _abs, hasNaN) : REDUX_F32_ID_IMPL(op, , hasNaN)

Definition at line 33 of file NVVMToLLVMIRTranslation.cpp.

◆ 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

mlir::presburger::abs

Fraction abs(const Fraction &f)

Definition: Fraction.h:107

Definition at line 29 of file NVVMToLLVMIRTranslation.cpp.

◆ TCGEN05LD

#define TCGEN05LD	(		SHAPE,
			NUM
	)		llvm::Intrinsic::nvvm_tcgen05_ld_##SHAPE##_##NUM

Definition at line 211 of file NVVMToLLVMIRTranslation.cpp.

◆ TCGEN05ST

#define TCGEN05ST	(		SHAPE,
			NUM
	)		llvm::Intrinsic::nvvm_tcgen05_st_##SHAPE##_##NUM

Definition at line 264 of file NVVMToLLVMIRTranslation.cpp.

Function Documentation

◆ getLdMatrixIntrinsicId()

static llvm::Intrinsic::ID getLdMatrixIntrinsicId	(	NVVM::MMALayout	layout,
		int32_t	num
	)

static

Return the intrinsic ID associated with ldmatrix for the given paramters.

Definition at line 139 of file NVVMToLLVMIRTranslation.cpp.

◆ getMatchSyncIntrinsicId()

static llvm::Intrinsic::ID getMatchSyncIntrinsicId	(	Type	valType,
		NVVM::MatchSyncKind	kind
	)

static

Definition at line 108 of file NVVMToLLVMIRTranslation.cpp.

References mlir::Type::isInteger(), and kind.

◆ getReduxIntrinsicId()

static llvm::Intrinsic::ID getReduxIntrinsicId	(	llvm::Type *	resultType,
		NVVM::ReduxKind	kind,
		bool	hasAbs,
		bool	hasNaN
	)

static

Definition at line 36 of file NVVMToLLVMIRTranslation.cpp.

References GET_REDUX_F32_ID, kind, max(), and min().

◆ getShflIntrinsicId()

static llvm::Intrinsic::ID getShflIntrinsicId	(	llvm::Type *	resultType,
		NVVM::ShflKind	kind,
		bool	withPredicate
	)

static

Definition at line 67 of file NVVMToLLVMIRTranslation.cpp.

References kind.

◆ getStBulkIntrinsicId()

static llvm::Intrinsic::ID getStBulkIntrinsicId ( LLVM::LLVMPointerType addrType )

static

Return the intrinsic ID associated with st.bulk for the given address type.

Definition at line 169 of file NVVMToLLVMIRTranslation.cpp.

References isSharedMemory(), and kSharedMemorySpace.

◆ getTcgen05LdIntrinsicID()

static llvm::Intrinsic::ID getTcgen05LdIntrinsicID	(	mlir::NVVM::Tcgen05LdStShape	shape,
		uint32_t	num
	)

static

Definition at line 214 of file NVVMToLLVMIRTranslation.cpp.

References TCGEN05LD.

◆ getTcgen05StIntrinsicID()

static llvm::Intrinsic::ID getTcgen05StIntrinsicID	(	mlir::NVVM::Tcgen05LdStShape	shape,
		uint32_t	num
	)

static

Definition at line 267 of file NVVMToLLVMIRTranslation.cpp.

References TCGEN05ST.

◆ getUnidirectionalFenceProxyID()

static unsigned getUnidirectionalFenceProxyID	(	NVVM::ProxyKind	fromProxy,
		NVVM::ProxyKind	toProxy,
		NVVM::MemScopeKind	scope,
		bool	isRelease
	)

static

Definition at line 176 of file NVVMToLLVMIRTranslation.cpp.

◆ getVoteSyncIntrinsicId()

static llvm::Intrinsic::ID getVoteSyncIntrinsicId ( NVVM::VoteSyncKind kind )

static

Definition at line 124 of file NVVMToLLVMIRTranslation.cpp.

References kind.