mlir::NVVM Namespace Reference

Classes
class	PtxBuilder
	A class to build PTX assembly automatically. More...
struct	NVVMCheckSMVersion
class	SerializeGPUModuleBase
	Base class for all NVVM serializations from GPU modules into binary strings. More...

Typedefs
using	IDArgPair = std::pair< llvm::Intrinsic::ID, llvm::SmallVector< llvm::Value * > >
	A pair type of LLVM's Intrinsic ID and args (which are llvm values). More...

Enumerations
enum class	PTXRegisterMod { Read = 0 , Write = 2 , ReadWrite = 1 }
	Register read/write modifier to build constraint string for PTX inline https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html#parameters. More...
enum	NVVMMemorySpace {   kGenericMemorySpace = 0 , kGlobalMemorySpace = 1 , kSharedMemorySpace = 3 , kConstantMemorySpace = 4 ,   kLocalMemorySpace = 5 , kTensorMemorySpace = 6 , kSharedClusterMemorySpace = 7 }
	NVVM memory space identifiers. More...

Functions
void	registerConvertGpuToNVVMInterface (DialectRegistry &registry)
	Registers the ConvertToLLVMAttrInterface interface on the NVVM::NVVMTargetAttr attribute. More...
llvm::raw_ostream &	operator<< (llvm::raw_ostream &os, PTXRegisterMod mod)
void	countPlaceholderNumbers (StringRef ptxCode, llvm::SmallDenseSet< unsigned > &seenRW, llvm::SmallDenseSet< unsigned > &seenW, llvm::SmallDenseSet< unsigned > &seenR, llvm::SmallVectorImpl< unsigned > &rwNums, llvm::SmallVectorImpl< unsigned > &wNums, llvm::SmallVectorImpl< unsigned > &rNums)
	Count the number of placeholder variables such as {$r}, {$w}, {$rw} in the PTX code. More...
std::pair< mlir::Type, unsigned >	inferMMAType (mlir::NVVM::MMATypes type, mlir::NVVM::MMAFrag frag, int nRow, int nCol, mlir::MLIRContext *context)
	Return the element type and number of elements associated with a wmma matrix of given chracteristics. More...
void	registerInlinerInterface (DialectRegistry &registry)
	Register the NVVMInlinerInterface implementation of DialectInlinerInterface with the NVVM dialect. More...
void	registerNVVMTargetInterfaceExternalModels (DialectRegistry &registry)
	Registers the TargetAttrInterface for the #nvvm.target attribute in the given registry. More...
void	registerNVVMTargetInterfaceExternalModels (MLIRContext &context)
	Registers the TargetAttrInterface for the #nvvm.target attribute in the registry associated with the given context. More...
StringRef	getCUDAToolkitPath ()
	Searches & returns the path CUDA toolkit path, the search order is: More...

Variables
constexpr int	kSharedMemoryAlignmentBit = 128

Typedef Documentation

IDArgPair

using mlir::NVVM::IDArgPair = typedef std::pair<llvm::Intrinsic::ID, llvm::SmallVector<llvm::Value *> >

A pair type of LLVM's Intrinsic ID and args (which are llvm values).

This type is returned by the getIntrinsicIDAndArgs() methods.

Definition at line 60 of file NVVMDialect.h.

Enumeration Type Documentation

NVVMMemorySpace

enum mlir::NVVM::NVVMMemorySpace

NVVM memory space identifiers.

Enumerator
kGenericMemorySpace	Generic memory space identifier.
kGlobalMemorySpace	Global memory space identifier.
kSharedMemorySpace	Shared memory space identifier.
kConstantMemorySpace	Constant memory space identifier.
kLocalMemorySpace	Local memory space identifier.
kTensorMemorySpace	Tensor memory space identifier. Tensor memory is available only in arch-accelerated variants from sm100 onwards.
kSharedClusterMemorySpace	Distributed shared memory space identifier. Distributed shared memory is available only in sm90+.

Definition at line 38 of file NVVMDialect.h.

PTXRegisterMod

enum mlir::NVVM::PTXRegisterMod

strong

Register read/write modifier to build constraint string for PTX inline https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html#parameters.

