MLIR  19.0.0git
Namespaces | Classes | Enumerations | Functions | Variables
mlir::gpu Namespace Reference

Namespaces

 amd
 

Classes

class  OffloadingTranslationAttrTrait
 This class indicates that the attribute associated with this trait is a GPU offloading translation attribute. More...
 
class  TargetOptions
 This class serves as an opaque interface for passing options to the TargetAttrInterface methods. More...
 
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  SparseDnTensorHandleType
 
class  SparseSpMatHandleType
 
class  SparseSpGEMMOpHandleType
 
struct  GPUToNVVMPipelineOptions
 Options for the gpu to nvvm pipeline. 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...
 

Enumerations

enum class  SparseHandleKind { SpMat , DnTensor , SpGEMMOp }
 

Functions

void addAsyncDependency (Operation *op, Value token)
 
void buildLowerToNVVMPassPipeline (OpPassManager &pm, const GPUToNVVMPipelineOptions &options)
 Adds the GPU to NVVM pipeline to the given pass manager. More...
 
void registerGPUToNVVMPipeline ()
 Register all pipeleines for the gpu dialect. More...
 
void registerTransformDialectExtension (DialectRegistry &registry)
 
void registerBufferDeallocationOpInterfaceExternalModels (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...
 
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. More...
 
std::string getDefaultGpuBinaryAnnotation ()
 Returns the default annotation name for GPU binary blobs. More...
 
vector::CombiningKind convertReductionKind (gpu::AllReduceOperation mode)
 Returns the matching vector combining kind. More...
 
void registerOffloadingLLVMTranslationInterfaceExternalModels (mlir::DialectRegistry &registry)
 Registers the offloading LLVM translation interfaces for gpu.select_object. 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
 

Enumeration Type Documentation

◆ SparseHandleKind

Enumerator
SpMat 
DnTensor 
SpGEMMOp 

Definition at line 174 of file GPUDialect.h.

Function Documentation

◆ addAsyncDependency()

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

Definition at line 625 of file GPUDialect.cpp.

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

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

◆ getDefaultGpuBinaryAnnotation()

std::string mlir::gpu::getDefaultGpuBinaryAnnotation ( )

Returns the default annotation name for GPU binary blobs.

Definition at line 35 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 74 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 110 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 61 of file ParallelLoopMapper.cpp.

◆ registerBufferDeallocationOpInterfaceExternalModels()

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

◆ registerGPUToNVVMPipeline()

void mlir::gpu::registerGPUToNVVMPipeline ( )

Register all pipeleines for the gpu dialect.

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 57 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 950 of file GPUTransformOps.cpp.

References mlir::DialectRegistry::addExtensions().

Referenced by mlir::registerAllExtensions().

◆ 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 141 of file ModuleToBinary.cpp.

Variable Documentation

◆ kNumHardwareIds

constexpr int mlir::gpu::kNumHardwareIds = 3
staticconstexpr

Definition at line 57 of file ParallelLoopMapper.cpp.

Referenced by getHardwareIdForMapping().