MLIR  16.0.0git
Namespaces | Functions
NvGpuSupport.cpp File Reference
#include "NvGpuSupport.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
+ Include dependency graph for NvGpuSupport.cpp:

Go to the source code of this file.

Namespaces

 mlir
 Include the generated interface declarations.
 
 mlir::nvgpu
 

Functions

FailureOr< WarpMatrixInfo > mlir::nvgpu::getWarpMatrixInfo (Operation *op)
 Given an op that operates on a VectorType representing a warp-level matrix operand, the function returns a struct containing relevant type information. More...
 
int64_t mlir::nvgpu::inferTileWidthInBits (const WarpMatrixInfo &type)
 Returns the number of bits in a single tile row. More...
 
FailureOr< FragmentElementInfo > mlir::nvgpu::getMmaSyncRegisterType (const WarpMatrixInfo &type)
 Returns a FragmentElementInfo struct describing the register types for the given matrix fragment type. More...
 
static AffineMap mlir::nvgpu::getRegisterIndexToTileOffsetMap (int64_t lineSize, Type elementType, ArrayRef< int64_t > operandShape, bool isAccumulator, int64_t elementsPerRegister, AffineExpr logicalValueId)
 
FailureOr< AffineMap > mlir::nvgpu::getLaneIdAndValueIdToOperandCoord (Location loc, OpBuilder &builder, 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< nvgpu::LdMatrixParams > mlir::nvgpu::getLdMatrixParams (const WarpMatrixInfo &type, bool transpose)
 
FailureOr< AffineMap > mlir::nvgpu::getLaneIdToLdMatrixMatrixCoord (Location loc, OpBuilder &builder, 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...