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, 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. 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

Definition at line 29 of file NVVMToLLVMIRTranslation.cpp.

TCGEN05LD

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

Definition at line 297 of file NVVMToLLVMIRTranslation.cpp.

TCGEN05ST

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

Definition at line 350 of file NVVMToLLVMIRTranslation.cpp.

Function Documentation

getLdMatrixIntrinsicId()

static llvm::Intrinsic::ID getLdMatrixIntrinsicId ( NVVM::MMALayout  layout, int32_t  num, NVVM::LdStMatrixShapeAttr  shape, NVVM::LdStMatrixEltType  eltType  )

static

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 255 of file NVVMToLLVMIRTranslation.cpp.

References isSharedMemory().

getStMatrixIntrinsicId()

static llvm::Intrinsic::ID getStMatrixIntrinsicId ( NVVM::MMALayout  layout, int32_t  num, NVVM::LdStMatrixShapeAttr  shape, NVVM::LdStMatrixEltType  eltType  )

static

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

Definition at line 219 of file NVVMToLLVMIRTranslation.cpp.

getTcgen05LdIntrinsicID()

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

static

Definition at line 300 of file NVVMToLLVMIRTranslation.cpp.

References TCGEN05LD.

getTcgen05StIntrinsicID()

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

static

Definition at line 353 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 262 of file NVVMToLLVMIRTranslation.cpp.

getVoteSyncIntrinsicId()

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

static

Definition at line 124 of file NVVMToLLVMIRTranslation.cpp.

References kind.