VectorToGPU.cpp File Reference

#include "mlir/Conversion/VectorToGPU/VectorToGPU.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/SCF/IR/SCF.h"
#include "mlir/Dialect/Utils/StructuredOpsUtils.h"
#include "mlir/Dialect/Vector/IR/VectorOps.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::ConvertVectorToGPUPassBase< 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-gpu"
#define	GEN_PASS_DEF_CONVERTVECTORTOGPUPASS

Functions
std::unique_ptr<::mlir::Pass >	mlir::impl::createConvertVectorToGPUPass ()
std::unique_ptr<::mlir::Pass >	mlir::createConvertVectorToGPUPass ()
static bool	contractSupportsMMAMatrixType (vector::ContractionOp contract)
static bool	isFirstResultLastMapDimension (AffineMap permutationMap)
static std::optional< int64_t >	getStaticallyKnownRowStride (ShapedType type, AffineMap permutationMap)
static bool	transferReadSupportsMMAMatrixType (vector::TransferReadOp readOp)
static bool	transferWriteSupportsMMAMatrixType (vector::TransferWriteOp writeOp)
static bool	supportsMMaMatrixType (Operation *op)
static SetVector< Operation * >	getOpToConvert (mlir::Operation *op)
static const char *	inferFragType (Operation *op)
static LogicalResult	convertTransferReadOp (RewriterBase &rewriter, vector::TransferReadOp op, llvm::DenseMap< Value, Value > &valueMapping)
static LogicalResult	convertTransferWriteOp (RewriterBase &rewriter, vector::TransferWriteOp op, llvm::DenseMap< Value, Value > &valueMapping)
static LogicalResult	convertContractOp (RewriterBase &rewriter, vector::ContractionOp op, llvm::DenseMap< Value, Value > &valueMapping)
static LogicalResult	convertConstantOp (RewriterBase &rewriter, arith::ConstantOp op, llvm::DenseMap< Value, Value > &valueMapping)
	Convert a 2D splat ConstantOp to a SubgroupMmaConstantMatrix op.
static LogicalResult	convertBroadcastOp (RewriterBase &rewriter, vector::BroadcastOp op, llvm::DenseMap< Value, Value > &valueMapping)
	Convert a vector.broadcast from scalar to a SubgroupMmaConstantMatrix op.
static LogicalResult	convertElementwiseOp (RewriterBase &rewriter, Operation *op, gpu::MMAElementwiseOp opType, llvm::DenseMap< Value, Value > &valueMapping)
	Convert an elementwise op to the equivalent elementwise op on MMA matrix.

Macro Definition Documentation

DEBUG_TYPE

#define DEBUG_TYPE   "vector-to-gpu"

Definition at line 33 of file VectorToGPU.cpp.

GEN_PASS_DEF_CONVERTVECTORTOGPUPASS

#define GEN_PASS_DEF_CONVERTVECTORTOGPUPASS

Definition at line 36 of file VectorToGPU.cpp.

Function Documentation

contractSupportsMMAMatrixType()

bool contractSupportsMMAMatrixType ( vector::ContractionOp contract )

static

Definition at line 43 of file VectorToGPU.cpp.

References mlir::bindDims(), contract(), mlir::AffineMap::inferFromExprList(), mlir::vector::isParallelIterator(), and mlir::vector::isReductionIterator().

Referenced by supportsMMaMatrixType().

convertBroadcastOp()

LogicalResult convertBroadcastOp ( RewriterBase & rewriter, vector::BroadcastOp op, llvm::DenseMap< Value, Value > & valueMapping )

static

Convert a vector.broadcast from scalar to a SubgroupMmaConstantMatrix op.

Definition at line 536 of file VectorToGPU.cpp.

References mlir::broadcastSupportsMMAMatrixType(), mlir::gpu::MMAMatrixType::get(), inferFragType(), mlir::OpBuilder::setInsertionPoint(), and success().

Referenced by mlir::convertVectorToMMAOps().

convertConstantOp()

LogicalResult convertConstantOp ( RewriterBase & rewriter, arith::ConstantOp op, llvm::DenseMap< Value, Value > & valueMapping )

static

Convert a 2D splat ConstantOp to a SubgroupMmaConstantMatrix op.

Definition at line 513 of file VectorToGPU.cpp.

References mlir::constantSupportsMMAMatrixType(), mlir::gpu::MMAMatrixType::get(), inferFragType(), mlir::OpBuilder::setInsertionPoint(), and success().

Referenced by mlir::convertVectorToMMAOps().

convertContractOp()

LogicalResult convertContractOp ( RewriterBase & rewriter, vector::ContractionOp op, llvm::DenseMap< Value, Value > & valueMapping )

static

Definition at line 491 of file VectorToGPU.cpp.

References mlir::Builder::getType(), mlir::RewriterBase::notifyMatchFailure(), mlir::OpBuilder::setInsertionPoint(), and success().

Referenced by mlir::convertVectorToMMAOps().

convertElementwiseOp()

LogicalResult convertElementwiseOp ( RewriterBase & rewriter, Operation * op, gpu::MMAElementwiseOp opType, llvm::DenseMap< Value, Value > & valueMapping )

