|
FailureOr< vector::ContractionOp > | mlir::nvgpu::getUserContract (Operation *op) |
| Returns the first user of the op that is vector.contract. More...
|
|
FailureOr< WarpMatrixInfo > | mlir::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< FragmentElementInfo > | mlir::nvgpu::getMmaSyncRegisterType (const WarpMatrixInfo &type) |
| Returns a FragmentElementInfo struct describing the register types for the given matrix fragment type. More...
|
|
FailureOr< AffineMap > | mlir::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< LdMatrixParams > | mlir::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< AffineMap > | mlir::nvgpu::getLaneIdToLdMatrixMatrixCoord (OpBuilder &builder, Location loc, const LdMatrixParams ¶ms) |
| 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...
|
|