mlir::gpu Namespace Reference

Namespaces
	amd

Classes
class	AsyncTokenType
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	SerializeToBlobPass
	Base pass class to serialize kernel functions through LLVM into user-specified IR and add the resulting blob as module attribute. More...

Functions
void	addAsyncDependency (Operation *op, Value token)
StringRef	getMappingAttrName ()
	Name of the mapping attribute produced by loop mappers. More...
Processor	getProcessor (ParallelLoopDimMapping attr)
	Get the value of the processor in the ParallelLoopDimMapping attribute. More...
ParallelLoopDimMapping	getParallelLoopDimMappingAttr (Processor processor, AffineMap map, AffineMap bound)
	Helper function to create a ParallelDimMapperAttr. More...
LogicalResult	setMappingAttr (scf::ParallelOp ploopOp, ArrayRef< ParallelLoopDimMapping > mapping)
	Sets the mapping attribute of a scf.parallel operation. More...
std::string	getDefaultGpuBinaryAnnotation ()
	Returns the default annotation name for GPU binary blobs. More...

Function Documentation

addAsyncDependency()

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

Definition at line 371 of file GPUDialect.cpp.

References mlir::Block::addArgument(), mlir::OperationState::addAttribute(), mlir::OperationState::addOperands(), mlir::OperationState::addRegion(), mlir::Operation::getContext(), mlir::Builder::getI32VectorAttr(), mlir::Builder::getIndexType(), mlir::OpTrait::AttrSizedOperandSegments< ConcreteType >::getOperandSegmentSizeAttr(), mlir::Builder::getType(), mlir::Operation::insertOperands(), mlir::OperationState::location, mlir::Region::push_back(), mlir::Operation::setAttr(), mlir::success(), and mlir::OperationState::types.

getDefaultGpuBinaryAnnotation()

std::string mlir::gpu::getDefaultGpuBinaryAnnotation

(

)

Returns the default annotation name for GPU binary blobs.

Definition at line 30 of file SerializeToBlob.cpp.

Referenced by mlir::populateGpuRewritePatterns().

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(), processParallelLoop(), and setMappingAttr().

getParallelLoopDimMappingAttr()

ParallelLoopDimMapping mlir::gpu::getParallelLoopDimMappingAttr	(	Processor	processor,
		AffineMap	map,
		AffineMap	bound
	)

Helper function to create a ParallelDimMapperAttr.

TODO: Replace its uses with an auto-gened method.

Definition at line 33 of file ParallelLoopMapper.cpp.

References mlir::Builder::getContext(), and mlir::AffineMap::getContext().

Referenced by getProcessor(), and mapParallelOp().

getProcessor()

Processor mlir::gpu::getProcessor ( ParallelLoopDimMapping  attr )

inline

Get the value of the processor in the ParallelLoopDimMapping attribute.

Definition at line 45 of file ParallelLoopMapper.h.

References getParallelLoopDimMappingAttr(), mlir::greedilyMapParallelSCFToGPU(), and setMappingAttr().

Referenced by processParallelLoop(), and setMappingAttr().

setMappingAttr()

LogicalResult mlir::gpu::setMappingAttr	(	scf::ParallelOp	ploopOp,
		ArrayRef< ParallelLoopDimMapping >	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.

Definition at line 43 of file ParallelLoopMapper.cpp.

References getMappingAttrName(), getProcessor(), and mlir::success().

Referenced by getProcessor(), and mapParallelOp().