#include "mlir/Conversion/VectorToNVGPU/VectorToNVGPU.h"#include "mlir/Analysis/SliceAnalysis.h"#include "mlir/Analysis/TopologicalSortUtils.h"#include "mlir/Conversion/VectorToGPUCommon/VectorToGPUUtils.h"#include "mlir/Dialect/Affine/IR/AffineOps.h"#include "mlir/Dialect/Arith/IR/Arith.h"#include "mlir/Dialect/GPU/IR/GPUDialect.h"#include "mlir/Dialect/MemRef/IR/MemRef.h"#include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h"#include "mlir/Dialect/NVGPU/Utils/MMAUtils.h"#include "mlir/Dialect/SCF/IR/SCF.h"#include "mlir/Dialect/Utils/StructuredOpsUtils.h"#include "mlir/Dialect/Vector/IR/VectorOps.h"#include "mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h"#include "mlir/IR/Builders.h"#include "mlir/IR/Region.h"#include "mlir/Pass/Pass.h"#include "mlir/Transforms/GreedyPatternRewriteDriver.h"#include "llvm/ADT/STLExtras.h"#include "llvm/ADT/TypeSwitch.h"#include "llvm/Support/DebugLog.h"#include "mlir/Conversion/Passes.h.inc"Go to the source code of this file.
Classes | |
| class | mlir::impl::ConvertVectorToNVGPUPassBase< DerivedT > |
Namespaces | |
| namespace | mlir |
| Include the generated interface declarations. | |
| namespace | mlir::impl |
| Attribute collections provide a dictionary-like interface. | |
Macros | |
| #define | DEBUG_TYPE "vector-to-nvgpu" |
| #define | GEN_PASS_DEF_CONVERTVECTORTONVGPUPASS |
Functions | |
| std::unique_ptr<::mlir::Pass > | mlir::impl::createConvertVectorToNVGPUPass () |
| std::unique_ptr<::mlir::Pass > | mlir::createConvertVectorToNVGPUPass () |
| template<typename TransferOpType> | |
| static void | getXferIndices (RewriterBase &rewriter, TransferOpType xferOp, AffineMap offsetMap, ArrayRef< Value > dimValues, SmallVector< Value, 4 > &indices) |
| For a vector TransferOpType xferOp, an empty indices vector, and an AffineMap representing offsets to apply to indices, the function fills indices with the original indices plus the offsets. | |
| static bool | contractSupportsMMAMatrixType (vector::ContractionOp contract) |
| static bool | extractStridedSliceSupportsMMAMatrixType (vector::ExtractStridedSliceOp op) |
| Returns true if the extract strided slice op is supported with mma.sync path. | |
| static bool | supportsMMaMatrixType (Operation *op) |
| static SetVector< Operation * > | getOpToConvert (mlir::Operation *op) |
| static VectorType | getMmaSyncVectorOperandType (const nvgpu::FragmentElementInfo ®Info) |
| Returns the vector type which represents a matrix fragment. | |
| static LogicalResult | convertConstantOpMmaSync (RewriterBase &rewriter, arith::ConstantOp op, llvm::DenseMap< Value, Value > &valueMapping) |
| Convert a 2D splat ConstantOp to a SubgroupMmaConstantMatrix op. | |
| static FailureOr< bool > | isTransposed (vector::TransferReadOp op) |
| Check if the loaded matrix operand requires transposed. | |
| static LogicalResult | creatLdMatrixCompatibleLoads (RewriterBase &rewriter, vector::TransferReadOp op, llvm::DenseMap< Value, Value > &valueMapping) |
| static LogicalResult | createNonLdMatrixLoads (RewriterBase &rewriter, vector::TransferReadOp op, llvm::DenseMap< Value, Value > &valueMapping) |
| static bool | isSharedMemory (MemRefType type) |
| Return true if this is a shared memory memref type. | |
| static LogicalResult | convertTransferReadToLoads (RewriterBase &rewriter, vector::TransferReadOp op, llvm::DenseMap< Value, Value > &valueMapping) |
| Converts a vector.transfer_read operation directly to either a vector.load or a nvgpu.ldmatrix operation. | |
| static LogicalResult | convertTransferWriteToStores (RewriterBase &rewriter, vector::TransferWriteOp op, llvm::DenseMap< Value, Value > &valueMapping) |
| static void | populateFromInt64AttrArray (ArrayAttr arrayAttr, SmallVectorImpl< int64_t > &results) |
| static LogicalResult | convertExtractStridedSlice (RewriterBase &rewriter, vector::ExtractStridedSliceOp op, llvm::DenseMap< Value, Value > &valueMapping) |
| static LogicalResult | convertContractOpToMmaSync (RewriterBase &rewriter, vector::ContractionOp op, llvm::DenseMap< Value, Value > &valueMapping) |
Macro Definition Documentation
◆ DEBUG_TYPE
| #define DEBUG_TYPE "vector-to-nvgpu" |
Definition at line 37 of file VectorToNVGPU.cpp.
◆ GEN_PASS_DEF_CONVERTVECTORTONVGPUPASS
| #define GEN_PASS_DEF_CONVERTVECTORTONVGPUPASS |
Definition at line 40 of file VectorToNVGPU.cpp.
Function Documentation
◆ contractSupportsMMAMatrixType()
|
static |
Definition at line 73 of file VectorToNVGPU.cpp.
References mlir::bindDims(), contract(), mlir::AffineMap::inferFromExprList(), mlir::vector::isParallelIterator(), and mlir::vector::isReductionIterator().
Referenced by supportsMMaMatrixType().
◆ convertConstantOpMmaSync()
|
static |
Convert a 2D splat ConstantOp to a SubgroupMmaConstantMatrix op.
Definition at line 261 of file VectorToNVGPU.cpp.
References mlir::DenseElementsAttr::get(), getMmaSyncVectorOperandType(), mlir::nvgpu::getWarpMatrixInfo(), mlir::RewriterBase::notifyMatchFailure(), result, mlir::OpBuilder::setInsertionPoint(), and success().
Referenced by mlir::convertVectorToNVVMCompatibleMMASync().
◆ convertContractOpToMmaSync()
|
static |
Definition at line 648 of file VectorToNVGPU.cpp.
References mlir::Builder::getI64ArrayAttr(), mlir::RewriterBase::notifyMatchFailure(), mlir::OpBuilder::setInsertionPoint(), and success().
Referenced by mlir::convertVectorToNVVMCompatibleMMASync().
◆ convertExtractStridedSlice()
|
static |
Definition at line 572 of file VectorToNVGPU.cpp.
References mlir::nvgpu::getWarpMatrixInfo(), mlir::RewriterBase::notifyMatchFailure(), populateFromInt64AttrArray(), mlir::OpBuilder::setInsertionPoint(), and success().
Referenced by mlir::convertVectorToNVVMCompatibleMMASync().
◆ convertTransferReadToLoads()
|
static |
Converts a vector.transfer_read operation directly to either a vector.load or a nvgpu.ldmatrix operation.
This function should only be used when converting to nvgpu.mma.sync operations.
Definition at line 488 of file VectorToNVGPU.cpp.
References createNonLdMatrixLoads(), creatLdMatrixCompatibleLoads(), mlir::nvgpu::getWarpMatrixInfo(), isSharedMemory(), mlir::RewriterBase::notifyMatchFailure(), and mlir::OpBuilder::setInsertionPoint().
Referenced by mlir::convertVectorToNVVMCompatibleMMASync().
◆ convertTransferWriteToStores()
|
static |
Definition at line 520 of file VectorToNVGPU.cpp.
References mlir::RewriterBase::eraseOp(), mlir::Builder::getIndexAttr(), mlir::Builder::getIndexType(), getMmaSyncVectorOperandType(), mlir::nvgpu::getWarpMatrixInfo(), getXferIndices(), mlir::RewriterBase::notifyMatchFailure(), mlir::OpBuilder::setInsertionPoint(), and success().
Referenced by mlir::convertVectorToNVVMCompatibleMMASync().
◆ createNonLdMatrixLoads()
|
static |
Definition at line 389 of file VectorToNVGPU.cpp.
References mlir::Builder::getIndexAttr(), mlir::Builder::getIndexType(), getMmaSyncVectorOperandType(), mlir::nvgpu::getWarpMatrixInfo(), getXferIndices(), mlir::Builder::getZeroAttr(), mlir::RewriterBase::notifyMatchFailure(), result, mlir::OpBuilder::setInsertionPoint(), and success().
Referenced by convertTransferReadToLoads().
◆ creatLdMatrixCompatibleLoads()
|
static |
Definition at line 328 of file VectorToNVGPU.cpp.
References getMmaSyncVectorOperandType(), mlir::nvgpu::getWarpMatrixInfo(), getXferIndices(), indices, isTransposed(), mlir::RewriterBase::notifyMatchFailure(), mlir::OpBuilder::setInsertionPoint(), and success().
Referenced by convertTransferReadToLoads().
◆ extractStridedSliceSupportsMMAMatrixType()
|
static |
Returns true if the extract strided slice op is supported with mma.sync path.
Definition at line 97 of file VectorToNVGPU.cpp.
References mlir::nvgpu::B, mlir::nvgpu::C, mlir::nvgpu::getUserContract(), and mlir::nvgpu::getWarpMatrixInfo().
Referenced by supportsMMaMatrixType().
◆ getMmaSyncVectorOperandType()
|
static |
Returns the vector type which represents a matrix fragment.
Definition at line 250 of file VectorToNVGPU.cpp.
Referenced by convertConstantOpMmaSync(), convertTransferWriteToStores(), createNonLdMatrixLoads(), and creatLdMatrixCompatibleLoads().
◆ getOpToConvert()
|
static |
Definition at line 151 of file VectorToNVGPU.cpp.
References mlir::elementwiseSupportsMMAMatrixType(), mlir::SliceOptions::filter, mlir::getBackwardSlice(), mlir::getForwardSlice(), mlir::Operation::getOperandTypes(), mlir::Operation::getResultTypes(), inserted(), result, supportsMMaMatrixType(), mlir::topologicalSort(), and mlir::Operation::walk().
Referenced by mlir::convertVectorToNVVMCompatibleMMASync().
◆ getXferIndices()
|
static |
For a vector TransferOpType xferOp, an empty indices vector, and an AffineMap representing offsets to apply to indices, the function fills indices with the original indices plus the offsets.
The offsets are applied by taking into account the permutation map of the transfer op. If the offsetMap has dimension placeholders, those should be provided in dimValues.
Definition at line 53 of file VectorToNVGPU.cpp.
References mlir::Builder::getAffineDimExpr(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getResult(), indices, and mlir::affine::makeComposedAffineApply().
Referenced by convertTransferWriteToStores(), createNonLdMatrixLoads(), and creatLdMatrixCompatibleLoads().
◆ isSharedMemory()
|
static |
Return true if this is a shared memory memref type.
Definition at line 477 of file VectorToNVGPU.cpp.
Referenced by convertTransferReadToLoads().
◆ isTransposed()
|
static |
Check if the loaded matrix operand requires transposed.
Transposed Map Example: Example 1 : (..., d0, d1) -> (d1 * 1, d0 * 2) Example 2 : (d0, d1, d2, d3) -> (d3, d2) The code below checks if the output 2D is transposed using a generalized version : (d0, d1, dn, ..., dm, ...) -> (dm, dn) Returns : true; if m > n, false o.w.
Definition at line 301 of file VectorToNVGPU.cpp.
References mlir::AffineMap::getNumResults(), mlir::AffineMap::getPermutationMap(), and mlir::AffineMap::getResult().
Referenced by creatLdMatrixCompatibleLoads().
◆ populateFromInt64AttrArray()
|
static |
Definition at line 565 of file VectorToNVGPU.cpp.
References ArrayAttr().
Referenced by convertExtractStridedSlice().
◆ supportsMMaMatrixType()
Definition at line 121 of file VectorToNVGPU.cpp.
References broadcast(), mlir::broadcastSupportsMMAMatrixType(), mlir::nvgpu::canLowerToWarpMatrixOperation(), mlir::constantSupportsMMAMatrixType(), contract(), contractSupportsMMAMatrixType(), mlir::elementwiseSupportsMMAMatrixType(), extractStridedSliceSupportsMMAMatrixType(), mlir::integerExtendSupportsMMAMatrixType< arith::ExtSIOp >(), and mlir::integerExtendSupportsMMAMatrixType< arith::ExtUIOp >().
Referenced by getOpToConvert().
Generated on for MLIR by
1.14.0