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 {   kGlobalMemorySpace = 1 , kSharedMemorySpace = 3 , kConstantMemorySpace = 4 , kTensorMemorySpace = 6 ,   kSharedClusterMemorySpace = 7 }
	NVVM memory space identifiers. More...

Functions
void	registerConvertGpuToNVVMInterface (DialectRegistry &registry)
	Registers the ConvertToLLVMAttrInterface interface on the NVVM::NVVMTargetAttr attribute. 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 56 of file NVVMDialect.h.

Enumeration Type Documentation

NVVMMemorySpace

enum mlir::NVVM::NVVMMemorySpace

NVVM memory space identifiers.

Enumerator
kGlobalMemorySpace	Global memory space identifier.
kSharedMemorySpace	Shared memory space identifier.
kConstantMemorySpace	Constant 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	Read register with '+' modifier.
ReadWrite	Read register with '=' modifier. Note that, this is not natively supported by LLVM, but it is possible to set read and write for the same operand.

Definition at line 26 of file BasicPtxBuilderInterface.h.

Function Documentation

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 87 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().

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 682 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().

Referenced by mlir::registerAllDialects().

registerNVVMTargetInterfaceExternalModels() [1/2]

void mlir::NVVM::registerNVVMTargetInterfaceExternalModels

(

DialectRegistry &

registry

)

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

Definition at line 72 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 79 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().