LinalgTransformOps.cpp File Reference

#include "mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.h"
#include "mlir/AsmParser/AsmParser.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Arith/Utils/Utils.h"
#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
#include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h"
#include "mlir/Dialect/CommonFolders.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
#include "mlir/Dialect/Linalg/TransformOps/GPUHeuristics.h"
#include "mlir/Dialect/Linalg/TransformOps/Syntax.h"
#include "mlir/Dialect/Linalg/Transforms/Hoisting.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
#include "mlir/Dialect/Linalg/Utils/Utils.h"
#include "mlir/Dialect/SCF/Transforms/TileUsingInterface.h"
#include "mlir/Dialect/Tensor/IR/Tensor.h"
#include "mlir/Dialect/Transform/IR/TransformTypes.h"
#include "mlir/Dialect/Transform/Interfaces/TransformInterfaces.h"
#include "mlir/Dialect/Transform/Utils/Utils.h"
#include "mlir/Dialect/UB/IR/UBOps.h"
#include "mlir/Dialect/Utils/IndexingUtils.h"
#include "mlir/Dialect/Utils/StaticValueUtils.h"
#include "mlir/Dialect/Vector/Transforms/LoweringPatterns.h"
#include "mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h"
#include "mlir/IR/BuiltinTypeInterfaces.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/IR/TypeUtilities.h"
#include "mlir/Interfaces/ParallelCombiningOpInterface.h"
#include "mlir/Interfaces/TilingInterface.h"
#include "mlir/Support/LLVM.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/ScopeExit.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/TypeSwitch.h"
#include "llvm/Support/DebugLog.h"
#include "llvm/Support/LogicalResult.h"
#include <type_traits>
#include "mlir/Dialect/Linalg/TransformOps/LinalgTransformOpsEnums.cpp.inc"
#include "mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp.inc"

Go to the source code of this file.

Macros
#define	DEBUG_TYPE   "linalg-transforms"
#define	DOWNSCALE(trans)
#define	DOWNSCALE_CALL(a, b)
#define	DOWNSCALE_NORMAL(a, b)
#define	GET_OP_CLASSES

Functions
template<typename PatternTy, typename... Args>
static FailureOr< LinalgOp >	tryApply (Operation *operation, Args &&...args)
	Attempts to apply the pattern specified as template argument to the given operation.
static DiagnosedSilenceableFailure	unpackSingleIndexResultPayloadOperations (transform::TransformState &state, TransformOpInterface transformOp, SmallVector< OpFoldResult > &result, ArrayRef< OpFoldResult > ofrs)
	Assuming that ofr is an index attr or a param of index type or a transform dialect handle mapped to exactly one op with one index result, return that value.
static DiagnosedSilenceableFailure	unpackSingleIndexResultPayloadOperations (transform::TransformState &state, TransformOpInterface transformOp, SmallVector< OpFoldResult > &result, Value packedHandle)
static DiagnosedSilenceableFailure	reifyMixedParamAndHandleResults (TransformState &state, TransformOpInterface &transformOp, ArrayRef< OpFoldResult > mixedResults, SmallVectorImpl< int64_t > &reified)
	When possible, converts each OpFoldResult in mixedResult to an integer if the value can be statically inferred.
static bool	mayBeRead (OpOperand &operand)
	Return true if the operand may be read from by its owner.
static bool	mayBeRead (Value value)
	Return true if the value may be read through any of its uses.
template<typename Range>
static LogicalResult	applyTilingToAll (RewriterBase &rewriter, Operation *transformOp, Range &&payloadOps, unsigned numLoops, transform::TransformResults &transformResults, function_ref< FailureOr< scf::SCFTileAndFuseResult >(TilingInterface)> applyFn)
	Apply a tiling transformation to all payload ops and store both the tiled operation as well as the created tile loops.
static Operation *	replaceForAllWithNewSignature (RewriterBase &rewriter, Diagnostic &diag, Operation *producerOp, Operation *containingOp, TilingResult &tileAndFuseResult, int64_t resultNumber, SmallVector< OpFoldResult > &offsets, SmallVector< OpFoldResult > &sizes)
	Add new operands to the forall op for users of the producerOp that are dominated by the containing scf.forall op.
static bool	sameOrEquivalentIterArg (Value src, Value dst)
	Given two operands coming from a loop iter arg, 'src' and 'dst', return true if the operand 'src' is equal to 'dst' or equal to a iter arg present in a outer loop.
static std::tuple< SmallVector< Operation * >, Operation * >	tileAndFuseFirstExtractUse (RewriterBase &rewriter, Diagnostic &diag, Operation *producerOp, Operation *containingOp)
	Find the first "extract" user of producerOp and tile it right before its use.
static SmallVector< Operation * >	tileAndFuseFirstExtractUseThroughContainingOpBlockArgument (RewriterBase &rewriter, Diagnostic &diag, Operation *producerOp, Operation *containingOp)
	First, find the first "scf::ForallOp" user of producerOp and ensure it is exactly the containingOp, otherwise bail.
static Operation *	cloneAndFuseFirstUse (RewriterBase &rewriter, Diagnostic &diag, Operation *producerOp, Operation *containingOp)
static void	printMultitileSizesTypes (OpAsmPrinter &printer, Operation *op, Type targetType, Type lowSizeType, Type, Type)
static ParseResult	parseMultitileSizesTypes (OpAsmParser &parser, Type &targetType, Type &lowSizeType, Type &highSizeType, Type &splitPointType)
	ArrayAttr ()
b	getI64ArrayAttr (paddingDimensions)
static void	printContinuousTileSizeTypes (OpAsmPrinter &printer, Operation *op, Type targetType, Type tileSizes, Type)
static ParseResult	parseContinuousTileSizeTypes (OpAsmParser &parser, Type &targetType, Type &tileSizesType, Type &chunkSizesType)
static SmallVector< OpFoldResult >	normalizeUpperBounds (RewriterBase &rewriter, Location loc, ArrayRef< OpFoldResult > lbs, ArrayRef< OpFoldResult > ubs, ArrayRef< OpFoldResult > steps)
	Given lbs, ubs and steps of loops, return (for each loop), the normalized upper bound.
static SmallVector< Value >	denormalizeIndVar (RewriterBase &rewriter, Location loc, ValueRange ivs, ArrayRef< OpFoldResult > lbs, ArrayRef< OpFoldResult > steps)
	When a loop is normalized, the uses of the induction variable within the loop need to replaced with original_lb + old_iv * original_step.
static scf::ForallOp	normalizeForallLoopOp (RewriterBase &rewriter, scf::ForallOp loop)
	Given a scf.forall loop return a loop op with the loop bounds normalized. TODO: Replace this with a general utility to normalize scf.forall. At the time of writing, this wasnt done since adding this to scf dialect would disallow using of affine.apply operations due to cyclic dependencies. To avoid churn in lit tests with the change this was added with, defer that to a follow up.
template<typename OpTy>
static DiagnosedSilenceableFailure	doit (RewriterBase &rewriter, OpTy target, transform::ApplyToEachResultList &results, transform::TransformState &state)

Variables
	b
	Return true if permutation is a valid permutation of the outer_dims_perm (case OuterOrInnerPerm::Outer) or inner_dims_pos (OuterOrInnerPerm::Inner) of the tensor.pack or tensor.unpack op.
	result
	TypeRange {resultType, resultType}
	target
b	ValueRange {}

Macro Definition Documentation

DEBUG_TYPE

#define DEBUG_TYPE   "linalg-transforms"

Definition at line 55 of file LinalgTransformOps.cpp.

DOWNSCALE

#define DOWNSCALE

(

trans

)

Value:

  {                                                                            \
    FailureOr<LinalgOp> res = tryApply<trans>(target);                         \
    if (succeeded(res)) {                                                      \
      results.push_back(*res);                                                 \
      return DiagnosedSilenceableFailure::success();                           \
    }                                                                          \
  }

DOWNSCALE_CALL

#define DOWNSCALE_CALL	(		a,
			b )

Value:

DownscaleSizeOneWindowed2DConvolution<a, b>

DOWNSCALE_NORMAL

#define DOWNSCALE_NORMAL	(		a,
			b )

Value:

DOWNSCALE(DOWNSCALE_CALL(a, b))

GET_OP_CLASSES

#define GET_OP_CLASSES

Definition at line 4529 of file LinalgTransformOps.cpp.

Function Documentation

applyTilingToAll()

template<typename Range>

LogicalResult applyTilingToAll ( RewriterBase & rewriter, Operation * transformOp, Range && payloadOps, unsigned numLoops, transform::TransformResults & transformResults, function_ref< FailureOr< scf::SCFTileAndFuseResult >(TilingInterface)> applyFn )

static

Apply a tiling transformation to all payload ops and store both the tiled operation as well as the created tile loops.

Definition at line 664 of file LinalgTransformOps.cpp.

References mlir::Operation::emitError(), mlir::RewriterBase::eraseOp(), mlir::Operation::getOpResult(), push_back(), mlir::RewriterBase::replaceAllUsesWith(), replacement(), mlir::transform::TransformResults::set(), mlir::OpBuilder::setInsertionPoint(), success(), and target.

ArrayAttr()

ArrayAttr

(

)