static

Convert an elementwise op to the equivalent elementwise op on MMA matrix.

Definition at line 555 of file VectorToGPU.cpp.

References mlir::gpu::MMAMatrixType::get(), mlir::Operation::getLoc(), mlir::Operation::getOperands(), mlir::Operation::getResult(), mlir::Operation::getResultTypes(), mlir::getType(), mlir::RewriterBase::notifyMatchFailure(), mlir::OpBuilder::setInsertionPoint(), and success().

Referenced by mlir::convertVectorToMMAOps().

convertTransferReadOp()

LogicalResult convertTransferReadOp ( RewriterBase & rewriter, vector::TransferReadOp op, llvm::DenseMap< Value, Value > & valueMapping )

static

Definition at line 401 of file VectorToGPU.cpp.

References mlir::gpu::MMAMatrixType::get(), mlir::Builder::getIndexAttr(), mlir::AffineMap::getPermutationMap(), getStaticallyKnownRowStride(), mlir::Builder::getUnitAttr(), inferFragType(), isFirstResultLastMapDimension(), load, mlir::RewriterBase::notifyMatchFailure(), mlir::OpBuilder::setInsertionPoint(), success(), and transferReadSupportsMMAMatrixType().

Referenced by mlir::convertVectorToMMAOps().

convertTransferWriteOp()

LogicalResult convertTransferWriteOp ( RewriterBase & rewriter, vector::TransferWriteOp op, llvm::DenseMap< Value, Value > & valueMapping )

static

Definition at line 451 of file VectorToGPU.cpp.

References mlir::RewriterBase::eraseOp(), mlir::Builder::getIndexAttr(), mlir::AffineMap::getPermutationMap(), getStaticallyKnownRowStride(), mlir::Builder::getUnitAttr(), isFirstResultLastMapDimension(), mlir::RewriterBase::notifyMatchFailure(), mlir::OpBuilder::setInsertionPoint(), success(), and transferWriteSupportsMMAMatrixType().

Referenced by mlir::convertVectorToMMAOps().

getOpToConvert()

SetVector< Operation * > getOpToConvert ( mlir::Operation * op )

static

Definition at line 210 of file VectorToGPU.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::convertVectorToMMAOps().

getStaticallyKnownRowStride()

std::optional< int64_t > getStaticallyKnownRowStride ( ShapedType type, AffineMap permutationMap )

static

Definition at line 93 of file VectorToGPU.cpp.

References mlir::AffineMap::getNumResults(), mlir::AffineMap::getResults(), and result.

Referenced by convertTransferReadOp(), convertTransferWriteOp(), transferReadSupportsMMAMatrixType(), and transferWriteSupportsMMAMatrixType().

inferFragType()

const char * inferFragType ( Operation * op )

static

Definition at line 380 of file VectorToGPU.cpp.

References contract(), mlir::Operation::getNumResults(), mlir::Operation::getResult(), mlir::Operation::getUsers(), mlir::Operation::hasOneUse(), mlir::Operation::hasTrait(), inferFragType(), and mlir::Operation::user_begin().

Referenced by convertBroadcastOp(), convertConstantOp(), convertTransferReadOp(), and inferFragType().

isFirstResultLastMapDimension()

bool isFirstResultLastMapDimension ( AffineMap permutationMap )

static

Definition at line 70 of file VectorToGPU.cpp.

References mlir::getAffineDimExpr(), mlir::AffineMap::getContext(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getResult(), and mlir::AffineMap::getResults().

Referenced by convertTransferReadOp(), and convertTransferWriteOp().

supportsMMaMatrixType()

bool supportsMMaMatrixType ( Operation * op )

static

Definition at line 182 of file VectorToGPU.cpp.

References broadcast(), mlir::broadcastSupportsMMAMatrixType(), mlir::constantSupportsMMAMatrixType(), contract(), contractSupportsMMAMatrixType(), mlir::elementwiseSupportsMMAMatrixType(), mlir::integerExtendSupportsMMAMatrixType< arith::ExtSIOp >(), mlir::integerExtendSupportsMMAMatrixType< arith::ExtUIOp >(), transferReadSupportsMMAMatrixType(), and transferWriteSupportsMMAMatrixType().

Referenced by getOpToConvert().

transferReadSupportsMMAMatrixType()

bool transferReadSupportsMMAMatrixType ( vector::TransferReadOp readOp )

static

Definition at line 139 of file VectorToGPU.cpp.

References mlir::getAffineDimExpr(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getPermutationMap(), mlir::AffineMap::getResults(), and getStaticallyKnownRowStride().

Referenced by convertTransferReadOp(), and supportsMMaMatrixType().

transferWriteSupportsMMAMatrixType()

bool transferWriteSupportsMMAMatrixType ( vector::TransferWriteOp writeOp )

static

Definition at line 161 of file VectorToGPU.cpp.

References mlir::getAffineDimExpr(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getPermutationMap(), mlir::AffineMap::getResults(), and getStaticallyKnownRowStride().

Referenced by convertTransferWriteOp(), and supportsMMaMatrixType().