MLIR  15.0.0git
Public Member Functions | Static Public Member Functions | List of all members
mlir::gpu::MMAMatrixType Class Reference

MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply accumulate operations. More...

#include "mlir/Dialect/GPU/GPUDialect.h"

+ Inheritance diagram for mlir::gpu::MMAMatrixType:
+ Collaboration diagram for mlir::gpu::MMAMatrixType:

Public Member Functions

unsigned getNumDims () const
 Get number of dims. More...
 
ArrayRef< int64_t > getShape () const
 Get shape of the matrix. More...
 
Type getElementType () const
 Get elementType of a single element. More...
 
StringRef getOperand () const
 The general form of operation this type supports is given by the equation C += A*B. More...
 

Static Public Member Functions

static MMAMatrixType get (ArrayRef< int64_t > shape, Type elementType, StringRef operand)
 Get MMAMatrixType and verify construction Invariants. More...
 
static MMAMatrixType getChecked (function_ref< InFlightDiagnostic()> emitError, ArrayRef< int64_t > shape, Type elementType, StringRef operand)
 Get MMAMatrixType at a particular location and verify construction Invariants. More...
 
static bool isValidElementType (Type elementType)
 Check if a type is valid a MMAMatrixType elementType. More...
 
static LogicalResult verify (function_ref< InFlightDiagnostic()> emitError, ArrayRef< int64_t > shape, Type elementType, StringRef operand)
 Verify that shape and elementType are actually allowed for the MMAMatrixType. More...
 
- Static Public Member Functions inherited from mlir::detail::StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits >
static TypeID getTypeID ()
 Return a unique identifier for the concrete type. More...
 
template<typename T >
static bool classof (T val)
 Provide an implementation of 'classof' that compares the type id of the provided value with that of the concrete type. More...
 
static detail::InterfaceMap getInterfaceMap ()
 Returns an interface map for the interfaces registered to this storage user. More...
 
static HasTraitFn getHasTraitFn ()
 Returns the function that returns true if the given Trait ID matches the IDs of any of the traits defined by the storage user. More...
 
template<typename... IfaceModels>
static void attachInterface (MLIRContext &context)
 Attach the given models as implementations of the corresponding interfaces for the concrete storage user class. More...
 
template<typename... Args>
static ConcreteT get (MLIRContext *ctx, Args... args)
 Get or create a new ConcreteT instance within the ctx. More...
 
template<typename... Args>
static ConcreteT getChecked (const Location &loc, Args... args)
 Get or create a new ConcreteT instance within the ctx, defined at the given, potentially unknown, location. More...
 
template<typename... Args>
static ConcreteT getChecked (function_ref< InFlightDiagnostic()> emitErrorFn, MLIRContext *ctx, Args... args)
 Get or create a new ConcreteT instance within the ctx. More...
 
static ConcreteT getFromOpaquePointer (const void *ptr)
 Get an instance of the concrete type from a void pointer. More...
 

Additional Inherited Members

- Public Types inherited from mlir::detail::StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits >
using Base = StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits... >
 Utility declarations for the concrete attribute class. More...
 
using ImplType = StorageT
 
using HasTraitFn = bool(*)(TypeID)
 
- Protected Member Functions inherited from mlir::detail::StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits >
template<typename... Args>
LogicalResult mutate (Args &&...args)
 Mutate the current storage instance. More...
 
ImplTypegetImpl () const
 Utility for easy access to the storage instance. More...
 
- Static Protected Member Functions inherited from mlir::detail::StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits >
template<typename... Args>
static LogicalResult verify (Args... args)
 Default implementation that just returns success. More...
 

Detailed Description

MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply accumulate operations.

MMAMatrices are taken as direct operands by these operations and are also produced as results. These matrices are meant to reside in the registers. A limited number of pointwise operations can be performed on these matrices, i.e., operations which operate uniformly on all the elements in the matrix and do not change the order of matrix elements. The above conditions exist because the layout of matrix elements inside the matrix is opaque i.e., the elements may be present in the matrix in any order. The general usage of this type is shown as follows:-

%0 = gpu.subgroup_mma_load_matrix arg0[c0, c0] {leadDimension = 16 : index} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "AOp">

The MMAMatrixType describes the shape of the matrix being loaded and the operand being loaded too. The operand needs to be specified to aid the lowering of this type to dialects such as NVVM where each workitem may hold different amount of elements depending on the elementType of the matrix. For e.g., Each workitem holds 4 vector<2xf16>s for f16 data type and 8 f32s for f32 data type of MMAMatrix. Some other instances of usage are:-

%3 = gpu.subgroup_mma_compute %0, %1, %2 : !gpu.mma_matrix<16x16xf16, "AOp">, !gpu.mma_matrix<16x16xf16, "BOp"> -> !gpu.mma_matrix<16x16xf32, "COp">

gpu.subgroup_mma_store_matrix %3, arg22[c0, c0] {leadDimension = 16 : index}: !gpu.mma_matrix<16x16xf32, "COp">, memref<16x16xf32>

Definition at line 122 of file GPUDialect.h.

Member Function Documentation

◆ get()

MMAMatrixType MMAMatrixType::get ( ArrayRef< int64_t >  shape,
Type  elementType,
StringRef  operand 
)
static

◆ getChecked()

MMAMatrixType MMAMatrixType::getChecked ( function_ref< InFlightDiagnostic()>  emitError,
ArrayRef< int64_t >  shape,
Type  elementType,
StringRef  operand 
)
static

Get MMAMatrixType at a particular location and verify construction Invariants.

Definition at line 46 of file GPUDialect.cpp.

References mlir::emitError(), mlir::detail::StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits >::getChecked(), and mlir::Type::getContext().

◆ getElementType()

Type MMAMatrixType::getElementType ( ) const

Get elementType of a single element.

Definition at line 59 of file GPUDialect.cpp.

References mlir::detail::StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits >::getImpl().

◆ getNumDims()

unsigned MMAMatrixType::getNumDims ( ) const

◆ getOperand()

StringRef MMAMatrixType::getOperand ( ) const

The general form of operation this type supports is given by the equation C += A*B.

This function returns which operand in the given equation is held by this type. String returned can be one of"AOp", "BOp" and "COp".

Definition at line 61 of file GPUDialect.cpp.

References mlir::detail::StorageUserBase< ConcreteT, BaseT, StorageT, UniquerT, Traits >::getImpl().

Referenced by mlir::convertMMAToLLVMType().

◆ getShape()

ArrayRef< int64_t > MMAMatrixType::getShape ( ) const

◆ isValidElementType()

bool MMAMatrixType::isValidElementType ( Type  elementType)
static

Check if a type is valid a MMAMatrixType elementType.

Definition at line 63 of file GPUDialect.cpp.

References mlir::Type::isF16(), and mlir::Type::isF32().

Referenced by verify().

◆ verify()

LogicalResult MMAMatrixType::verify ( function_ref< InFlightDiagnostic()>  emitError,
ArrayRef< int64_t >  shape,
Type  elementType,
StringRef  operand 
)
static

Verify that shape and elementType are actually allowed for the MMAMatrixType.

Definition at line 68 of file GPUDialect.cpp.

References mlir::emitError(), isValidElementType(), and mlir::success().


The documentation for this class was generated from the following files: