Namespaces
namespace	amd
namespace	index_lowering

Classes
class	AsyncTokenType
struct	GPUToNVVMPipelineOptions
	Options for the gpu to nvvm pipeline. More...
struct	GPUToXeVMPipelineOptions
struct	KernelDim3
	Utility class for the GPU dialect to represent triples of Values accessible through .x, .y, and .z similarly to CUDA notation. More...
struct	MMAMatrixStorageType
	MMAMatrixType storage and uniquing. More...
class	MMAMatrixType
	MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply accumulate operations. More...
class	OffloadingTranslationAttrTrait
	This class indicates that the attribute associated with this trait is a GPU offloading translation attribute. More...
class	SerializedObject
	This class represents a serialized object (GPU binary) with metadata (e.g. More...
class	SparseDnTensorHandleType
class	SparseSpGEMMOpHandleType
class	SparseSpMatHandleType
class	TargetOptions
	This class serves as an opaque interface for passing options to the TargetAttrInterface methods. More...
struct	WarpDistributionPattern

Enumerations
enum class	SparseHandleKind { SpMat , DnTensor , SpGEMMOp }

Functions
void	registerConvertGpuToLLVMInterface (DialectRegistry &registry)
	Registers the ConvertToLLVMOpInterface interface on the gpu::GPUModuleOP operation.
void	addAsyncDependency (Operation *op, Value token)
std::optional< uint32_t >	getKnownDimensionSizeAround (Operation *op, DimensionKind kind, Dimension dim)
	Retrieve the constant bounds for a given dimension and dimension kind from the context surrounding op, if known, and return them.
void	registerValueBoundsOpInterfaceExternalModels (DialectRegistry &registry)
void	buildLowerToNVVMPassPipeline (OpPassManager &pm, const GPUToNVVMPipelineOptions &options)
	Adds the GPU to NVVM pipeline to the given pass manager.
void	buildLowerToXeVMPassPipeline (OpPassManager &pm, const GPUToXeVMPipelineOptions &options)
	Adds the GPU to XeVM pipeline to the given pass manager.
void	registerGPUToNVVMPipeline ()
	Register all pipelines for the gpu dialect.
void	registerGPUToXeVMPipeline ()
void	registerTransformDialectExtension (DialectRegistry &registry)
void	registerBufferDeallocationOpInterfaceExternalModels (DialectRegistry &registry)
StringRef	getMappingAttrName ()
	Name of the mapping attribute produced by loop mappers.
LogicalResult	setMappingAttr (scf::ParallelOp ploopOp, ArrayRef< ParallelLoopDimMappingAttr > mapping)
	Sets the mapping attribute of a scf.parallel operation.
LogicalResult	transformGpuModulesToBinaries (Operation *op, OffloadingLLVMTranslationAttrInterface handler=nullptr, const gpu::TargetOptions &options={})
	Searches for all GPU modules in op and transforms them into GPU binary operations.
vector::CombiningKind	convertReductionKind (gpu::AllReduceOperation mode)
	Returns the matching vector combining kind.
void	registerOffloadingLLVMTranslationInterfaceExternalModels (mlir::DialectRegistry &registry)
	Registers the offloading LLVM translation interfaces for gpu.select_object.
static MappingLevel &	operator++ (MappingLevel &mappingLevel)
	Bounded increment on MappingLevel.
static FailureOr< MappingPolicy >	getMappingPolicyFromStr (StringRef policy)
static Processor	getHardwareIdForMapping (MappingLevel level, int dimension)
	Computed the hardware id to use for a given mapping level.
static void	mapParallelOp (ParallelOp parallelOp, MappingLevel mappingLevel=MapGrid, MappingPolicy mappingPolicy=MappingPolicy::OutermostFirst)
	Add mapping information to the given parallel loop.

Variables
constexpr StringLiteral	elfSectionName = "section"
static constexpr int	kNumHardwareIds = 3

Enumeration Type Documentation

◆ SparseHandleKind

enum class mlir::gpu::SparseHandleKind

strong

Enumerator
SpMat
DnTensor
SpGEMMOp

Definition at line 176 of file GPUDialect.h.

Function Documentation

◆ addAsyncDependency()

void mlir::gpu::addAsyncDependency	(	Operation *	op,
		Value	token )

Definition at line 668 of file GPUDialect.cpp.

References mlir::Operation::getContext(), mlir::Builder::getDenseI32ArrayAttr(), mlir::OpTrait::AttrSizedOperandSegments< ConcreteType >::getOperandSegmentSizeAttr(), mlir::Operation::insertOperands(), and mlir::Operation::setAttr().

◆ buildLowerToNVVMPassPipeline()

void mlir::gpu::buildLowerToNVVMPassPipeline	(	OpPassManager &	pm,
		const GPUToNVVMPipelineOptions &	options )

Adds the GPU to NVVM pipeline to the given pass manager.

Transforms main dialects into NVVM targets. Begins with GPU code regions, then handles host code.

Definition at line 103 of file GPUToNVVMPipeline.cpp.

References options.

Referenced by registerGPUToNVVMPipeline().

◆ buildLowerToXeVMPassPipeline()

void mlir::gpu::buildLowerToXeVMPassPipeline	(	OpPassManager &	pm,
		const GPUToXeVMPipelineOptions &	options )

Adds the GPU to XeVM pipeline to the given pass manager.

Transforms main dialects into XeVM targets. Begins with GPU code regions, then handles host code.

Definition at line 143 of file GPUToXeVMPipeline.cpp.

References options.

Referenced by registerGPUToXeVMPipeline().

◆ convertReductionKind()

vector::CombiningKind mlir::gpu::convertReductionKind ( gpu::AllReduceOperation mode )

Returns the matching vector combining kind.

Definition at line 18 of file Utils.cpp.

References MAP_CASE, and MINUI.

◆ getHardwareIdForMapping()

Processor mlir::gpu::getHardwareIdForMapping	(	MappingLevel	level,
		int	dimension )

static

Computed the hardware id to use for a given mapping level.

Will assign x,y and z hardware ids for the first 3 dimensions and use sequential after.

Definition at line 88 of file ParallelLoopMapper.cpp.

References kNumHardwareIds.

Referenced by mapParallelOp().

◆ getKnownDimensionSizeAround()

std::optional< uint32_t > mlir::gpu::getKnownDimensionSizeAround	(	Operation *	op,
		DimensionKind	kind,
		Dimension	dim )

Retrieve the constant bounds for a given dimension and dimension kind from the context surrounding op, if known, and return them.

This will check the bounds on an enclosing gpu.launch, an enclosing gpu.func, and any gpu.known_*_size on other function-like operations, in that order.

Definition at line 80 of file InferIntRangeInterfaceImpls.cpp.

References getKnownLaunchAttr(), mlir::Operation::getParentOfType(), mlir::m_ConstantInt(), mlir::matchPattern(), and valueByDim().

Referenced by mlir::gpu::index_lowering::getIndexOpRange(), and getKnownOrOcklDim().

◆ getMappingAttrName()

StringRef mlir::gpu::getMappingAttrName ( )

Name of the mapping attribute produced by loop mappers.

Definition at line 31 of file ParallelLoopMapper.cpp.

Referenced by mlir::configureParallelLoopToGPULegality(), mapParallelOp(), and processParallelLoop().

◆ getMappingPolicyFromStr()

FailureOr< MappingPolicy > mlir::gpu::getMappingPolicyFromStr ( StringRef policy )

static

Definition at line 71 of file ParallelLoopMapper.cpp.

◆ mapParallelOp()

void mlir::gpu::mapParallelOp	(	ParallelOp	parallelOp,
		MappingLevel	mappingLevel = MapGrid,
		MappingPolicy	mappingPolicy = MappingPolicy::OutermostFirst )

static

Add mapping information to the given parallel loop.

Do not add mapping information if the loop already has it. Also, don't start a mapping at a nested loop.

Definition at line 126 of file ParallelLoopMapper.cpp.

References b, getHardwareIdForMapping(), getMappingAttrName(), kNumHardwareIds, mapParallelOp(), and setMappingAttr().

Referenced by mapParallelOp().

◆ operator++()

MappingLevel & mlir::gpu::operator++ ( MappingLevel & mappingLevel )

static

Bounded increment on MappingLevel.

Increments to the next level unless Sequential was already reached.

Definition at line 62 of file ParallelLoopMapper.cpp.

◆ registerBufferDeallocationOpInterfaceExternalModels()

void mlir::gpu::registerBufferDeallocationOpInterfaceExternalModels ( DialectRegistry & registry )

Definition at line 31 of file BufferDeallocationOpInterfaceImpl.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllDialects().

◆ registerConvertGpuToLLVMInterface()

void mlir::gpu::registerConvertGpuToLLVMInterface ( DialectRegistry & registry )

Registers the ConvertToLLVMOpInterface interface on the gpu::GPUModuleOP operation.

Definition at line 1868 of file GPUToLLVMConversion.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllExtensions().

◆ registerGPUToNVVMPipeline()

void mlir::gpu::registerGPUToNVVMPipeline ( )

Register all pipelines for the gpu dialect.

Definition at line 115 of file GPUToNVVMPipeline.cpp.

References buildLowerToNVVMPassPipeline().

Referenced by mlir::registerAllPasses().

◆ registerGPUToXeVMPipeline()

void mlir::gpu::registerGPUToXeVMPipeline ( )

Definition at line 155 of file GPUToXeVMPipeline.cpp.

References buildLowerToXeVMPassPipeline().

Referenced by mlir::registerAllPasses().

◆ registerOffloadingLLVMTranslationInterfaceExternalModels()

void mlir::gpu::registerOffloadingLLVMTranslationInterfaceExternalModels ( mlir::DialectRegistry & registry )

Registers the offloading LLVM translation interfaces for gpu.select_object.

Definition at line 467 of file SelectObjectAttr.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllGPUToLLVMIRTranslations(), and mlir::registerAllToLLVMIRTranslations().

◆ registerTransformDialectExtension()

void mlir::gpu::registerTransformDialectExtension ( DialectRegistry & registry )

Definition at line 978 of file GPUTransformOps.cpp.

References mlir::DialectRegistry::addExtensions().

Referenced by mlir::registerAllExtensions().

◆ registerValueBoundsOpInterfaceExternalModels()

void mlir::gpu::registerValueBoundsOpInterfaceExternalModels ( DialectRegistry & registry )

Definition at line 115 of file ValueBoundsOpInterfaceImpl.cpp.

References mlir::DialectRegistry::addExtension(), and REGISTER.

Referenced by mlir::registerAllDialects().

◆ setMappingAttr()

LogicalResult mlir::gpu::setMappingAttr	(	scf::ParallelOp	ploopOp,
		ArrayRef< ParallelLoopDimMappingAttr >	mapping )

Sets the mapping attribute of a scf.parallel operation.

Verifies that the mapping passed is valid.

the number of DimMapperAttr provided is same as the number of loops of the ploopOp.
the mapping does not map multiple loops to the same processor.

Referenced by mapParallelOp().

◆ transformGpuModulesToBinaries()

LogicalResult mlir::gpu::transformGpuModulesToBinaries	(	Operation *	op,
		OffloadingLLVMTranslationAttrInterface	handler = nullptr,
		const gpu::TargetOptions &	options = {} )

Searches for all GPU modules in op and transforms them into GPU binary operations.

The resulting gpu.binary has handler as its offloading handler attribute.

Definition at line 123 of file ModuleToBinary.cpp.

References mlir::Operation::getRegions(), and success().

Variable Documentation

◆ elfSectionName

StringLiteral mlir::gpu::elfSectionName = "section"

constexpr

Definition at line 31 of file CompilationInterfaces.h.

Referenced by llvm::embedBinaryImpl().

◆ kNumHardwareIds

int mlir::gpu::kNumHardwareIds = 3

staticconstexpr

Definition at line 58 of file ParallelLoopMapper.cpp.

Referenced by getHardwareIdForMapping(), and mapParallelOp().

Namespaces

Classes

Enumerations

Functions

Variables

Enumeration Type Documentation

◆ SparseHandleKind

Function Documentation

◆ addAsyncDependency()

◆ buildLowerToNVVMPassPipeline()

◆ buildLowerToXeVMPassPipeline()

◆ convertReductionKind()

◆ getHardwareIdForMapping()

◆ getKnownDimensionSizeAround()

◆ getMappingAttrName()

◆ getMappingPolicyFromStr()

◆ mapParallelOp()

◆ operator++()

◆ registerBufferDeallocationOpInterfaceExternalModels()

◆ registerConvertGpuToLLVMInterface()

◆ registerGPUToNVVMPipeline()

◆ registerGPUToXeVMPipeline()

◆ registerOffloadingLLVMTranslationInterfaceExternalModels()

◆ registerTransformDialectExtension()

◆ registerValueBoundsOpInterfaceExternalModels()

◆ setMappingAttr()

◆ transformGpuModulesToBinaries()

Variable Documentation

◆ elfSectionName

◆ kNumHardwareIds