NVVMToLLVMIRTranslation.cpp File Reference

#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Dialect/Utils/StaticValueUtils.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	getLdMatrixIntrinsicId (NVVM::MMALayout layout, int32_t num)
	Return the intrinsic ID associated with ldmatrix for the given paramters. 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 34 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 30 of file NVVMToLLVMIRTranslation.cpp.

TCGEN05LD

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

Definition at line 173 of file NVVMToLLVMIRTranslation.cpp.

TCGEN05ST

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

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

getReduxIntrinsicId()

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

static

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

References kind.

getTcgen05LdIntrinsicID()

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

static

Definition at line 176 of file NVVMToLLVMIRTranslation.cpp.

References TCGEN05LD.

getTcgen05StIntrinsicID()

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

static

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