Referenced by adjustIter(), appendCallOpAliasScopes(), mlir::spirv::attributeName< FPFastMathMode >(), mlir::spirv::attributeName< ImageSamplerUseInfo >(), mlir::spirv::attributeName< ImageSamplingInfo >(), mlir::spirv::attributeName< MemoryModel >(), mlir::spirv::attributeName< PackedVectorFormat >(), mlir::NVVM::PtxBuilder::build(), mlir::spirv::VerCapExtAttr::cap_iterator::cap_iterator(), checkGlobalXtorData(), concatArrayAttr(), mlir::LLVM::ModuleTranslation::convertArgAndResultAttrs(), mlir::LLVM::convertArrayToIndices(), mlir::convertFuncOpToLLVMFuncOp(), convertLinkerOptionsOp(), convertMLIRAttributesToLLVM(), convertModuleFlagsOp(), convertModuleFlagValue(), convertOperandBundles(), mlir::irdl::createVerifier(), deepCloneAliasScopes(), mlir::function_interface_impl::eraseFunctionArguments(), mlir::function_interface_impl::eraseFunctionResults(), mlir::spirv::VerCapExtAttr::ext_iterator::ext_iterator(), extractVector(), extractVector(), findSegment(), foldExtractStridedOpFromInsertChain(), mlir::spirv::VerCapExtAttr::get(), mlir::LLVM::detail::LoopAnnotationTranslation::getAccessGroups(), mlir::Builder::getAffineMapArrayAttr(), mlir::linalg::getAffineResultPositions(), mlir::function_interface_impl::getArgAttrDict(), getArgResAttrs(), mlir::Builder::getArrayAttr(), mlir::getAsOpFoldResult(), mlir::Builder::getBoolArrayAttr(), mlir::spirv::TargetEnvAttr::getCapabilitiesAttr(), mlir::spirv::VerCapExtAttr::getCapabilitiesAttr(), mlir::spirv::getDefaultResourceLimits(), getDimMap(), mlir::spirv::TargetEnvAttr::getExtensionsAttr(), mlir::spirv::VerCapExtAttr::getExtensionsAttr(), mlir::Builder::getF32ArrayAttr(), mlir::Builder::getF64ArrayAttr(), getFirstIntValue(), getFirstOrLastMappedMemberPtr(), mlir::Builder::getI32ArrayAttr(), mlir::Builder::getI64ArrayAttr(), mlir::getI64SubArray(), mlir::Builder::getIndexArrayAttr(), getLLVMAlignParamForCompressExpand(), mlir::getReassociationIndicesAttribute(), getReductionIndex(), mlir::function_interface_impl::getResultAttrDict(), mlir::Builder::getStrArrayAttr(), mlir::spirv::getStrArrayAttrForEnumList(), mlir::Builder::getTypeArrayAttr(), mlir::vector::getVectorSubscriptAttr(), handleArgumentImpl(), handleResultImpl(), inferStridedSliceOpResultType(), mlir::function_interface_impl::insertFunctionArguments(), mlir::function_interface_impl::insertFunctionResults(), inverseTransposeInBoundsAttr(), isArrayOf(), mlir::isBatchMatvec(), mlir::isBatchVecmat(), mlir::sparse_tensor::isCertainZero(), mlir::isColumnMajorMatmul(), mlir::linalg::BatchMatmulTransposeAOp::isDefaultIndexingMaps(), mlir::linalg::BatchMatmulTransposeBOp::isDefaultIndexingMaps(), mlir::linalg::MatmulTransposeAOp::isDefaultIndexingMaps(), mlir::linalg::MatmulTransposeBOp::isDefaultIndexingMaps(), isIntegerArrayAttrConfinedToRange(), isIntegerArrayAttrConfinedToShape(), isIntegerArrayAttrSmallerThanShape(), mlir::isMatvec(), mlir::isRowMajorBatchMatmul(), mlir::isRowMajorMatmul(), isSumOfIntegerArrayAttrConfinedToShape(), IsValidMatrixOpParams(), mlir::isVecmat(), makeArrayAttr(), makeI64ArrayAttr(), mlir::GPUFuncOpLowering::matchAndRewrite(), mlirLinalgGetIndexingMapsAttribute(), mlir::x86vector::avx2::inline_asm::mm256BlendPsAsm(), parseAlignedClause(), parseAttributesOp(), parseClauseWithRegionArgs(), parseCopyprivate(), parseCreateOperationOpAttributes(), parseDependVarList(), parseForeachMatchSymbols(), mlir::linalg::parseIndexingMapsAttr(), parseInReductionPrivateReductionRegion(), parseInReductionPrivateRegion(), parseNamedValueList(), parseNamedValueListImpl(), parseNamedValueListWithVariadicity(), parseOpBundles(), parseOperationOpAttributes(), mlir::AsmParser::parseOptionalAttribute(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalAttribute(), mlir::detail::Parser::parseOptionalAttribute(), parsePrivateReductionRegion(), parsePrivateRegion(), parseTargetOpRegion(), parseTaskReductionRegion(), parseVerCapExtAttr(), populateFromInt64AttrArray(), populateFromInt64AttrArray(), mlir::vector::populateVectorToVectorCanonicalizationPatterns(), printAttributesOp(), printAttributions(), printClauseWithRegionArgs(), printCreateOperationOpAttributes(), printForeachMatchSymbols(), printFunctionResultList(), mlir::call_interface_impl::printFunctionSignature(), printInReductionPrivateReductionRegion(), printInReductionPrivateRegion(), printMembersIndex(), printNamedValueList(), printNamedValueListImpl(), printNamedValueListWithVariadicity(), printOperationOpAttributes(), printPrivateReductionRegion(), printPrivateRegion(), printTargetOpRegion(), printTaskReductionRegion(), processParallelLoop(), processPassthroughAttrs(), processTargetSpecificAttrs(), mlir::spirv::Deserializer::processTensorARMType(), mlir::LLVM::ModuleTranslation::setAliasScopeMetadata(), setArgResAttrDict(), setArgResAttrs(), setAttributionAttrs(), mlir::function_interface_impl::setFunctionType(), mlir::LLVM::ModuleTranslation::setTBAAMetadata(), mlir::LLVM::detail::verifyAccessGroupOpInterface(), mlir::LLVM::detail::verifyAliasAnalysisOpInterface(), verifyDeviceTypeAndSegmentCountMatch(), verifyDeviceTypeCountMatch(), mlir::spirv::VerCapExtAttr::verifyInvariants(), verifyNames(), verifyStructArrayConstant(), mlir::function_interface_impl::verifyTrait(), and verifyTransferOp().

cloneAndFuseFirstUse()

Operation * cloneAndFuseFirstUse ( RewriterBase & rewriter, Diagnostic & diag, Operation * producerOp, Operation * containingOp )

static

Definition at line 1181 of file LinalgTransformOps.cpp.

References mlir::OpBuilder::clone(), diag(), mlir::IROperand< DerivedT, IRValueT >::get(), mlir::Operation::getLoc(), mlir::Operation::getOpResults(), mlir::detail::IROperandBase::getOwner(), mlir::Operation::isProperAncestor(), mlir::RewriterBase::modifyOpInPlace(), result, and mlir::OpBuilder::setInsertionPoint().

denormalizeIndVar()

SmallVector< Value > denormalizeIndVar ( RewriterBase & rewriter, Location loc, ValueRange ivs, ArrayRef< OpFoldResult > lbs, ArrayRef< OpFoldResult > steps )

static

When a loop is normalized, the uses of the induction variable within the loop need to replaced with original_lb + old_iv * original_step.

Definition at line 3787 of file LinalgTransformOps.cpp.

References mlir::bindDims(), mlir::bindSymbols(), mlir::Builder::getContext(), mlir::getValueOrCreateConstantIndexOp(), and mlir::affine::makeComposedFoldedAffineApply().

Referenced by normalizeForallLoopOp().

doit()

template<typename OpTy>

DiagnosedSilenceableFailure doit ( RewriterBase & rewriter, OpTy target, transform::ApplyToEachResultList & results, transform::TransformState & state )

static

Definition at line 4329 of file LinalgTransformOps.cpp.

References mlir::Value::getDefiningOp(), mlir::transform::ApplyToEachResultList::push_back(), mlir::RewriterBase::replaceOpWithNewOp(), mlir::OpBuilder::setInsertionPoint(), mlir::DiagnosedSilenceableFailure::success(), and target.

getI64ArrayAttr()

b getI64ArrayAttr

(

paddingDimensions

)

References b.

mayBeRead() [1/2]

bool mayBeRead ( OpOperand & operand )

static

Return true if the operand may be read from by its owner.

This is currently very conservative and only looks inside linalg operations to prevent unintentional data loss.

Definition at line 393 of file LinalgTransformOps.cpp.

References mlir::detail::IROperandBase::getOwner(), and mlir::Value::use_empty().

Referenced by mayBeRead().

mayBeRead() [2/2]

bool mayBeRead ( Value value )

static

Return true if the value may be read through any of its uses.

Definition at line 406 of file LinalgTransformOps.cpp.

References mlir::Value::getType(), mlir::Value::getUses(), and mayBeRead().

normalizeForallLoopOp()

scf::ForallOp normalizeForallLoopOp ( RewriterBase & rewriter, scf::ForallOp loop )

static

Given a scf.forall loop return a loop op with the loop bounds normalized. TODO: Replace this with a general utility to normalize scf.forall. At the time of writing, this wasnt done since adding this to scf dialect would disallow using of affine.apply operations due to cyclic dependencies. To avoid churn in lit tests with the change this was added with, defer that to a follow up.

Definition at line 3814 of file LinalgTransformOps.cpp.

References denormalizeIndVar(), mlir::Builder::getIndexAttr(), mlir::isOneInteger(), mlir::isZeroInteger(), mlir::RewriterBase::mergeBlocks(), normalizeUpperBounds(), mlir::RewriterBase::replaceOp(), and mlir::OpBuilder::setInsertionPointToStart().

Referenced by mlir::transform::tileToForallOpImpl().

normalizeUpperBounds()

SmallVector< OpFoldResult > normalizeUpperBounds ( RewriterBase & rewriter, Location loc, ArrayRef< OpFoldResult > lbs, ArrayRef< OpFoldResult > ubs, ArrayRef< OpFoldResult > steps )

static

Given lbs, ubs and steps of loops, return (for each loop), the normalized upper bound.

Definition at line 3770 of file LinalgTransformOps.cpp.

References mlir::bindSymbols(), mlir::Builder::getContext(), and mlir::affine::makeComposedFoldedAffineApply().

Referenced by normalizeForallLoopOp().

parseContinuousTileSizeTypes()

ParseResult parseContinuousTileSizeTypes ( OpAsmParser & parser, Type & targetType, Type & tileSizesType, Type & chunkSizesType )

static

Definition at line 3412 of file LinalgTransformOps.cpp.

References mlir::AsmParser::emitError(), mlir::AsmParser::getCurrentLocation(), mlir::AsmParser::parseType(), and success().

parseMultitileSizesTypes()

ParseResult parseMultitileSizesTypes ( OpAsmParser & parser, Type & targetType, Type & lowSizeType, Type & highSizeType, Type & splitPointType )

static

Definition at line 1690 of file LinalgTransformOps.cpp.

References mlir::AsmParser::emitError(), mlir::AsmParser::getCurrentLocation(), mlir::AsmParser::parseType(), and success().

printContinuousTileSizeTypes()

void printContinuousTileSizeTypes ( OpAsmPrinter & printer, Operation * op, Type targetType, Type tileSizes, Type  )

static

Definition at line 3406 of file LinalgTransformOps.cpp.

References mlir::OpAsmPrinter::printFunctionalType().

printMultitileSizesTypes()

void printMultitileSizesTypes ( OpAsmPrinter & printer, Operation * op, Type targetType, Type lowSizeType, Type , Type  )

static

Definition at line 1684 of file LinalgTransformOps.cpp.

References mlir::OpAsmPrinter::printFunctionalType().

reifyMixedParamAndHandleResults()

DiagnosedSilenceableFailure reifyMixedParamAndHandleResults ( TransformState & state, TransformOpInterface & transformOp, ArrayRef< OpFoldResult > mixedResults, SmallVectorImpl< int64_t > & reified )

static

When possible, converts each OpFoldResult in mixedResult to an integer if the value can be statically inferred.

If a result is a Value then it must be either a ParamType or a handle to an a constant like op.

Definition at line 172 of file LinalgTransformOps.cpp.

References mlir::Operation::getNumResults(), mlir::transform::TransformState::getParams(), mlir::transform::TransformState::getPayloadOps(), mlir::Operation::getResult(), mlir::getType(), mlir::Value::getType(), mlir::Type::isIndex(), mlir::m_Constant(), mlir::matchPattern(), and mlir::DiagnosedSilenceableFailure::success().

replaceForAllWithNewSignature()

Operation * replaceForAllWithNewSignature ( RewriterBase & rewriter, Diagnostic & diag, Operation * producerOp, Operation * containingOp, TilingResult & tileAndFuseResult, int64_t resultNumber, SmallVector< OpFoldResult > & offsets, SmallVector< OpFoldResult > & sizes )

static

Add new operands to the forall op for users of the producerOp that are dominated by the containing scf.forall op.

Definition at line 824 of file LinalgTransformOps.cpp.

References diag(), mlir::DominanceInfo::dominates(), mlir::RewriterBase::eraseBlock(), mlir::Builder::getIndexAttr(), mlir::Operation::getLoc(), mlir::Operation::getResult(), mlir::Value::getUsers(), mlir::Operation::isAncestor(), mlir::RewriterBase::replaceAllUsesWith(), mlir::RewriterBase::replaceUsesWithIf(), result, mlir::OpBuilder::setInsertionPoint(), and mlir::TilingResult::tiledValues.

Referenced by tileAndFuseFirstExtractUse().

sameOrEquivalentIterArg()

bool sameOrEquivalentIterArg ( Value src, Value dst )

static

Given two operands coming from a loop iter arg, 'src' and 'dst', return true if the operand 'src' is equal to 'dst' or equal to a iter arg present in a outer loop.

To determine the second condition, this function iterates using a worklist over the enclosing loops, trying to find 'src' in any of the parent loop's iter args.

Definition at line 906 of file LinalgTransformOps.cpp.

References mlir::OpOperand::getOperandNumber(), and mlir::Block::getParentOp().

Referenced by tileAndFuseFirstExtractUse().

tileAndFuseFirstExtractUse()

std::tuple< SmallVector< Operation * >, Operation * > tileAndFuseFirstExtractUse ( RewriterBase & rewriter, Diagnostic & diag, Operation * producerOp, Operation * containingOp )

static

Find the first "extract" user of producerOp and tile it right before its use.

The tiled op is fused under the containingOp. Return this fused op on success or nullptr if anything fails. If tiled op has uses that are dominated by containingOp, return a new containingOp with results of the fused op appended to results of the containingOp or nullptr if there are no dominated uses.

Definition at line 956 of file LinalgTransformOps.cpp.

References mlir::clone(), mlir::OpBuilder::clone(), diag(), mlir::RewriterBase::eraseOp(), mlir::Operation::getLoc(), mlir::OpOperand::getOperandNumber(), mlir::RewriterBase::modifyOpInPlace(), replaceForAllWithNewSignature(), mlir::RewriterBase::replaceOp(), sameOrEquivalentIterArg(), and mlir::OpBuilder::setInsertionPoint().

tileAndFuseFirstExtractUseThroughContainingOpBlockArgument()

SmallVector< Operation * > tileAndFuseFirstExtractUseThroughContainingOpBlockArgument ( RewriterBase & rewriter, Diagnostic & diag, Operation * producerOp, Operation * containingOp )

static

First, find the first "scf::ForallOp" user of producerOp and ensure it is exactly the containingOp, otherwise bail.

Then, find the first "extract" user of the tied block argument and tile it right before its "extract" use. The tiled op is fused under the containingOp. Return this fused op on success or nullptr if anything fails.

Definition at line 1079 of file LinalgTransformOps.cpp.

References mlir::OpBuilder::clone(), diag(), mlir::RewriterBase::eraseOp(), mlir::IROperand< DerivedT, IRValueT >::get(), mlir::Operation::getLoc(), mlir::OpOperand::getOperandNumber(), mlir::tensor::getOrCreateDestinations(), mlir::Value::getUsers(), mlir::IRMapping::map(), mlir::RewriterBase::modifyOpInPlace(), mlir::RewriterBase::replaceOp(), mlir::OpBuilder::setInsertionPoint(), and mlir::Operation::setOperand().

tryApply()

template<typename PatternTy, typename... Args>

FailureOr< LinalgOp > tryApply ( Operation * operation, Args &&... args )

static

Attempts to apply the pattern specified as template argument to the given operation.

The pattern is expected to have a returningMatchAndRewrite function that returns the "main" result or failure. Returns failure if the pattern failed to apply. Extra arguments are forwarded to the pattern constructor.

Definition at line 63 of file LinalgTransformOps.cpp.

References mlir::Operation::getContext(), result, and mlir::OpBuilder::setInsertionPoint().

unpackSingleIndexResultPayloadOperations() [1/2]

DiagnosedSilenceableFailure unpackSingleIndexResultPayloadOperations ( transform::TransformState & state, TransformOpInterface transformOp, SmallVector< OpFoldResult > & result, ArrayRef< OpFoldResult > ofrs )

static

Assuming that ofr is an index attr or a param of index type or a transform dialect handle mapped to exactly one op with one index result, return that value.

Definition at line 86 of file LinalgTransformOps.cpp.

References diag(), mlir::Operation::getLoc(), mlir::Value::getLoc(), mlir::Operation::getNumResults(), mlir::transform::TransformState::getParams(), mlir::transform::TransformState::getPayloadOps(), mlir::Operation::getResult(), mlir::Value::getType(), mlir::Type::isIndex(), result, and mlir::DiagnosedSilenceableFailure::success().

unpackSingleIndexResultPayloadOperations() [2/2]

DiagnosedSilenceableFailure unpackSingleIndexResultPayloadOperations ( transform::TransformState & state, TransformOpInterface transformOp, SmallVector< OpFoldResult > & result, Value packedHandle )

static

Definition at line 138 of file LinalgTransformOps.cpp.

References diag(), mlir::transform::TransformState::getParams(), mlir::transform::TransformState::getPayloadOps(), mlir::Value::getType(), result, and mlir::DiagnosedSilenceableFailure::success().

Variable Documentation

b

b

Return true if permutation is a valid permutation of the outer_dims_perm (case OuterOrInnerPerm::Outer) or inner_dims_pos (OuterOrInnerPerm::Inner) of the tensor.pack or tensor.unpack op.

/ This is the case when the permutation rank matches the rank expected by / op and permutation is itself a permutation vector. / Return true if either op or permutation` are empty to allow a simpler / polymorphic implementation. template <typename RelayoutOpTy> static bool isValidPackingPermutation( RelayoutOpTy op, ArrayRef<int64_t> permutation, OuterOrInnerPerm outerOrInnerPerm = OuterOrInnerPerm::Outer) { static_assert( llvm::is_one_of<RelayoutOpTy, linalg::PackOp, linalg::UnPackOp>::value, "applies to only pack or unpack operations"); if (!op || permutation.empty()) return true; size_t innerRank = op.getInnerDimsPos().size(); if (outerOrInnerPerm == OuterOrInnerPerm::Inner) return permutation.size() == innerRank && isPermutationVector(permutation); op.getOuterDimsPerm() may be empty, in which case it is identity. Don't rely on it. if (std::is_same<RelayoutOpTy, linalg::PackOp>::value) { return permutation.size() == op.getSourceRank() && isPermutationVector(permutation); } return permutation.size() == op.getDestRank() && isPermutationVector(permutation); }

DiagnosedSilenceableFailure transform::PackTransposeOp::apply(transform::TransformRewriter &rewriter, transform::TransformResults &transformResults, transform::TransformState &state) { auto packOrUnpackOps = state.getPayloadOps(getTargetPackOrUnPackOp()); auto linalgOps = state.getPayloadOps(getTargetLinalgOp()); Step 1. If nothing to pack, propagate success. if (std::empty(packOrUnpackOps)) { transformResults.set(cast<OpResult>(getPackedOp()), {}); transformResults.set(cast<OpResult>(getPackOp()), {}); transformResults.set(cast<OpResult>(getUnPackOp()), {}); return DiagnosedSilenceableFailure::success(); }

Step 2. Bunch of runtime sanity check and error messages. Step 2.1. Fail on multi-op handles. if (!llvm::hasSingleElement(packOrUnpackOps) || !llvm::hasSingleElement(linalgOps)) { return emitSilenceableError() << "requires target to map to exactly 1 " "packing op and 1 packed op (" << "got " << llvm::range_size(packOrUnpackOps) << " and " << llvm::range_size(linalgOps) << ")"; }

Step 2.2. Fail on wrong type. auto packOp = dyn_cast<linalg::PackOp>(*packOrUnpackOps.begin()); auto unPackOp = dyn_cast<linalg::UnPackOp>(*packOrUnpackOps.begin()); if ((!packOp && !unPackOp)) { return emitSilenceableError() << "requires target to map to a " "linalg.pack or linalg.unpack"; } LinalgOp linalgOpTarget = dyn_cast<LinalgOp>(*linalgOps.begin()); if (!linalgOpTarget) return emitSilenceableError() << "requires a LinalgOp target";

Step 2.3. Fail if we can't get the producer / consumer Linalg op. LinalgOp linalgOp; if (packOp && packOp.getResult().hasOneUse()) linalgOp = dyn_cast<LinalgOp>(*(packOp.getResult().getUsers().begin())); else if (unPackOp) linalgOp = unPackOp.getSource().getDefiningOp<LinalgOp>(); if (linalgOp != linalgOpTarget) { auto errorMsg = packOp ? StringLiteral{"not a single use by the LinalgOp target"} : StringLiteral{"not produced by the LinalgOp target"}; return emitSilenceableError() << errorMsg; }

Step 2.4. If we have an UnPackOp, we need to fetch the symmetrical PackOp. if (unPackOp) { assert(!packOp && "packOp must be null on entry when unPackOp is not null"); OpOperand *packUse = linalgOp.getDpsInitOperand( cast<OpResult>(unPackOp.getSource()).getResultNumber()); packOp = packUse->get().getDefiningOp<linalg::PackOp>(); if (!packOp || !packOp.getResult().hasOneUse()) return emitSilenceableError() << "could not find matching pack op"; }

Step 2.5. Fail if any permutation does not validate. for (auto permType : {OuterOrInnerPerm::Outer, OuterOrInnerPerm::Inner}) { ArrayRef<int64_t> perm = (permType == OuterOrInnerPerm::Outer) ? getOuterPerm() : getInnerPerm(); auto errorMsg = (permType == OuterOrInnerPerm::Outer) ? StringLiteral{"invalid outer_perm"} : StringLiteral{"invalid inner_perm"}; if (!isValidPackingPermutation(packOp, perm, permType) || !isValidPackingPermutation(unPackOp, perm, permType)) { Operation *packOrUnpackOp = unPackOp ? unPackOp.getOperation() : packOp.getOperation(); return emitSilenceableError() << errorMsg << ": " << *packOrUnpackOp; } }

From here on, packOp and linalgOp are always present, unPackOp may or may not be present. assert(packOp && linalgOp && "unexpected null op");

Step 3. Actually transpose the ops. FailureOr<PackTransposeResult> res = packTranspose( rewriter, packOp, linalgOp, unPackOp, getOuterPerm(), getInnerPerm()); Preconditions have been checked, it is an error to fail here. assert(succeeded(res) && "unexpected packTranspose failure");

Step 4. Return results. transformResults.set(cast<OpResult>(getPackOp()), {res->transposedPackOp}); transformResults.set(cast<OpResult>(getPackedOp()), {res->transposedLinalgOp}); if (unPackOp) { transformResults.set(cast<OpResult>(getUnPackOp()), {res->transposedUnPackOp}); } else { transformResults.set(cast<OpResult>(getUnPackOp()), {}); }

return DiagnosedSilenceableFailure::success(); }

===------------------------------------------------------------------—===// PadOp ===------------------------------------------------------------------—===//

void transform::PadOp::build(OpBuilder &b, OperationState &result, Value target, ArrayRef<int64_t> paddingDimensions, ArrayRef<int64_t> padToMultipleOf, ArrayRef<int64_t> nofoldFlags, ArrayRef<Attribute> transposePaddings, StringRef copyBackOp, bool usePrescribedTensorShapes) { auto resultType = transform::AnyOpType::get(b.getContext()); return build(/*odsBuilder=

Definition at line 2096 of file LinalgTransformOps.cpp.

Referenced by addBodyWithPayloadOp(), mlir::FlatLinearValueConstraints::addBound(), addInitOperandsToLoopNest(), mlir::sparse_tensor::Merger::addLat(), mlir::scf::addLoopRangeConstraints(), addShardOp(), addShardOp(), mlir::xegpu::addWithRightAligned(), mlir::affine::affineForOpBodySkew(), mlir::arm_sme::allocateSMETiles(), allocBuffer(), annotateConflict(), annotateNonWritableTensor(), annotateOpsWithAliasSets(), mlir::OpToFuncCallLowering< SourceOp >::appendOrGetFuncOp(), areAllOpsInTheBlockListInvariant(), mlir::ValueBoundsConstraintSet::areOverlappingSlices(), areSameBitwidthScalarType(), areVarsAligned(), mlir::remark::detail::Remark::Arg::Arg(), mlir::spirv::attributeName< ImageArrayedInfo >(), mlir::spirv::attributeName< ImageSamplingInfo >(), augmentMapAndBounds(), mlir::affine::boundCheckLoadOrStoreOp(), broadcastIfNeeded(), mlir::tensor::bubbleUpPadSlice(), buildArithValue(), buildBatchMatmulOp(), buildBatchReduceMatmulOp(), buildDivOp(), buildIdentityRegion(), mlir::memref::buildIndependentOp(), mlir::tensor::buildIndependentOp(), mlir::tensor::buildIndependentOp(), buildLLVMFunctionType(), buildMatmulOp(), mlir::tosa::buildMatMulOpQuantizationAttr(), buildMatMulOpWithQuantInfo(), buildMax(), buildMin(), buildMultiDimReduce(), buildNumReadElements(), buildStructuredOp(), buildSubAndExpOp(), calculateTileOffsetsAndSizes(), mlir::bufferization::castOrReallocMemRefValue(), ValueEquivalenceCache::checkCommutativeEquivalent(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::classof(), mlir::nvgpu::impl::OptimizeSharedMemoryBase< DerivedT >::classof(), mlir::clone(), mlir::clone(), mlir::arith::impl::ArithExpandOpsPassBase< DerivedT >::clonePass(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::clonePass(), mlir::cloneWithoutRegions(), mlir::sparse_tensor::SparseTensorLevel::collapseRangeBetween(), composedAffineMultiply(), mlir::ValueBoundsConstraintSet::computeBound(), mlir::ValueBoundsConstraintSet::computeConstantDelta(), mlir::linalg::computeContinuousTileSizes(), mlir::linalg::computeMultiTileSizes(), mlir::affine::computeSliceUnion(), mlir::linalg::computeTileOffsets(), mlir::linalg::computeTileSizes(), mlir::linalg::concat(), mlir::sparse_tensor::constantI1(), mlir::constFoldBinaryOp(), mlir::constFoldBinaryOp(), constructParametricallyTiledIndexSetHyperRect(), constructTiledIndexSetHyperRect(), constructTiledLoopNest(), convertAcoshOp(), convertAsinhOp(), convertAtanhOp(), convertCeilOp(), convertExp2fOp(), convertFmaFOp(), convertFPowIOp(), convertOperationImpl(), convertPowfOp(), convertRoundEvenOp(), convertRoundOp(), convertScalarToComplexDtype(), mlir::convertScalarToDtype(), convertScalarToFpDtype(), convertScalarToIntDtype(), convertSinhOp(), convertTanOp(), mlir::linalg::copyToGPUPrivateMemory(), mlir::linalg::copyToWorkgroupMemory(), mlir::tensor::ExtractSliceFromCollapseHelper::create(), mlir::tensor::ExtractSliceFromCollapseHelper::create(), create4x128BitSuffle(), mlir::arith::createArithExpandOpsPass(), mlir::arith::createArithExpandOpsPass(), createAsyncDispatchFunction(), mlir::createAsyncParallelForPass(), mlir::OpBuilder::createBlock(), mlir::tensor::createCanonicalRankReducingExtractSliceOp(), mlir::tensor::createCanonicalRankReducingInsertSliceOp(), mlir::memref::createCanonicalRankReducingSubViewOp(), createChoosePivot(), createCompareThenSwap(), mlir::impl::ConvertLinalgToAffineLoopsPassBase< DerivedT >::createConvertLinalgToAffineLoopsPass, createCtlzFunc(), mlir::tensor::createDynamicDimValues(), createFloatConst(), mlir::linalg::createFoldedDimOp(), createFullTiles(), createInBoundsCond(), createLinalgBodyCalculationForElementwiseOp(), mlir::math::createMathExpandOpsPass(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::createMathExpandOpsPass, createMemcpy(), createNewDynamicSizes(), mlir::linalg::createOrFoldDimOp(), mlir::vector::createOrFoldDimOp(), mlir::createOutlineShapeComputationPass(), mlir::tensor::createPadHighOp(), createParallelComputeFunction(), createPrivateMemRef(), createSeparationCondition(), createSplitPart(), createTypeCanonicalizedMemRefOperands(), createUnpackHiPd(), createUnpackHiPs(), createUnpackLoPd(), createUnpackLoPs(), mlir::shard::detail::defaultAddShardingAnnotations(), defaultAllocBufferCallBack(), defaultDeallocBufferCallBack(), delinearize(), dependsOn(), mlir::sparse_tensor::SparseIterator::deref(), mlir::sparse_tensor::SparseIterator::derefImpl(), mlir::presburger::FracMatrix::determinant(), mlir::affine::AffineValueMap::difference(), doAsyncDispatch(), mlir::linalg::GenerateLoopNest< LoopTy >::doit(), mlir::linalg::GenerateLoopNest< LoopTy >::doit(), mlir::linalg::GenerateLoopNest< LoopTy >::doit(), mlir::DominanceInfo::dominates(), mlir::DominanceInfo::dominates(), mlir::DominanceInfo::dominates(), doSequentialDispatch(), mlir::presburger::dotProduct(), mlir::Region::dropAllReferences(), mlir::tensor::dropGivenUnitDims(), dropRefIfNoUses(), mlir::linalg::dropUnitDims(), mlir::sparse_tensor::Merger::dumpBits(), emitIsPositiveIndexAssertion(), emitScalarImplementation(), mlir::sparse_tensor::SparseIterationSpace::extractIterator(), fillStructuredOpRegion(), mlir::sparse_tensor::I64BitSet::find_next(), findAncestorIteratorInRegion(), mlir::affine::findInnermostCommonBlockInScope(), mlir::detail::DominanceInfoBase< IsPostDom >::findNearestCommonDominator(), mlir::presburger::IntegerRelation::findSymbolicIntegerLexMax(), mlir::foldAttributesIntoMap(), foldExtractAfterInsert(), foldExtractAfterInsertSlice(), foldExtractFromExtractStrided(), foldExtractFromShapeCast(), foldExtractOpFromExtractChain(), foldExtractStridedOpFromInsertChain(), foldExtractStridedOpFromInsertChain(), foldIdentityOffsetSizeAndStrideOpInterface(), foldInsertAfterExtractSlice(), foldInsertAfterInsertSlice(), foldTransferInBoundsAttribute(), mlir::sparse_tensor::Merger::foreachTensorLoopId(), mlir::sparse_tensor::SparseIterator::forward(), mlir::sparse_tensor::SparseIterator::forwardIf(), mlir::sparse_tensor::SparseIterator::forwardImpl(), mlir::affine::fullyComposeAndComputeConstantDelta(), fuse(), fuse(), fuseIfLegal(), mlir::affine::fuseLoops(), mlir::linalg::fuseProducerOfTensor(), mlir::linalg::fuseProducerOfTensor(), mlir::xegpu::genBinOp(), genBuffers(), generateCopy(), generateLoopNestingName(), mlir::linalg::generateParallelLoopNest(), generateShiftedLoop(), mlir::sparse_tensor::SparseIterator::genForCond(), genIf(), mlir::sparse_tensor::SparseIterator::genInit(), genLoopWithIterator(), mlir::sparse_tensor::SparseIterator::genNotEnd(), mlir::sparse_tensor::SparseIterator::genNotEndImpl(), genWhenInBound(), mlir::sparse_tensor::SparseIterator::genWhileCond(), getAllTidLvlsInLatPoints(), mlir::arith::impl::ArithExpandOpsPassBase< DerivedT >::getArgument(), mlir::impl::SCFForLoopSpecializationBase< DerivedT >::getArgument(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::getArgument(), mlir::impl::SCFForLoopSpecializationBase< DerivedT >::getArgumentName(), mlir::getBlocksSortedByDominance(), getBoundedTileSize(), getCleanupLoopLowerBound(), mlir::tensor::getCollapsedExtractSliceInfo(), getCommonConstraints(), mlir::linalg::getConvolvedExpr(), mlir::sparse_tensor::SparseIterator::getCursorValTypes(), mlir::spirv::getDefaultResourceLimits(), mlir::impl::ConvertLinalgToLoopsPassBase< DerivedT >::getDependentDialects(), mlir::impl::OutlineShapeComputationPassBase< DerivedT >::getDependentDialects(), mlir::impl::SCFForLoopPeelingBase< DerivedT >::getDependentDialects(), mlir::nvgpu::impl::OptimizeSharedMemoryBase< DerivedT >::getDependentDialects(), mlir::arith::impl::ArithExpandOpsPassBase< DerivedT >::getDescription(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::getDescription(), mlir::tensor::getExpandedExtractSliceInfo(), getFirstOrLastMappedMemberPtr(), getFullRankPaddingSizes(), getGenericOpLoopRange(), getI64ArrayAttr(), getIndicesForAccess(), getInputAndOutputIndices(), getInstAtPosition(), mlir::sparse_tensor::Merger::getLoopDependentLevel(), mlir::sparse_tensor::Merger::getLoopDependentLevelType(), mlir::FlatLinearConstraints::getLowerAndUpperBound(), mlir::sparse_tensor::Merger::getLvl(), mlir::sparse_tensor::Merger::getLvlType(), getMaskFormat(), mlir::shard::getMixedAsValues(), mlir::getMixedValues(), mlir::impl::SCFForLoopPeelingBase< DerivedT >::getName(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::getName(), mlir::arith::getNeutralElement(), mlir::affine::getNumCommonSurroundingLoops(), mlir::tensor::getOrCreateDestination(), mlir::tensor::getOrCreateDestinations(), getOrCreateOperandsMatchingResultTypes(), mlir::getOrCreateRanges(), mlir::getOrCreateStringConstant(), mlir::getOrDefineFunction(), mlir::arith::impl::ArithExpandOpsPassBase< DerivedT >::getPassName(), mlir::math::impl::MathExtendToSupportedTypesBase< DerivedT >::getPassName(), getRangeFromOperandShape(), mlir::getReassociationIndicesAttribute(), getReturnOps(), mlir::bufferization::getReturnOps(), getSetDifference(), getSharedAddressSpaceAttribute(), mlir::getSimplifiedOfrAndStaticSizePair(), getSlice(), mlir::xegpu::uArch::SubgroupMatrixMultiplyAcc::getSupportedShapes(), mlir::linalg::getUnPackInverseSrcPerm(), mlir::getValueOrCreateCastToIndexLike(), mlir::getValueOrCreateConstantIndexOp(), mlir::getValueOrCreateConstantIndexOp(), mlir::getValueOrCreateConstantIntOp(), getZero(), happensBefore(), happensBefore(), mlir::sparse_tensor::Merger::hasAnySparse(), hasSameBitwidth(), mlir::sparse_tensor::Merger::hasSparseIdxReduction(), hoistAffineIfOp(), mlir::linalg::hoistRedundantVectorTransfers(), hoistSubsetAtIterArg(), if(), mlir::intrange::inferAdd(), mlir::intrange::inferAnd(), inferContractionDimsImpl(), inferDivSRange(), inferDivURange(), mlir::inferExpandShapeOutputShape(), mlir::intrange::inferMul(), mlir::intrange::inferOr(), mlir::intrange::inferSub(), inlinePayload(), inlineRegionAndEmitStore(), mlir::affine::insertBackwardComputationSlice(), insertCopies(), insertCopyLoops(), mlir::scf::insideMutuallyExclusiveBranches(), invertCollapseShapeIndexing(), invertSliceIndexing(), mlir::isBatchMatvec(), mlir::isBatchVecmat(), isDivZeroOrOverflow(), isFunctionArgument(), mlir::sparse_tensor::Merger::isLvlWithNonTrivialIdxExp(), mlir::sparse_tensor::Merger::isOutTensor(), mlir::isRowMajorBatchMatmul(), TosaProfileCompliance::isSameTypeInfo(), mlir::memref::isSameViewOrTrivialAlias(), mlir::sparse_tensor::Merger::isSparseLvlWithNonTrivialIdxExp(), isTransposeMatrixLoadMap(), mlir::sparse_tensor::Merger::latGT(), mlir::sparse_tensor::SparseIterator::locate(), mlir::sparse_tensor::SparseIterator::locateImpl(), mlir::LLVM::lookupOrCreateAlignedAllocFn(), mlir::LLVM::lookupOrCreateFn(), mlir::LLVM::lookupOrCreateFreeFn(), mlir::LLVM::lookupOrCreateGenericAlignedAllocFn(), mlir::LLVM::lookupOrCreateGenericAllocFn(), mlir::LLVM::lookupOrCreateGenericFreeFn(), mlir::LLVM::lookupOrCreateMallocFn(), mlir::LLVM::lookupOrCreateMemRefCopyFn(), mlir::LLVM::lookupOrCreatePrintBF16Fn(), mlir::LLVM::lookupOrCreatePrintCloseFn(), mlir::LLVM::lookupOrCreatePrintCommaFn(), mlir::LLVM::lookupOrCreatePrintF16Fn(), mlir::LLVM::lookupOrCreatePrintF32Fn(), mlir::LLVM::lookupOrCreatePrintF64Fn(), mlir::LLVM::lookupOrCreatePrintI64Fn(), mlir::LLVM::lookupOrCreatePrintNewlineFn(), mlir::LLVM::lookupOrCreatePrintOpenFn(), mlir::LLVM::lookupOrCreatePrintStringFn(), mlir::LLVM::lookupOrCreatePrintU64Fn(), lookupOrCreateReservedFn(), lookupOrCreateSPIRVFn(), mlir::sparse_tensor::Merger::loop(), mlir::affine::loopUnrollByFactor(), mlir::loopUnrollByFactor(), mlir::affine::loopUnrollJamByFactor(), mlir::loopUnrollJamByFactor(), mlir::sparse_tensor::CodegenEnv::lt(), mlir::vector::makeArithReduction(), makeCanonicalAffineApplies(), mlir::affine::makeComposedAffineApply(), mlir::affine::makeComposedAffineApply(), mlir::affine::makeComposedAffineMin(), mlir::affine::makeComposedFoldedAffineApply(), mlir::affine::makeComposedFoldedAffineApply(), mlir::affine::makeComposedFoldedAffineMax(), mlir::affine::makeComposedFoldedAffineMin(), makeComposedFoldedMinMax(), mlir::affine::makeComposedFoldedMultiResultAffineApply(), makeComposedMinMax(), mlir::linalg::makeComposedPadHighOp(), makeIndependent(), mlir::linalg::makeMemRefCopyOp(), mlir::sparse_tensor::makeNonEmptySubSectIterator(), mlir::presburger::Simplex::makeProduct(), mlir::makeRegionIsolatedFromAbove(), mlir::sparse_tensor::makeSimpleIterator(), mlir::sparse_tensor::makeSparseTensorLevel(), mlir::sparse_tensor::makeSparseTensorLevel(), mlir::linalg::makeTiledLoopRanges(), mlir::sparse_tensor::makeTraverseSubSectIterator(), mlir::gpu::mapParallelOp(), CanonicalizeContractAdd< AddOpType >::matchAndRewrite(), DropUnitDimsFromScfForOp::matchAndRewrite(), IAddCarryFold::matchAndRewrite(), MulExtendedFold< spirv::SMulExtendedOp, true >::matchAndRewrite(), OuterProductOpLowering::matchAndRewrite(), mlir::linalg::RegionMatcher::matchAsScalarBinaryOp(), matchFullMask(), mlir::affine::materializeComputedBound(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::MathExpandOpsPassBase(), mayAlias(), mergeAndAlignVars(), mfmaOpToIntrinsic(), mfmaOpToScaledIntrinsic(), mlirBlockCreate(), mlirBlockDetach(), mlirBlockGetNumPredecessors(), mlirBlockGetPredecessor(), mlirLinalgFillBuiltinNamedOpRegion(), mlir::x86vector::avx2::intrin::mm256BlendPs(), mlir::x86vector::avx2::inline_asm::mm256BlendPsAsm(), mlir::x86vector::avx2::intrin::mm256Permute2f128Ps(), mlir::x86vector::avx2::intrin::mm256ShufflePs(), mlir::x86vector::avx2::intrin::mm256UnpackHiPs(), mlir::x86vector::avx2::intrin::mm256UnpackLoPs(), mlir::presburger::multiplyPolynomials(), mlir::scf::naivelyFuseParallelOps(), notifyBlockInsertions(), offsetFromMinCrd(), mlir::linalg::offsetIndices(), mlir::linalg::offsetIndices(), mlir::bufferization::OneShotAnalysisState::OneShotAnalysisState(), mlir::detail::DominanceInfoBase< IsPostDom >::operator=(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::operator=(), mlir::outlineIfOp(), padThroughLoopIterArg(), mlir::parallelLoopUnrollByFactors(), mlir::arith::parseFloatType(), mlir::shard::partitionFuncOp(), mlir::RewritePattern::Pattern(), mlir::sparse_tensor::SparseTensorLevel::peekCrdAt(), mlir::sparse_tensor::SparseTensorLevel::peekRangeAt(), peelForLoop(), mlir::bufferization::populateDynamicDimSizes(), mlir::math::populateExtendToSupportedTypesTypeConverter(), mlir::vector::populateVectorToVectorCanonicalizationPatterns(), mlir::PostDominanceInfo::postDominates(), mlir::PostDominanceInfo::postDominates(), printArgs(), printTransferAttrs(), promoteSingleIterReductionLoop(), mlir::linalg::promoteSubviewAsNewBuffer(), mlir::linalg::promoteSubViews(), promoteSubViews(), promoteSubViews(), mlir::DominanceInfo::properlyDominates(), mlir::DominanceInfo::properlyDominates(), mlir::DominanceInfo::properlyDominates(), mlir::PostDominanceInfo::properlyPostDominates(), mlir::PostDominanceInfo::properlyPostDominates(), reduce(), reduceIfNeeded(), mlir::nvgpu::registerOptimizeSharedMemory(), mlir::nvgpu::registerOptimizeSharedMemoryPass(), mlir::affine::reifyIndexValueBound(), mlir::arith::reifyIndexValueBound(), reifyOrComputeDynamicSizes(), mlir::reifyResultShapes(), ReifyCollapseShapeOp::reifyResultShapes(), mlir::affine::reifyShapedValueDimBound(), mlir::arith::reifyShapedValueDimBound(), mlir::affine::reifyValueBound(), mlir::arith::reifyValueBound(), rewriteOneForallCommonImpl(), mlir::detail::sameOffsetsSizesAndStrides(), scaleAndAddForAssert(), mlir::impl::SCFForLoopPeelingBase< DerivedT >::SCFForLoopPeelingBase(), mlir::impl::SCFForLoopSpecializationBase< DerivedT >::SCFForLoopSpecializationBase(), mlir::impl::SCFForLoopSpecializationBase< DerivedT >::SCFForLoopSpecializationBase(), selectShardingOption(), mlir::sparse_tensor::ir_detail::DimSpec::setElideExpr(), mlir::sparse_tensor::ir_detail::LvlSpec::setElideVar(), setInsertionPointAfter(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), mlir::linalg::PadTilingInterfaceOptions::setPadToMultipleOf(), mlir::linalg::LinalgTilingOptions::setTileSizes(), signedCeilNonnegInputs(), mlir::sparse_tensor::Merger::simplifyCond(), mlir::sparse_tensor::SparseIterationSpace::SparseIterationSpace(), mlir::sparse_tensor::SparseIterationSpace::SparseIterationSpace(), mlir::sparse_tensor::SparseIterationSpace::SparseIterationSpace(), specializeForLoopForUnrolling(), specializeParallelLoopForUnrolling(), mlir::linalg::splitReduction(), mlir::linalg::splitReductionByScaling(), stableTopologicalSort(), stripmineSink(), stripmineSink(), mlir::sparse_tensor::Merger::tensor(), mlir::linalg::tileLinalgOp(), tileLinalgOpImpl(), tileLinalgOpImpl(), mlir::linalg::tileReductionUsingForall(), mlir::linalg::transformIndexOps(), transposeToShuffle16x16(), transposeToShuffle1D(), truncToI32(), tryGetBlocksInSameRegion(), unpackOperandVector(), mlir::linalg::updateBoundsForCyclicDistribution(), mlir::sparse_tensor::SparseIterator::upperBound(), mlir::ValueBoundsConstraintSet::Variable::Variable(), mlir::ValueBoundsConstraintSet::Variable::Variable(), mlir::SimpleAffineExprFlattener::visitMulExpr(), wrapInExecuteRegion(), mlir::math::impl::MathExpandOpsPassBase< DerivedT >::~MathExpandOpsPassBase(), and mlir::impl::SCFForLoopSpecializationBase< DerivedT >::~SCFForLoopSpecializationBase().

result

result

Definition at line 2097 of file LinalgTransformOps.cpp.

Referenced by mlir::call_interface_impl::addArgAndResultAttrs(), mlir::call_interface_impl::addArgAndResultAttrs(), addBodyWithPayloadOp(), mlir::tracing::TagBreakpointManager::addBreakpoint(), mlir::affine::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::bufferization::OneShotAnalysisState::addExtension(), mlir::transform::TransformState::addExtension(), mlir::presburger::Simplex::addInequality(), mlir::detail::AttrTypeReplacerBase< Concrete >::addReplacement(), addShardOp(), mlir::AsmParser::addTypesToList(), mlir::AsmParser::addTypeToList(), mlir::AsmParserState::addUses(), mlir::AliasAnalysis::alias(), mlir::LocalAliasAnalysis::alias(), mlir::LocalAliasAnalysis::aliasImpl(), mlir::transform::TransformEachOpTrait< OpTy >::apply(), mlir::applyPermutationMap(), applySequenceBlock(), mlir::presburger::LinearTransform::applyTo(), mlir::transform::TransformState::applyTransform(), mlir::linalg::areElementwiseOpsFusable(), mlir::ROCDL::SerializeGPUModuleBase::assembleIsa(), mlir::linalg::asShapeWithAnyValueAsDynamic(), mlir::python::PyMlirContext::attachDiagnosticHandler(), mlir::barePtrFuncArgTypeConverter(), binaryFolder(), bottomUpFromTerminatorsHeuristic(), bubbleDownCastsPassthroughOpImpl(), mlir::tensor::bubbleUpPadSlice(), mlir::bufferization::func_ext::CallOpInterface::bufferize(), mlir::linalg::bufferizeToAllocation(), mlir::linalg::bufferizeToAllocation(), mlir::affine::AffineDmaStartOp::build(), mlir::affine::AffineDmaWaitOp::build(), mlir::affine::NestedMatch::build(), mlir::arith::ConstantFloatOp::build(), mlir::arith::ConstantIndexOp::build(), mlir::arith::ConstantIntOp::build(), mlir::arith::ConstantIntOp::build(), mlir::arith::ConstantIntOp::build(), mlir::linalg::BatchMatmulTransposeAOp::build(), mlir::linalg::BatchMatmulTransposeAOp::build(), mlir::linalg::BatchMatmulTransposeAOp::build(), mlir::linalg::BatchMatmulTransposeBOp::build(), mlir::linalg::BatchMatmulTransposeBOp::build(), mlir::linalg::BatchMatmulTransposeBOp::build(), mlir::linalg::MatmulTransposeAOp::build(), mlir::linalg::MatmulTransposeAOp::build(), mlir::linalg::MatmulTransposeAOp::build(), mlir::linalg::MatmulTransposeBOp::build(), mlir::linalg::MatmulTransposeBOp::build(), mlir::linalg::MatmulTransposeBOp::build(), buildAttributeAPInt(), buildAvgPool2dOpWithQuantInfo(), buildConvOpWithQuantInfo(), buildDefaultRegistryFn(), mlir::tblgen::buildErrorStreamingString(), buildMatMulOpWithQuantInfo(), buildNegateOpWithQuantInfo(), buildPadOpWithQuantInfo(), buildSequentialConstant(), buildTransConvOpWithQuantInfo(), buildVariableOp(), cachedLookup(), calculateCycleEdges(), canonicalizeRegisterConstraints(), checkAssumptionForFusingConsumer(), checkAssumptionForLoop(), checkImplementationStatus(), mlir::affine::checkMemrefAccessDependence(), checkTransformationPreconditions(), cloneAndFuseFirstUse(), mlir::xegpu::impl::XeGPUBlockingBase< DerivedT >::clonePass(), mlir::coalescePerfectlyNestedSCFForLoops(), collectUnderlyingAddressValues(), mlir::spirv::combine(), mlir::presburger::PresburgerRelation::compose(), composeLegalityCallbacks(), computeBackwardSlice(), computeElementwiseMulImpl(), mlir::OperationEquivalence::computeHash(), mlir::presburger::IntegerRelation::computeReprWithOnlyDivLocals(), mlir::presburger::PresburgerRelation::computeReprWithOnlyDivLocals(), mlir::computeShapeRatio(), mlir::presburger::computeSingleVarRepr(), mlir::affine::computeSliceUnion(), computeStrides(), mlir::presburger::SymbolicLexSimplex::computeSymbolicIntegerLexMin(), mlir::presburger::PresburgerRelation::computeVolume(), concatArrayAttr(), mlir::detail::AffineBinaryOpExprStorage::construct(), mlir::detail::AffineConstantExprStorage::construct(), mlir::detail::AffineDimExprStorage::construct(), mlir::detail::FileLineColRangeAttrStorage::construct(), mlir::detail::TupleTypeStorage::construct(), mlir::query::matcher::RegistryManager::constructMatcher(), mlir::LLVMTypeConverter::convertCallingConventionType(), convertConstantOpMmaSync(), mlir::lsp::MLIRServer::convertFromBytecode(), mlir::convertFuncOpToLLVMFuncOp(), convertFuncOpTypes(), mlir::LLVMTypeConverter::convertFunctionSignature(), mlir::LLVMTypeConverter::convertFunctionSignature(), convertHostOrTargetOperation(), convertIntrinsicResult(), convertMixedValuesToInt(), convertMLIRAttributeToLLVM(), convertOmpTargetData(), convertOperationImpl(), convertRoundEvenOp(), convertRoundOp(), mlir::convertScalarToDtype(), convertTargetOpsInNest(), copyAPIntToArrayForBEmachine(), copyArrayRefInto(), copyArrayToAPIntForBEmachine(), mlir::StorageUniquer::StorageAllocator::copyInto(), mlir::StorageUniquer::StorageAllocator::copyInto(), mlir::affine::AffineDmaStartOp::create(), mlir::affine::AffineDmaWaitOp::create(), mlir::arith::ConstantFloatOp::create(), mlir::arith::ConstantIndexOp::create(), mlir::arith::ConstantIntOp::create(), mlir::arith::ConstantIntOp::create(), mlir::arith::ConstantIntOp::create(), mlir::OpBuilder::create(), mlir::Operation::create(), mlir::python::PyOperation::create(), mlir::arith::createArithExpandOpsPass(), mlir::createArithToAMDGPUConversionPass(), mlir::createArithToAMDGPUConversionPass(), mlir::emitc::createExpression(), createInlinedCompareImplementation(), createLinalgBodyCalculationForElementwiseOp(), createNonLdMatrixLoads(), mlir::shard::createShardingPropagation(), createSingleEntryBlock(), createSplitPart(), mlir::tosa::impl::TosaInferShapesPassBase< DerivedT >::createTosaInferShapesPass, mlir::xegpu::createVectorWithShapeFromValues(), mlir::irdl::createVerifier(), mlir::LivenessBlockInfo::currentlyLiveValues(), mlir::linalg::deduplicateOperandsAndRemoveDeadResults(), deduplicateOutputOperands(), mlir::shard::detail::defaultAddShardingAnnotations(), mlir::shard::detail::defaultGetShardingAnnotations(), dependsOnCarriedVals(), deriveStaticUpperBound(), dispatchParse(), mlir::Inliner::doInlining(), mlir::xegpu::doSCFStructuralTypeConversionWithTensorType(), mlir::Operation::dropAllUses(), mlir::linalg::dropUnitDims(), mlir::ValueBoundsConstraintSet::dump(), emit(), mlir::transform::expandTargetSpecification(), expandValue(), mlir::smt::exportSMTLIB(), extractBeneficiaryOps(), extractConvFilterSlices(), extractConvInputSlices(), extractConvResultSlices(), extractOperandTypes(), mlir::xegpu::extractVectorsWithShapeFromValue(), filterByValRefArgAttrs(), filterFuncAttributes(), mlir::LocationAttr::findInstanceOf(), mlir::LocationAttr::findInstanceOfOrUnknown(), mlir::presburger::IntegerRelation::findIntegerSample(), mlir::transform::TrackingListener::findReplacementOp(), mlir::presburger::IntegerRelation::findSymbolicIntegerLexMin(), findSymbolicIntegerLexOpt(), mlir::query::matcher::MatchFinder::flattenMatchedOps(), flattenValues(), flattenValues(), mlir::xegpu::flattenValues(), foldBinaryOpUnchecked(), foldCastOp(), mlir::op_definition_impl::foldTrait(), mlir::scf::forallToParallelLoop(), mlir::transform::detail::forwardTerminatorOperands(), forwardToUsers(), llvm::lsp::fromJSON(), mlir::lsp::fromJSON(), mlir::lsp::fromJSON(), mlir::linalg::fuseElementwiseOps(), mlir::sparse_tensor::FuncCallOrInlineGenerator< SubClass >::genCallOrInline(), mlir::OpState::genericParseProperties(), mlir::AffineMap::get(), mlir::DynamicOpDefinition::get(), mlir::LLVM::detail::LoopAnnotationTranslation::getAccessGroup(), mlir::linalg::getAffineResultPositions(), mlir::bufferization::OpWithUnstructuredControlFlowBufferizableOpInterfaceExternalModel< ConcreteModel, ConcreteOp >::getAliasingBranchOpOperands(), mlir::bufferization::BranchOpBufferizableOpInterfaceExternalModel< ConcreteModel, ConcreteOp >::getAliasingValues(), mlir::bufferization::func_ext::CallOpInterface::getAliasingValues(), mlir::linalg::getArityGroupAndKind(), mlir::presburger::MultiAffineFunction::getAsRelation(), mlir::impl::getAttrFromSortedRange(), mlir::MLIRContext::getAvailableDialects(), mlir::getBackwardSlice(), mlir::sparse_tensor::getBlockSize(), mlir::quant::UniformQuantizedSubChannelType::getBlockSizeInfo(), getBroadcastingMap(), mlir::bufferization::detail::getCallerOpOperands(), mlir::presburger::IntegerRelation::getConstantBoundOnDimSize64(), mlir::AffineMap::getConstantResults(), getConstraintPredicates(), mlir::tblgen::AttrOrTypeParameter::getDefaultValue(), mlir::Value::getDefiningOp(), mlir::affine::getDependenceComponents(), mlir::bufferization::impl::OwnershipBasedBufferDeallocationPassBase< DerivedT >::getDependentDialects(), mlir::tosa::impl::TosaInferShapesPassBase< DerivedT >::getDependentDialects(), mlir::vector::impl::LowerVectorMaskPassBase< DerivedT >::getDependentDialects(), mlir::xegpu::impl::XeGPUBlockingBase< DerivedT >::getDependentDialects(), mlir::memref::impl::NormalizeMemRefsPassBase< DerivedT >::getDescription(), mlir::xegpu::getDistributedVectorType(), mlir::xegpu::getDistributeLayoutAttr(), mlir::presburger::PresburgerRelation::getDomainSet(), mlir::presburger::getDynamicAPIntVec(), mlir::query::matcher::VariantMatcher::VariadicOpPayload::getDynMatcher(), mlir::presburger::IntegerRelation::getEmpty(), mlir::Token::getFloatingPointValue(), getForwardSliceImpl(), mlir::sparse_tensor::getFunc(), mlir::bufferization::func_ext::getFuncAnalysisState(), mlir::Token::getHashIdentifierNumber(), getIndicesForAccess(), mlir::getInnermostParallelLoops(), mlir::presburger::getInt64Vec(), mlir::Token::getIntTypeBitwidth(), mlir::tosa::ValueKnowledge::getKnowledgeFromType(), mlir::dataflow::AbstractSparseBackwardDataFlowAnalysis::getLatticeElements(), mlir::xegpu::getLayoutName(), mlir::presburger::MultiAffineFunction::getLexSet(), mlir::getLinearizedDimensions(), mlir::LLVM::detail::getLLVMConstant(), mlir::MLIRContext::getLoadedDialects(), mlir::tblgen::Pattern::getLocation(), getMangledSortHelperFunc(), getMaxLoopDepth(), getMemoryFootprintBytes(), mlir::memref::getMixedSizes(), mlir::tensor::getMixedSizes(), mlir::AliasAnalysis::getModRef(), mlir::LocalAliasAnalysis::getModRef(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineMap::getMultiDimMapWithTargets(), mlir::arith::impl::ArithExpandOpsPassBase< DerivedT >::getName(), mlir::impl::ArithToAMDGPUConversionPassBase< DerivedT >::getName(), mlir::impl::getNamedAttrFromSortedRange(), mlir::detail::OpResultImpl::getNextResultAtOffset(), mlir::tensor::getOrCreateDestinations(), getOrCreateFuncAnalysisState(), mlir::detail::OpResultImpl::getOwner(), mlir::linalg::getPackedOuterShapeWithoutTransposition(), mlir::arith::impl::ArithExpandOpsPassBase< DerivedT >::getPassName(), mlir::memref::impl::NormalizeMemRefsPassBase< DerivedT >::getPassName(), getPipelineStages(), mlir::presburger::PresburgerRelation::getRangeSet(), mlir::pdl_to_pdl_interp::PredicateBuilder::getResult(), mlir::tblgen::Operator::getResultDecorators(), mlir::shard::getResultShardings(), mlir::bufferization::getReturnOps(), getReturnTypes(), getSetDifference(), getSharding(), mlir::shard::getSharding(), mlir::getSlice(), getSliceContract(), mlir::Token::getStringValue(), mlir::xegpu::uArch::SubgroupMatrixMultiplyAcc::getSupportedShapes(), getSymbolUsesImpl(), getTransferFoldableInnerUnitDims(), mlir::Token::getUInt64IntegerValue(), getUnderlyingObjectSet(), mlir::presburger::PresburgerRelation::getUniverse(), mlir::presburger::PresburgerSet::getUniverse(), mlir::Token::getUnsignedIntegerValue(), mlir::SideEffects::EffectInstance< Effect >::getValue(), getValueTypes(), handleError(), handleError(), handleMultidimensionalVectors(), mlir::DialectInlinerInterface::handleResult(), mlir::InlinerInterface::handleResult(), handleResultImpl(), mlir::affine::hasDependence(), mlir::arm_sme::hasTileResult(), mlir::intrange::inferAdd(), mlir::intrange::inferCeilDivS(), mlir::intrange::inferCeilDivU(), mlir::intrange::inferDivS(), inferDivSRange(), mlir::intrange::inferDivU(), mlir::intrange::inferFloorDivS(), mlir::intrange::inferMul(), mlir::python::PyInferShapedTypeOpInterface::inferReturnTypeComponents(), mlir::python::PyInferTypeOpInterface::inferReturnTypes(), mlir::intrange::inferShapedDimOpInterface(), mlir::intrange::inferShl(), mlir::intrange::inferSub(), initTargetDefaultAttrs(), insertCasts(), mlir::bufferization::insertTensorCopies(), mlir::insideMutuallyExclusiveRegions(), mlir::presburger::IntegerRelation::intersect(), mlir::presburger::PresburgerRelation::intersect(), mlir::detail::AnalysisModel< AnalysisT >::invalidate(), mlir::sparse_tensor::inverseBlockSparsity(), mlir::sparse_tensor::isBlockSparsity(), isCompatibleImpl(), isContiguousLoadIdx(), mlir::Tester::isInteresting(), isLoopInvariantIdx(), mlir::affine::isOpwiseShiftValid(), isResultValueDead(), mlir::affine::isTilingValid(), mlir::Operation::isUsedOutsideOfBlock(), mlir::dataflow::PredecessorState::join(), mlir::tosa::ValueKnowledge::join(), linalgBroadcastAndMaybeExt(), linalgIntBroadcastExtSIAdd(), loadModule(), loadOperation(), mlir::IRMapping::lookup(), mlir::LLVM::ModuleTranslation::lookupOMPLoop(), mlir::CyclicReplacerCache< InT, OutT >::lookupOrInit(), mlir::CyclicReplacerCache< void *, const void * >::lookupOrInit(), mlir::detail::ConversionPatternRewriterImpl::lookupOrNull(), mlir::presburger::Simplex::makeProduct(), makeReductionGen(), mlir::LLVM::ModuleImport::mapBlock(), mlir::LLVM::ModuleTranslation::mapBlock(), mlir::LLVM::ModuleTranslation::mapBlockAddress(), mlir::LLVM::ModuleTranslation::mapBranch(), mlir::LLVM::ModuleTranslation::mapCall(), mlir::LLVM::ModuleTranslation::mapFunction(), mlir::LLVM::ModuleTranslation::mapUnresolvedBlockAddress(), mlir::detail::constant_op_binder< AttrT >::match(), mlir::query::matcher::PredicateBackwardSliceMatcher< BaseMatcher, Filter >::match(), matchAndReplaceDepthwiseConv(), Convert1DExtractStridedSliceIntoExtractInsertChain::matchAndRewrite(), GPUAllReduceConversion::matchAndRewrite(), GPUSubgroupReduceConversion::matchAndRewrite(), mlir::ComposeCollapseOfExpandOp< CollapseOpTy, ExpandOpTy, CastOpTy, DimOpTy, TensorTy >::matchAndRewrite(), mlir::PatternApplicator::matchAndRewrite(), OuterProductOpLowering::matchAndRewrite(), TransposeIsReshape::matchAndRewrite(), ContractOpToElementwise::matchAndRewriteMaskableOp(), materializeBinaryNanCheckIfRequired(), mlir::shard::maybeInsertTargetShardingAnnotation(), maybePickPermanentLayout(), mayHaveEffect(), mlir::tosa::ValueKnowledge::meet(), mlir::presburger::IntegerRelation::mergeAndCompose(), mgpuModuleLoadJIT(), minMaxBy(), mlirAffineMapCompressUnusedSymbols(), mlirAsmStateCreateForValue(), mlirLinalgInferContractionDimensions(), mlirLinalgInferConvolutionDimensions(), mlirMergeSymbolsIntoFromClone(), mlirStringRefCreate(), moveAndReset(), mlir::moveOperationDependencies(), mlir::moveValueDefinitions(), mustReachAtInnermost(), mlir::AnalysisManager::nest(), mlir::affine::noDependence(), nvvmInferResultRanges(), GpuAsyncRegionPass::SingleTokenUseCallback::operator()(), mlir::CachedCyclicReplacer< InT, OutT >::operator()(), mlir::operator<<(), mlir::operator<<(), mlir::OptionalParseResult::OptionalParseResult(), mlir::OptionalParseResult::OptionalParseResult(), mlir::bufferization::impl::OwnershipBasedBufferDeallocationPassBase< DerivedT >::OwnershipBasedBufferDeallocationPassBase(), mlir::linalg::pack(), packFunctionArguments(), packValInto64Bits(), padTileShapeToSize(), mlir::affine::AffineDmaStartOp::parse(), mlir::affine::AffineDmaWaitOp::parse(), mlir::detail::DenseArrayAttrImpl< T >::parse(), mlir::FieldParser< std::optional< AttributeT >, std::enable_if_t< std::is_base_of< Attribute, AttributeT >::value, std::optional< AttributeT > > >::parse(), mlir::FieldParser< std::optional< IntT >, std::enable_if_t< std::is_integral< IntT >::value, std::optional< IntT > > >::parse(), mlir::OpState::parse(), parseAffineMinMaxOp(), mlir::OpAsmParser::parseArgument(), mlir::OpAsmParser::parseArgumentList(), parseArithmeticExtendedBinaryOp(), mlir::AsmParser::parseArrowTypeList(), mlir::detail::AsmParserImpl< BaseT >::parseArrowTypeList(), mlir::OpAsmParser::parseAssignmentList(), mlir::AsmParser::parseAttribute(), mlir::AsmParser::parseAttribute(), mlir::AsmParser::parseAttribute(), mlir::AsmParser::parseAttribute(), mlir::AsmParser::parseAttribute(), mlir::detail::AsmParserImpl< BaseT >::parseAttribute(), mlir::detail::Parser::parseAttribute(), parseAttributions(), parseBound(), parseCallTypeAndResolveOperands(), parseCmpOp(), mlir::AsmParser::parseColonType(), mlir::AsmParser::parseColonType(), mlir::detail::AsmParserImpl< BaseT >::parseColonType(), mlir::AsmParser::parseColonTypeList(), mlir::detail::AsmParserImpl< BaseT >::parseColonTypeList(), parseCommonGlobalAndAlias(), parseCommonStructuredOpParts(), mlir::AsmParser::parseCustomAttributeWithFallback(), mlir::AsmParser::parseCustomAttributeWithFallback(), mlir::AsmParser::parseCustomAttributeWithFallback(), mlir::detail::AsmParserImpl< BaseT >::parseCustomAttributeWithFallback(), mlir::AsmParser::parseCustomTypeWithFallback(), mlir::AsmParser::parseCustomTypeWithFallback(), mlir::detail::AsmParserImpl< BaseT >::parseCustomTypeWithFallback(), mlir::AsmParser::parseDecimalInteger(), mlir::detail::Parser::parseDecOrHexAttr(), parseDstStyleOp(), parseElementAttrHexValues(), mlir::AsmParser::parseFloat(), mlir::AsmParser::parseFloat(), mlir::detail::AsmParserImpl< BaseT >::parseFloat(), mlir::detail::AsmParserImpl< BaseT >::parseFloat(), mlir::detail::Parser::parseFloatFromIntegerLiteral(), mlir::detail::Parser::parseFloatFromLiteral(), mlir::function_interface_impl::parseFunctionOp(), mlir::OpAsmParser::parseGenericOperationAfterOpName(), mlir::AsmParser::parseInteger(), mlir::AsmParser::parseKeywordOrString(), mlir::AsmParser::parseKeywordType(), mlir::query::matcher::internal::Parser::parseMatcherExpression(), parseNamedStructuredOp(), parseNamedStructuredOpRegion(), parseOneResultSameOperandTypeOp(), mlir::OpAsmParser::parseOperand(), mlir::OpAsmParser::parseOperandList(), mlir::OpAsmParser::parseOperandList(), mlir::OpAsmParser::parseOptionalArgument(), mlir::AsmParser::parseOptionalArrowTypeList(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalArrowTypeList(), mlir::AsmParser::parseOptionalAttrDict(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalAttrDict(), mlir::AsmParser::parseOptionalAttrDictWithKeyword(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalAttrDictWithKeyword(), mlir::AsmParser::parseOptionalAttribute(), mlir::AsmParser::parseOptionalAttribute(), mlir::AsmParser::parseOptionalAttribute(), mlir::AsmParser::parseOptionalAttribute(), mlir::AsmParser::parseOptionalAttribute(), mlir::AsmParser::parseOptionalAttribute(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalAttribute(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalAttribute(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalAttribute(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalAttribute(), mlir::detail::Parser::parseOptionalAttribute(), mlir::detail::Parser::parseOptionalAttribute(), mlir::AsmParser::parseOptionalColonTypeList(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalColonTypeList(), mlir::AsmParser::parseOptionalDecimalInteger(), mlir::AsmParser::parseOptionalDecimalInteger(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalDecimalInteger(), mlir::detail::Parser::parseOptionalDecimalInteger(), mlir::AsmParser::parseOptionalInteger(), mlir::AsmParser::parseOptionalInteger(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalInteger(), mlir::detail::Parser::parseOptionalInteger(), mlir::AsmParser::parseOptionalKeywordOrString(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalKeywordOrString(), mlir::detail::Parser::parseOptionalKeywordOrString(), mlir::OpAsmParser::parseOptionalLocationSpecifier(), mlir::OpAsmParser::parseOptionalOperand(), parseOptionalStaticSlice(), mlir::AsmParser::parseOptionalSymbolName(), mlir::AsmParser::parseOptionalSymbolName(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalSymbolName(), mlir::AsmParser::parseOptionalType(), mlir::detail::AsmParserImpl< BaseT >::parseOptionalType(), mlir::Op< ConcreteType, Traits >::parseProperties(), mlir::AsmParser::parseResourceHandle(), mlir::detail::Parser::parseResourceHandle(), parseSameOperandTypeVariadicToBoolOp(), parseStorageType(), mlir::detail::Parser::parseStridedLayoutAttr(), mlir::AsmParser::parseSymbolName(), mlir::AsmParser::parseSymbolName(), parseTargetAllocMemOp(), mlir::OpAsmParser::parseTrailingOperandList(), mlir::AsmParser::parseType(), mlir::AsmParser::parseType(), mlir::detail::AsmParserImpl< BaseT >::parseType(), mlir::AsmParser::parseTypeList(), mlir::linalg::partitionLinalgOpWithShardedReduction(), mlir::shard::partitionOperation(), populateDialectLLVMSubmodule(), populateDialectSMTSubmodule(), populateTransformInterpreterSubmodule(), mlir::presburger::Matrix< T >::postMultiplyWithColumn(), mlir::presburger::Matrix< T >::preMultiplyWithRow(), mlir::pdl_to_pdl_interp::OptimalBranching::preOrderTraversal(), mlir::Liveness::print(), mlir::pdll::ods::Context::print(), mlir::Value::printAsOperand(), printConstantOp(), mlir::OpAsmPrinter::printFunctionalType(), printFunctionBody(), printOperation(), produceSliceErrorMsg(), produceSubViewErrorMsg(), pruneNonTransposedDims(), pyListToVector(), mlir::quant::UniformQuantizedValueConverter::quantizeFloatToInt(), mlir::presburger::IntegerRelation::rangeProduct(), mlir::DialectBytecodeReader::readAttribute(), mlir::DialectBytecodeReader::readAttribute(), readBits(), mlir::DialectBytecodeReader::readBlob(), mlir::DialectBytecodeReader::readBool(), mlir::DialectBytecodeReader::readList(), mlir::DialectBytecodeReader::readOptionalAttribute(), mlir::DialectBytecodeReader::readResourceHandle(), mlir::readResourceHandle(), mlir::DialectBytecodeReader::readSignedVarInt(), mlir::DialectBytecodeReader::readSignedVarInts(), mlir::DialectBytecodeReader::readString(), mlir::DialectBytecodeReader::readType(), mlir::DialectBytecodeReader::readType(), mlir::DialectBytecodeReader::readVarInt(), mlir::DialectBytecodeReader::readVarIntWithFlag(), reduceMatchAndRewriteHelper(), mlir::reifyResultShapes(), mlir::xegpu::removeLayoutAttrs(), replaceConstantUsesOf(), replaceElementImpl(), mlir::detail::AttrTypeReplacerBase< Concrete >::replaceElementsIn(), replaceForAllWithNewSignature(), mlir::detail::ConversionPatternRewriterImpl::replaceOp(), replaceSubElements(), reshapeLoad(), reshapeStore(), mlir::Liveness::resolveLiveness(), mlir::OpAsmParser::resolveOperand(), mlir::OpAsmParser::resolveOperands(), mlir::OpAsmParser::resolveOperands(), mlir::OpAsmParser::resolveOperands(), resolveValues(), mlir::ResultRange::ResultRange(), mlir::linalg::rewriteInDestinationPassingStyle(), mlir::linalg::rewriteInIm2Col(), mlir::linalg::rewriteInIm2Col(), mlir::linalg::rewriteInIm2Col(), rewriteOneForallCommonImpl(), mlir::PassManager::run(), mlir::dataflow::RunLivenessAnalysis::RunLivenessAnalysis(), scalarizeVectorOpHelper(), mlir::xegpu::setDistributeLayoutAttrs(), mlir::MutableAffineMap::setResult(), shouldBeInlined(), mlir::LLVM::DIExpressionRewriter::simplify(), mlir::presburger::PresburgerRelation::simplify(), mlir::tblgen::FmtObjectBase::sstr(), mlir::StateStack::stackWalk(), mlir::dataflow::staticallyNonNegative(), staticallyOutOfBounds(), mlir::tblgen::FmtObjectBase::str(), mlir::LLVMTypeConverter::structFuncArgTypeConverter, mlir::structFuncArgTypeConverter(), mlir::presburger::PresburgerRelation::subtract(), subtractExprs(), tiebreakLex(), tileLinalgOpImpl(), mlir::linalg::tileReductionUsingForall(), transformToStructuredCFBranches(), mlir::LLVM::detail::DebugImporter::translate(), translateDataLayout(), mlir::LLVM::detail::DebugImporter::translateLoc(), tryApply(), tryToParseIntListImpl(), mlir::TypeRange::TypeRange(), mlir::presburger::PresburgerRelation::unionSet(), unpackOperandVector(), unpackSingleIndexResultPayloadOperations(), unpackSingleIndexResultPayloadOperations(), mlir::vector::unrollVectorOp(), unwrap(), mlir::ReductionNode::update(), updateCalls(), updateControlFlowOps(), updateOp(), mlir::ResultRange::use_empty(), mlir::presburger::MultiAffineFunction::valueAt(), vectorizeAsLinalgGeneric(), vectorizeOneOp(), mlir::irdl::ConstraintVerifier::verify(), verifyAllToAllOperandAndResultShape(), mlir::detail::verifyDestinationStyleOpInterface(), verifyGatherOperandAndResultShape(), mlir::verifyInBoundsSlice(), mlir::detail::verifyInferredResultTypes(), mlir::transform::detail::verifyParamProducerTransformOpTrait(), verifyQuantifierRegions(), mlir::OpTrait::impl::verifySameOperandsAndResultElementType(), verifyScatterOrSliceOperandAndResultShape(), mlir::detail::verifySymbolTable(), mlir::transform::detail::verifyTransformOpInterface(), visit(), mlir::dataflow::LivenessAnalysis::visitBranchOperand(), mlir::dataflow::AbstractSparseBackwardDataFlowAnalysis::visitCallableOperation(), visitOp(), mlir::dataflow::IntegerRangeAnalysis::visitOperation(), mlir::dataflow::StridedMetadataRangeAnalysis::visitOperation(), mlir::Block::walk(), mlir::detail::walk(), mlir::detail::walk(), mlir::detail::walk(), mlir::Region::walk(), walkSymbolTable(), walkSymbolTable(), widenOp(), wrapExternalFunction(), xegpu::removeLayoutAttr< mlir::OpResult >(), xegpu::setDistributeLayoutAttr< mlir::OpResult >(), and mlir::SimpleAffineExprFlattener::~SimpleAffineExprFlattener().

target

target

Definition at line 2099 of file LinalgTransformOps.cpp.

Referenced by mlir::transform::gpu::alterGpuLaunch(), applyConversion(), applyTilingToAll(), mlir::transform::detail::applyTransformToEach(), buildPredicateList(), mlir::bufferization::castOrReallocMemRefValue(), mlir::impl::LinalgDetensorizePassBase< DerivedT >::classof(), mlir::math::impl::MathExtendToSupportedTypesBase< DerivedT >::classof(), mlir::arith::impl::ArithIntRangeOptsBase< DerivedT >::clonePass(), mlir::impl::LinalgDetensorizePassBase< DerivedT >::clonePass(), mlir::impl::LowerSparseOpsToForeachBase< DerivedT >::clonePass(), mlir::impl::PreSparsificationRewriteBase< DerivedT >::clonePass(), mlir::impl::SparseBufferRewriteBase< DerivedT >::clonePass(), mlir::configureAMXLegalizeForExportTarget(), mlir::configureArmSMEToLLVMConversionLegality(), mlir::configureArmSVELegalizeForExportTarget(), mlir::configureGpuToNVVMConversionLegality(), mlir::configureGpuToROCDLConversionLegality(), mlir::configureOpenMPToLLVMConversionLegality(), mlir::configureParallelLoopToGPULegality(), mlir::configureX86VectorLegalizeForExportTarget(), contract(), mlir::convertAffineLoopNestToGPULaunch(), mlir::OperationConverter::convertOperations(), mlir::memref::createExpandOpsPass(), mlir::quant::createNormalizeQuantTypes(), mlir::xegpu::createXeGPUWgToSgDistribute(), DEFINE_C_API_STRUCT(), definiteFailureHelper(), doit(), mlir::xegpu::doSCFStructuralTypeConversionWithTensorType(), mlir::transform::expandTargetSpecification(), findPayloadRoot(), mlir::transform::gpu::findTopLevelForallOp(), mlir::fuseIndependentSiblingForallLoops(), mlir::fuseIndependentSiblingForLoops(), mlir::SPIRVConversionTarget::get(), mlir::spirv::impl::SPIRVCompositeTypeLayoutPassBase< DerivedT >::getArgument(), mlir::impl::LowerSparseOpsToForeachBase< DerivedT >::getArgumentName(), mlir::impl::PreSparsificationRewriteBase< DerivedT >::getArgumentName(), mlir::impl::SparseBufferRewriteBase< DerivedT >::getArgumentName(), mlir::tosa::impl::TosaConvertIntegerTypeToSignlessBase< DerivedT >::getArgumentName(), mlir::impl::LowerSparseOpsToForeachBase< DerivedT >::getDependentDialects(), mlir::impl::PreSparsificationRewriteBase< DerivedT >::getDependentDialects(), mlir::memref::impl::ExpandReallocPassBase< DerivedT >::getDependentDialects(), mlir::quant::impl::NormalizeQuantTypesBase< DerivedT >::getDependentDialects(), mlir::impl::LowerSparseOpsToForeachBase< DerivedT >::getDescription(), mlir::impl::SparseBufferRewriteBase< DerivedT >::getDescription(), mlir::spirv::impl::SPIRVCompositeTypeLayoutPassBase< DerivedT >::getDescription(), mlir::spirv::getMemorySpaceToStorageClassTarget(), mlir::impl::SparseBufferRewriteBase< DerivedT >::getName(), mlir::spirv::impl::SPIRVCompositeTypeLayoutPassBase< DerivedT >::getName(), mlir::LLVM::ModuleToObject::getOrCreateTargetMachine(), mlir::impl::SparseBufferRewriteBase< DerivedT >::getPassName(), mlir::LLVM::detail::getTargetMachine(), isForallWithIdenticalConfiguration(), isForWithIdenticalConfiguration(), isOpSibling(), mlir::transform::gpu::mapNestedForallToThreadsImpl(), mlir::transform::detail::mergeSymbolsInto(), MLIR_DECLARE_CAPI_DIALECT_REGISTRATION(), mlirGPUObjectAttrGet(), mlirGPUObjectAttrGetWithKernels(), mlirMergeSymbolsIntoFromClone(), mlirRegionTakeBody(), NB_MODULE(), mlir::OperationConverter::OperationConverter(), mlir::impl::SparseAssemblerBase< DerivedT >::operator=(), mlir::amdgpu::populateAmdgpuEmulateAtomicsPatterns(), mlir::populateAsyncFuncToAsyncRuntimeConversionPatterns(), mlir::populateAsyncStructuralTypeConversionsAndLegality(), mlir::cf::populateCFStructuralTypeConversionsAndLegality(), mlir::cf::populateCFStructuralTypeConversionTarget(), mlir::populateConversionTargetFromOperation(), mlir::populateConversionTargetFromOperation(), mlir::ConvertToEmitCPatternInterface::populateConvertToEmitCConversionPatterns(), mlir::ConvertToLLVMPatternInterface::populateConvertToLLVMConversionPatterns(), mlir::arith::populateEmulateUnsupportedFloatsLegality(), mlir::math::populateExtendToSupportedTypesConversionTarget(), mlir::math::populateExtendToSupportedTypesTypeConverter(), mlir::populateOpConvertToLLVMConversionPatterns(), mlir::scf::populateSCFStructuralTypeConversionsAndLegality(), mlir::scf::populateSCFStructuralTypeConversionTarget(), populateTransformInterpreterSubmodule(), mlir::populateXeVMToLLVMConversionPatterns(), mlir::tensor::preservesStaticInformation(), mlir::spirv::Deserializer::processBranch(), mlir::affine::promoteSingleIterationLoops(), remapArgumentEffects(), remapEffects(), mlir::shard::reshard(), mlir::shard::reshard(), ConvertMathToROCDLPass::runOnOperation(), setOptionalCommandlineArguments(), sortDependentLoops(), stripmineSink(), stripmineSink(), mlir::affine::tile(), mlir::tile(), mlir::transform::tileToForallOpImpl(), unpackOptionalValues(), updateSymbolAndAllUses(), mlir::wasmssa::detail::verifyLabelLevelInterface(), mlir::spirv::Deserializer::wireUpBlockArgument(), mlir::impl::LowerSparseOpsToForeachBase< DerivedT >::~LowerSparseOpsToForeachBase(), and mlir::impl::SparseAssemblerBase< DerivedT >::~SparseAssemblerBase().

TypeRange

TypeRange {resultType, resultType}

Definition at line 2098 of file LinalgTransformOps.cpp.

Referenced by buildStructuredOp(), mlir::detail::ConversionPatternRewriterImpl::buildUnresolvedMaterialization(), createAsyncDispatchFunction(), createHeapSortFunc(), mlir::LLVM::createPrintStrCall(), createQuickSortFunc(), doSequentialDispatch(), mlir::query::extractFunction(), genEndInsert(), mlir::linalg::generalizeNamedOp(), genExpand(), mlir::sparse_tensor::genTuple(), mlir::spirv::StructType::getElementTypes(), getParallelComputeFunctionType(), getReplacementValues(), legalizeUnresolvedMaterialization(), mlir::LLVMTypeConverter::LLVMTypeConverter(), mlir::detail::ConversionPatternRewriterImpl::lookupOrDefault(), mlir::detail::ConversionPatternRewriterImpl::lookupOrNull(), mlir::GPUReturnOpLowering::matchAndRewrite(), mlir::linalg::LinalgOpToLibraryCallRewrite::matchAndRewrite(), matchAndRewriteSortOp(), outlineKernelFuncImpl(), packRankedMemRefDesc(), packUnrankedMemRefDesc(), mlir::printSemiFunctionType(), mlir::detail::ConversionPatternRewriterImpl::replaceOp(), mlir::sparse_tensor::SparsificationAndBufferizationPass::runDenseBufferization(), mlir::SparseIterationTypeConverter::SparseIterationTypeConverter(), mlir::linalg::splitReduction(), and mlir::detail::verifyTypesAlongControlFlowEdges().

ValueRange

b ValueRange {}

Definition at line 2102 of file LinalgTransformOps.cpp.

Referenced by addNodeToMDG(), addResumeFunction(), mlir::affine::affineParallelize(), mlir::transform::TransformEachOpTrait< OpTy >::apply(), buildAffineLoopFromConstants(), buildAffineLoopNestImpl(), mlir::bufferization::buildDeallocationLibraryFunction(), mlir::scf::buildLoopNest(), checkAndNestUnderRewriteOp(), mlir::computeLinearIndex(), createAsyncDispatchFunction(), createParallelComputeFunction(), createQuickSortFunc(), doSequentialDispatch(), mlir::query::extractFunction(), foldExtractFromBroadcast(), fuseWithReshapeByExpansion(), genCoIterateBranchNest(), mlir::sparse_tensor::genCoIteration(), mlir::linalg::generalizeNamedOp(), genLoopWithIterator(), mlir::sparse_tensor::SparseIterator::getCursor(), mlir::impl::OutlineShapeComputationPassBase< DerivedT >::getDependentDialects(), getReplacementValues(), mlir::ValueShapeRange::getValues(), linalgBroadcastAndMaybeExt(), linalgIntBroadcastExtSIAdd(), mlir::detail::ConversionPatternRewriterImpl::lookupOrDefault(), mlir::detail::ConversionPatternRewriterImpl::lookupOrNull(), mlir::loopUnrollByFactor(), ExecuteRegionForwardingEliminator::matchAndRewrite(), movePaddingToFillOrGenericOp(), outlineKernelFuncImpl(), mlir::parallelLoopUnrollByFactors(), mlir::detail::ConversionPatternRewriterImpl::remapValues(), mlir::detail::ConversionPatternRewriterImpl::replaceOp(), mlir::linalg::rewriteAsPaddedOp(), mlir::linalg::rewriteAsPaddedOp(), mlir::linalg::rewriteInDestinationPassingStyle(), mlir::linalg::rewriteInDestinationPassingStyle(), mlir::linalg::splitReduction(), and mlir::linalg::splitReductionByScaling().