MLIR  20.0.0git
Classes | Namespaces | Enumerations | Functions
MMAUtils.h File Reference
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/IR/Types.h"

Go to the source code of this file.

Classes

struct  mlir::nvgpu::WarpMatrixInfo
 Collects information about a warp-level matrix operand represented by a VectorType. More...
 
struct  mlir::nvgpu::FragmentElementInfo
 Specifies information about the registers which compose a matrix fragment according to the PTX documentation. More...
 
struct  mlir::nvgpu::LdMatrixParams
 Encapsulates the parameters needed to lower a nvgpu.ldmatrix operation to nvvm.ldmatrix. More...
 

Namespaces

 mlir
 Include the generated interface declarations.
 
 mlir::nvgpu
 

Enumerations

enum class  mlir::nvgpu::MatMulOperandRole : int32_t { mlir::nvgpu::A = 0 , mlir::nvgpu::B , mlir::nvgpu::C }
 Represents the role of an operand in an MMA instruction: result := matmul(A, B) + C More...
 

Functions

FailureOr< vector::ContractionOp > mlir::nvgpu::getUserContract (Operation *op)
 Returns the first user of the op that is vector.contract. More...
 
FailureOr< WarpMatrixInfomlir::nvgpu::getWarpMatrixInfo (Operation *op)
 If op is a vector.transfer_write, return the WarpMatrixInfo for the vector operand. More...
 
int64_t mlir::nvgpu::inferTileWidthInBits (const WarpMatrixInfo &type)
 Returns the number of bits in a single tile row. More...
 
FailureOr< FragmentElementInfomlir::nvgpu::getMmaSyncRegisterType (const WarpMatrixInfo &type)
 Returns a FragmentElementInfo struct describing the register types for the given matrix fragment type. More...
 
FailureOr< AffineMapmlir::nvgpu::getLaneIdAndValueIdToOperandCoord (OpBuilder &builder, Location loc, const WarpMatrixInfo &fragmentType)
 Returns an AffineMap which maps a two dimensions representing (laneId, logicalValueId) and returns two results representing offsets within a matrix operand. More...
 
FailureOr< LdMatrixParamsmlir::nvgpu::getLdMatrixParams (const WarpMatrixInfo &type, bool transpose)
 Given type that contains info for a warp-matrix operand and whether or not the load is a transposed load, return the LdMatrixParams. More...
 
FailureOr< AffineMapmlir::nvgpu::getLaneIdToLdMatrixMatrixCoord (OpBuilder &builder, Location loc, const LdMatrixParams &params)
 Returns an AffineMap which maps a single dimension representing the laneId to two results representing offsets within the matrix operand that should be the pointer locations a thread should pass to the ldmatrix instruction. More...
 
bool mlir::nvgpu::canLowerToWarpMatrixOperation (vector::TransferReadOp op)
 Returns whether the vector.transfer_read instruction can be interpreted as a warp-level cooperative matrix load operation. More...
 
bool mlir::nvgpu::canLowerToWarpMatrixOperation (vector::TransferWriteOp op)
 Returns whether the vector.transfer_write instruction can be interpreted as a warp-level cooperative matrix store operation. More...