Enumerator
Read	Read register with no modifier.
Write	Write register with '=' modifier.
ReadWrite	ReadWrite register with '+' modifier. Note that, this is not natively supported by LLVM, the Interface does mapping

Definition at line 26 of file BasicPtxBuilderInterface.h.

Function Documentation

countPlaceholderNumbers()

void mlir::NVVM::countPlaceholderNumbers	(	StringRef	ptxCode,
		llvm::SmallDenseSet< unsigned > &	seenRW,
		llvm::SmallDenseSet< unsigned > &	seenW,
		llvm::SmallDenseSet< unsigned > &	seenR,
		llvm::SmallVectorImpl< unsigned > &	rwNums,
		llvm::SmallVectorImpl< unsigned > &	wNums,
		llvm::SmallVectorImpl< unsigned > &	rNums
	)

Count the number of placeholder variables such as {$r}, {$w}, {$rw} in the PTX code.

Referenced by rewriteAsmPlaceholders().

getCUDAToolkitPath()

StringRef mlir::NVVM::getCUDAToolkitPath

(

)

Searches & returns the path CUDA toolkit path, the search order is:

1. The CUDA_ROOT environment variable.
2. The CUDA_HOME environment variable.
3. The CUDA_PATH environment variable.
4. The CUDA toolkit path detected by CMake.
5. Returns an empty string.

Definition at line 88 of file Target.cpp.

References __DEFAULT_CUDATOOLKIT_PATH__.

Referenced by mlir::NVVM::SerializeGPUModuleBase::SerializeGPUModuleBase().

inferMMAType()

std::pair<mlir::Type, unsigned> mlir::NVVM::inferMMAType	(	mlir::NVVM::MMATypes	type,
		mlir::NVVM::MMAFrag	frag,
		int	nRow,
		int	nCol,
		mlir::MLIRContext *	context
	)

Return the element type and number of elements associated with a wmma matrix of given chracteristics.

This matches the logic in IntrinsicsNVVM.td WMMA_REGS structure.

Referenced by mlir::convertMMAToLLVMType(), and inferMMATypeFromMNK().

operator<<()

llvm::raw_ostream& mlir::NVVM::operator<< ( llvm::raw_ostream &  os, PTXRegisterMod  mod  )

inline

Definition at line 37 of file BasicPtxBuilderInterface.h.

References Read, ReadWrite, and Write.

registerConvertGpuToNVVMInterface()

void mlir::NVVM::registerConvertGpuToNVVMInterface

(

DialectRegistry &

registry

)

Registers the ConvertToLLVMAttrInterface interface on the NVVM::NVVMTargetAttr attribute.

This interface populates the conversion target, LLVM type converter, and pattern set for converting GPU operations to NVVM.

Definition at line 686 of file LowerGpuOpsToNVVMOps.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllExtensions().

registerInlinerInterface()

void mlir::NVVM::registerInlinerInterface

(

DialectRegistry &

registry

)

Register the NVVMInlinerInterface implementation of DialectInlinerInterface with the NVVM dialect.

Definition at line 837 of file InlinerInterfaceImpl.cpp.

References mlir::DialectRegistry::addExtension().

registerNVVMTargetInterfaceExternalModels() [1/2]

void mlir::NVVM::registerNVVMTargetInterfaceExternalModels

(

DialectRegistry &

registry

)

Registers the TargetAttrInterface for the #nvvm.target attribute in the given registry.

Definition at line 73 of file Target.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllDialects(), and registerNVVMTargetInterfaceExternalModels().

registerNVVMTargetInterfaceExternalModels() [2/2]

void mlir::NVVM::registerNVVMTargetInterfaceExternalModels

(

MLIRContext &

context

)

Registers the TargetAttrInterface for the #nvvm.target attribute in the registry associated with the given context.

Definition at line 80 of file Target.cpp.

References mlir::MLIRContext::appendDialectRegistry(), and registerNVVMTargetInterfaceExternalModels().

Variable Documentation

kSharedMemoryAlignmentBit

constexpr int mlir::NVVM::kSharedMemoryAlignmentBit = 128

constexpr

Definition at line 35 of file NVVMDialect.h.

Referenced by mlir::populateGpuToNVVMConversionPatterns().