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

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...
bool	operator== (unsigned as, NVVMMemorySpace memSpace)
	Utility functions to compare NVVMMemorySpace with unsigned values. More...
bool	operator== (NVVMMemorySpace memSpace, unsigned as)
bool	operator!= (unsigned as, NVVMMemorySpace memSpace)
bool	operator!= (NVVMMemorySpace memSpace, unsigned as)
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 53 of file NVVMDialect.h.

Enumeration Type Documentation

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 27 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!=() [1/2]

bool mlir::NVVM::operator!= ( NVVMMemorySpace  memSpace, unsigned  as  )

inline

Definition at line 44 of file NVVMDialect.h.

operator!=() [2/2]

bool mlir::NVVM::operator!= ( unsigned  as, NVVMMemorySpace  memSpace  )

inline

Definition at line 41 of file NVVMDialect.h.

operator<<()

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

inline

Definition at line 38 of file BasicPtxBuilderInterface.h.

References Read, ReadWrite, and Write.

operator==() [1/2]

bool mlir::NVVM::operator== ( NVVMMemorySpace  memSpace, unsigned  as  )

inline

Definition at line 38 of file NVVMDialect.h.

operator==() [2/2]

bool mlir::NVVM::operator== ( unsigned  as, NVVMMemorySpace  memSpace  )

inline

Utility functions to compare NVVMMemorySpace with unsigned values.

Definition at line 35 of file NVVMDialect.h.

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 782 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 835 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 49 of file NVVMDialect.h.

Referenced by mlir::populateGpuToNVVMConversionPatterns().