mlir::NVVM Namespace Reference

Classes
class	PtxBuilder
	A class to build PTX assembly automatically. More...
class	SerializeGPUModuleBase
	Base class for all NVVM serializations from GPU modules into binary strings. 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 }
	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...
std::unique_ptr< Pass >	createOptimizeForTargetPass ()
	Creates a pass that optimizes LLVM IR for the NVVM target. 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

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.

Definition at line 34 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

createOptimizeForTargetPass()

std::unique_ptr< Pass > mlir::NVVM::createOptimizeForTargetPass

(

)

Creates a pass that optimizes LLVM IR for the NVVM target.

Definition at line 103 of file OptimizeForNVVM.cpp.

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 75 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 502 of file LowerGpuOpsToNVVMOps.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllExtensions().

registerNVVMTargetInterfaceExternalModels() [1/2]

void mlir::NVVM::registerNVVMTargetInterfaceExternalModels

(

DialectRegistry &

registry

)

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

Definition at line 60 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 67 of file Target.cpp.

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

Variable Documentation

kSharedMemoryAlignmentBit

constexpr int mlir::NVVM::kSharedMemoryAlignmentBit = 128

constexpr

Definition at line 31 of file NVVMDialect.h.

Referenced by mlir::populateGpuToNVVMConversionPatterns().