MLIR  15.0.0git
Namespaces | Classes | Functions
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 
)

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