mlir::gpu Namespace Reference

Namespaces
	amd

Classes
struct	KernelDim3
	Utility class for the GPU dialect to represent triples of Values accessible through .x, .y, and .z similarly to CUDA notation. More...
class	AsyncTokenType
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	SparseHandleType
class	SerializeToBlobPass
	Base pass class to serialize kernel functions through LLVM into user-specified IR and add the resulting blob as module attribute. More...

Typedefs
using	SparseEnvHandleType = SparseHandleType< SparseHandleKind::Env >
using	SparseDnVecHandleType = SparseHandleType< SparseHandleKind::DnVec >
using	SparseDnMatHandleType = SparseHandleType< SparseHandleKind::DnMat >
using	SparseSpMatHandleType = SparseHandleType< SparseHandleKind::SpMat >

Enumerations
enum class	SparseHandleKind { Env , DnVec , DnMat , SpMat }

Functions
void	addAsyncDependency (Operation *op, Value token)
void	registerTransformDialectExtension (DialectRegistry &registry)
StringRef	getMappingAttrName ()
	Name of the mapping attribute produced by loop mappers. More...
LogicalResult	setMappingAttr (scf::ParallelOp ploopOp, ArrayRef< ParallelLoopDimMappingAttr > mapping)
	Sets the mapping attribute of a scf.parallel operation. More...
std::string	getDefaultGpuBinaryAnnotation ()
	Returns the default annotation name for GPU binary blobs. More...
static MappingLevel &	operator++ (MappingLevel &mappingLevel)
	Bounded increment on MappingLevel. More...
static Processor	getHardwareIdForMapping (MappingLevel level, int dimension)
	Computed the hardware id to use for a given mapping level. More...
static void	mapParallelOp (ParallelOp parallelOp, MappingLevel mappingLevel=MapGrid)
	Add mapping information to the given parallel loop. More...

Variables
static constexpr int	kNumHardwareIds = 3

Typedef Documentation

SparseDnMatHandleType

using mlir::gpu::SparseDnMatHandleType = typedef SparseHandleType<SparseHandleKind::DnMat>

Definition at line 181 of file GPUDialect.h.

SparseDnVecHandleType

using mlir::gpu::SparseDnVecHandleType = typedef SparseHandleType<SparseHandleKind::DnVec>

Definition at line 180 of file GPUDialect.h.

SparseEnvHandleType

using mlir::gpu::SparseEnvHandleType = typedef SparseHandleType<SparseHandleKind::Env>

Definition at line 179 of file GPUDialect.h.

SparseSpMatHandleType

using mlir::gpu::SparseSpMatHandleType = typedef SparseHandleType<SparseHandleKind::SpMat>

Definition at line 182 of file GPUDialect.h.

Enumeration Type Documentation

SparseHandleKind

enum mlir::gpu::SparseHandleKind

strong

Enumerator
Env
DnVec
DnMat
SpMat

Definition at line 168 of file GPUDialect.h.

Function Documentation

addAsyncDependency()

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

Definition at line 538 of file GPUDialect.cpp.

getDefaultGpuBinaryAnnotation()

std::string mlir::gpu::getDefaultGpuBinaryAnnotation

(

)

Returns the default annotation name for GPU binary blobs.

Definition at line 33 of file SerializeToBlob.cpp.

getHardwareIdForMapping()

static 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. TODO: Make this use x for the inner-most loop that is distributed to map to x, the next innermost to y and the next innermost to z.

Definition at line 73 of file ParallelLoopMapper.cpp.

References kNumHardwareIds.

Referenced by mapParallelOp().

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

mapParallelOp()

static void mlir::gpu::mapParallelOp ( ParallelOp  parallelOp, MappingLevel  mappingLevel = MapGrid  )

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 109 of file ParallelLoopMapper.cpp.

References mlir::Builder::getAttr(), mlir::Builder::getDimIdentityMap(), getHardwareIdForMapping(), getMappingAttrName(), and setMappingAttr().

operator++()

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

static

Bounded increment on MappingLevel.

Increments to the next level unless Sequential was already reached.

Definition at line 60 of file ParallelLoopMapper.cpp.

registerTransformDialectExtension()

void mlir::gpu::registerTransformDialectExtension

(

DialectRegistry &

registry

)

Definition at line 920 of file GPUTransformOps.cpp.

References mlir::DialectRegistry::addExtensions().

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

Variable Documentation

kNumHardwareIds

constexpr int mlir::gpu::kNumHardwareIds = 3

staticconstexpr

Definition at line 56 of file ParallelLoopMapper.cpp.

Referenced by getHardwareIdForMapping().