#include "mlir/Dialect/GPU/Utils/DistributionUtils.h"
Inheritance diagram for mlir::gpu::WarpDistributionPattern:Public Types | |
| using | Base = WarpDistributionPattern |
| Public Types inherited from mlir::OpRewritePattern< WarpExecuteOnLane0Op > | |
| using | Base = OpRewritePattern |
Type alias to allow derived classes to inherit constructors with using Base::Base;. More... | |
Public Member Functions | |
| virtual LogicalResult | matchAndRewrite (WarpExecuteOnLane0Op op, PatternRewriter &rewriter) const override=0 |
| OpRewritePattern (MLIRContext *context, PatternBenefit benefit=1, ArrayRef< StringRef > generatedNames={}) | |
| Patterns must specify the root operation name they match against, and can also specify the benefit of the pattern matching and a list of generated ops. More... | |
| Public Member Functions inherited from mlir::OpRewritePattern< WarpExecuteOnLane0Op > | |
| OpRewritePattern (MLIRContext *context, PatternBenefit benefit=1, ArrayRef< StringRef > generatedNames={}) | |
| Patterns must specify the root operation name they match against, and can also specify the benefit of the pattern matching and a list of generated ops. More... | |
| Public Member Functions inherited from mlir::detail::OpOrInterfaceRewritePatternBase< SourceOp > | |
| LogicalResult | matchAndRewrite (Operation *op, PatternRewriter &rewriter) const final |
| Wrapper around the RewritePattern method that passes the derived op type. More... | |
| virtual LogicalResult | matchAndRewrite (SourceOp op, PatternRewriter &rewriter) const =0 |
| Method that operates on the SourceOp type. More... | |
| Public Member Functions inherited from mlir::RewritePattern | |
| virtual | ~RewritePattern ()=default |
| Public Member Functions inherited from mlir::Pattern | |
| ArrayRef< OperationName > | getGeneratedOps () const |
| Return a list of operations that may be generated when rewriting an operation instance with this pattern. More... | |
| std::optional< OperationName > | getRootKind () const |
| Return the root node that this pattern matches. More... | |
| std::optional< TypeID > | getRootInterfaceID () const |
| Return the interface ID used to match the root operation of this pattern. More... | |
| std::optional< TypeID > | getRootTraitID () const |
| Return the trait ID used to match the root operation of this pattern. More... | |
| PatternBenefit | getBenefit () const |
| Return the benefit (the inverse of "cost") of matching this pattern. More... | |
| bool | hasBoundedRewriteRecursion () const |
| Returns true if this pattern is known to result in recursive application, i.e. More... | |
| MLIRContext * | getContext () const |
| Return the MLIRContext used to create this pattern. More... | |
| StringRef | getDebugName () const |
| Return a readable name for this pattern. More... | |
| void | setDebugName (StringRef name) |
| Set the human readable debug name used for this pattern. More... | |
| ArrayRef< StringRef > | getDebugLabels () const |
| Return the set of debug labels attached to this pattern. More... | |
| void | addDebugLabels (ArrayRef< StringRef > labels) |
| Add the provided debug labels to this pattern. More... | |
| void | addDebugLabels (StringRef label) |
Protected Member Functions | |
| OpOperand * | getWarpResult (WarpExecuteOnLane0Op warpOp, llvm::function_ref< bool(Operation *)> fn) const |
Return a value yielded by warpOp which statifies the filter lamdba condition and is not dead. More... | |
| WarpExecuteOnLane0Op | moveRegionToNewWarpOpAndReplaceReturns (RewriterBase &rewriter, WarpExecuteOnLane0Op warpOp, ValueRange newYieldedValues, TypeRange newReturnTypes) const |
| Helper to create a new WarpExecuteOnLane0Op with different signature. More... | |
| WarpExecuteOnLane0Op | moveRegionToNewWarpOpAndAppendReturns (RewriterBase &rewriter, WarpExecuteOnLane0Op warpOp, ValueRange newYieldedValues, TypeRange newReturnTypes, SmallVector< size_t > &indices) const |
| Helper to create a new WarpExecuteOnLane0Op region with extra outputs. More... | |
| bool | delinearizeLaneId (OpBuilder &builder, Location loc, ArrayRef< int64_t > originalShape, ArrayRef< int64_t > distributedShape, int64_t warpSize, Value laneId, SmallVectorImpl< Value > &delinearizedIds) const |
Delinearize the given laneId into multiple dimensions, where each dimension's size is determined by originalShape and distributedShape together. More... | |
| Protected Member Functions inherited from mlir::RewritePattern | |
| Pattern (StringRef rootName, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
Inherit the base constructors from Pattern. More... | |
| Pattern (MatchAnyOpTypeTag tag, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
Inherit the base constructors from Pattern. More... | |
| Pattern (MatchInterfaceOpTypeTag tag, TypeID interfaceID, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
Inherit the base constructors from Pattern. More... | |
| Pattern (MatchTraitOpTypeTag tag, TypeID traitID, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
Inherit the base constructors from Pattern. More... | |
| Protected Member Functions inherited from mlir::Pattern | |
| Pattern (StringRef rootName, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
| Construct a pattern with a certain benefit that matches the operation with the given root name. More... | |
| Pattern (MatchAnyOpTypeTag tag, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
| Construct a pattern that may match any operation type. More... | |
| Pattern (MatchInterfaceOpTypeTag tag, TypeID interfaceID, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
Construct a pattern that may match any operation that implements the interface defined by the provided interfaceID. More... | |
| Pattern (MatchTraitOpTypeTag tag, TypeID traitID, PatternBenefit benefit, MLIRContext *context, ArrayRef< StringRef > generatedNames={}) | |
Construct a pattern that may match any operation that implements the trait defined by the provided traitID. More... | |
| void | setHasBoundedRewriteRecursion (bool hasBoundedRecursionArg=true) |
| Set the flag detailing if this pattern has bounded rewrite recursion or not. More... | |
Additional Inherited Members | |
| Static Public Member Functions inherited from mlir::RewritePattern | |
| template<typename T , typename... Args> | |
| static std::unique_ptr< T > | create (Args &&...args) |
This method provides a convenient interface for creating and initializing derived rewrite patterns of the given type T. More... | |
Detailed Description
Definition at line 19 of file DistributionUtils.h.
Member Typedef Documentation
◆ Base
Definition at line 21 of file DistributionUtils.h.
Member Function Documentation
◆ delinearizeLaneId()
|
protected |
Delinearize the given laneId into multiple dimensions, where each dimension's size is determined by originalShape and distributedShape together.
This function expects the total numbers of threads needed for distribution is equal to warpSize. Returns true and updates delinearizedIds if so.
Definition at line 104 of file DistributionUtils.cpp.
References mlir::bindSymbols(), mlir::arith::ConstantIndexOp::create(), mlir::AffineExpr::floorDiv(), mlir::Builder::getContext(), and mlir::affine::makeComposedAffineApply().
◆ getWarpResult()
|
protected |
Return a value yielded by warpOp which statifies the filter lamdba condition and is not dead.
Definition at line 89 of file DistributionUtils.cpp.
References mlir::Value::getDefiningOp().
◆ matchAndRewrite()
|
overridepure virtual |
◆ moveRegionToNewWarpOpAndAppendReturns()
|
protected |
Helper to create a new WarpExecuteOnLane0Op region with extra outputs.
indices return the index of each new output.
Definition at line 54 of file DistributionUtils.cpp.
References mlir::detail::enumerate(), moveRegionToNewWarpOpAndReplaceReturns(), and mlir::RewriterBase::replaceOp().
◆ moveRegionToNewWarpOpAndReplaceReturns()
|
protected |
Helper to create a new WarpExecuteOnLane0Op with different signature.
Definition at line 26 of file DistributionUtils.cpp.
References mlir::Region::begin(), mlir::RewriterBase::eraseBlock(), mlir::Region::front(), mlir::Region::getBlocks(), mlir::RewriterBase::inlineRegionBefore(), mlir::RewriterBase::modifyOpInPlace(), and mlir::OpBuilder::setInsertionPoint().
Referenced by moveRegionToNewWarpOpAndAppendReturns().
◆ OpRewritePattern()
|
inline |
Patterns must specify the root operation name they match against, and can also specify the benefit of the pattern matching and a list of generated ops.
Definition at line 322 of file PatternMatch.h.
The documentation for this struct was generated from the following files:
- include/mlir/Dialect/GPU/Utils/DistributionUtils.h
- lib/Dialect/GPU/Utils/DistributionUtils.cpp
