OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 2
Code Issues Packages Projects Releases Wiki Activity

Revert "[MLIR] Update pass declarations to new autogenerated files" 

This reverts commit 2be8af8f0e.
This commit is contained in:
Michele Scuttari 2022-08-30 22:20:36 +02:00
parent 349e5bd24e
commit 039b969b32
No known key found for this signature in database
GPG Key ID: E79E7BDFEE4B62D4
301 changed files with 2061 additions and 2435 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
flang/include/flang/Optimizer/CodeGen/CGPasses.td
View File

@ -16,7 +16,7 @@
include "mlir/Pass/PassBase.td"
def FIRToLLVMLoweringPass : Pass<"fir-to-llvm-ir", "mlir::ModuleOp"> {
def FIRToLLVMLowering : Pass<"fir-to-llvm-ir", "mlir::ModuleOp"> {
let summary = "Convert FIR dialect to LLVM-IR dialect";
let description = [{
Convert the FIR dialect to the LLVM-IR dialect of MLIR. This conversion
@ -30,7 +30,7 @@ def FIRToLLVMLoweringPass : Pass<"fir-to-llvm-ir", "mlir::ModuleOp"> {
];
}
def CodeGenRewritePass : Pass<"cg-rewrite", "mlir::ModuleOp"> {
def CodeGenRewrite : Pass<"cg-rewrite", "mlir::ModuleOp"> {
let summary = "Rewrite some FIR ops into their code-gen forms.";
let description = [{
Fuse specific subgraphs into single Ops for code generation.
@ -44,7 +44,7 @@ def CodeGenRewritePass : Pass<"cg-rewrite", "mlir::ModuleOp"> {
];
}
def TargetRewritePass : Pass<"target-rewrite", "mlir::ModuleOp"> {
def TargetRewrite : Pass<"target-rewrite", "mlir::ModuleOp"> {
let summary = "Rewrite some FIR dialect into target specific forms.";
let description = [{
Certain abstractions in the FIR dialect need to be rewritten to reflect

6
flang/include/flang/Optimizer/CodeGen/CodeGen.h
View File

@ -20,12 +20,6 @@ namespace fir {
struct NameUniquer;
#define GEN_PASS_DECL_FIRTOLLVMLOWERINGPASS
#define GEN_PASS_DECL_CODEGENREWRITEPASS
#define GEN_PASS_DECL_TARGETREWRITEPASS
#define GEN_PASS_DECL_BOXEDPROCEDUREPASS
#include "flang/Optimizer/CodeGen/CGPasses.h.inc"
/// Prerequiste pass for code gen. Perform intermediate rewrites to perform
/// the code gen (to LLVM-IR dialect) conversion.
std::unique_ptr<mlir::Pass> createFirCodeGenRewritePass();

16
flang/include/flang/Optimizer/Transforms/Passes.h
View File

@ -28,22 +28,6 @@ namespace fir {
// Passes defined in Passes.td
//===----------------------------------------------------------------------===//
#define GEN_PASS_DECL_ABSTRACTRESULTONFUNCOPTPASS
#define GEN_PASS_DECL_ABSTRACTRESULTONGLOBALOPTPASS
#define GEN_PASS_DECL_AFFINEDIALECTPROMOTIONPASS
#define GEN_PASS_DECL_AFFINEDIALECTDEMOTIONPASS
#define GEN_PASS_DECL_ANNOTATECONSTANTOPERANDSPASS
#define GEN_PASS_DECL_ARRAYVALUECOPYPASS
#define GEN_PASS_DECL_CHARACTERCONVERSIONPASS
#define GEN_PASS_DECL_CFGCONVERSIONPASS
#define GEN_PASS_DECL_EXTERNALNAMECONVERSIONPASS
#define GEN_PASS_DECL_MEMREFDATAFLOWOPTPASS
#define GEN_PASS_DECL_SIMPLIFYINTRINSICSPASS
#define GEN_PASS_DECL_MEMORYALLOCATIONOPTPASS
#define GEN_PASS_DECL_SIMPLIFYREGIONLITEPASS
#define GEN_PASS_DECL_ALGEBRAICSIMPLIFICATIONPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
std::unique_ptr<mlir::Pass> createAbstractResultOnFuncOptPass();
std::unique_ptr<mlir::Pass> createAbstractResultOnGlobalOptPass();
std::unique_ptr<mlir::Pass> createAffineDemotionPass();

30
flang/include/flang/Optimizer/Transforms/Passes.td
View File

@ -16,7 +16,7 @@
include "mlir/Pass/PassBase.td"
class AbstractResultOptBase<string optExt, string operation>
class AbstractResultOptBase<string optExt, string operation>
: Pass<"abstract-result-on-" # optExt # "-opt", operation> {
let summary = "Convert fir.array, fir.box and fir.rec function result to "
"function argument";
@ -35,15 +35,15 @@ class AbstractResultOptBase<string optExt, string operation>
];
}
def AbstractResultOnFuncOptPass : AbstractResultOptBase<"func", "mlir::func::FuncOp"> {
def AbstractResultOnFuncOpt : AbstractResultOptBase<"func", "mlir::func::FuncOp"> {
let constructor = "::fir::createAbstractResultOnFuncOptPass()";
}
def AbstractResultOnGlobalOptPass : AbstractResultOptBase<"global", "fir::GlobalOp"> {
def AbstractResultOnGlobalOpt : AbstractResultOptBase<"global", "fir::GlobalOp"> {
let constructor = "::fir::createAbstractResultOnGlobalOptPass()";
}
def AffineDialectPromotionPass : Pass<"promote-to-affine", "::mlir::func::FuncOp"> {
def AffineDialectPromotion : Pass<"promote-to-affine", "::mlir::func::FuncOp"> {
let summary = "Promotes `fir.{do_loop,if}` to `affine.{for,if}`.";
let description = [{
Convert fir operations which satisfy affine constraints to the affine
@ -68,7 +68,7 @@ def AffineDialectPromotionPass : Pass<"promote-to-affine", "::mlir::func::FuncOp
];
}
def AffineDialectDemotionPass : Pass<"demote-affine", "::mlir::func::FuncOp"> {
def AffineDialectDemotion : Pass<"demote-affine", "::mlir::func::FuncOp"> {
let summary = "Converts `affine.{load,store}` back to fir operations";
let description = [{
Affine dialect's default lowering for loads and stores is different from
@ -82,7 +82,7 @@ def AffineDialectDemotionPass : Pass<"demote-affine", "::mlir::func::FuncOp"> {
];
}
def AnnotateConstantOperandsPass : Pass<"annotate-constant"> {
def AnnotateConstantOperands : Pass<"annotate-constant"> {
let summary = "Annotate constant operands to all FIR operations";
let description = [{
The MLIR canonicalizer makes a distinction between constants based on how
@ -99,7 +99,7 @@ def AnnotateConstantOperandsPass : Pass<"annotate-constant"> {
let dependentDialects = [ "fir::FIROpsDialect" ];
}
def ArrayValueCopyPass : Pass<"array-value-copy", "::mlir::func::FuncOp"> {
def ArrayValueCopy : Pass<"array-value-copy", "::mlir::func::FuncOp"> {
let summary = "Convert array value operations to memory operations.";
let description = [{
Transform the set of array value primitives to a memory-based array
@ -119,7 +119,7 @@ def ArrayValueCopyPass : Pass<"array-value-copy", "::mlir::func::FuncOp"> {
let dependentDialects = [ "fir::FIROpsDialect" ];
}
def CharacterConversionPass : Pass<"character-conversion"> {
def CharacterConversion : Pass<"character-conversion"> {
let summary = "Convert CHARACTER entities with different KINDs";
let description = [{
Translates entities of one CHARACTER KIND to another.
@ -137,7 +137,7 @@ def CharacterConversionPass : Pass<"character-conversion"> {
];
}
def CFGConversionPass : Pass<"cfg-conversion", "::mlir::func::FuncOp"> {
def CFGConversion : Pass<"cfg-conversion", "::mlir::func::FuncOp"> {
let summary = "Convert FIR structured control flow ops to CFG ops.";
let description = [{
Transform the `fir.do_loop`, `fir.if`, and `fir.iterate_while` ops into
@ -157,7 +157,7 @@ def CFGConversionPass : Pass<"cfg-conversion", "::mlir::func::FuncOp"> {
];
}
def ExternalNameConversionPass : Pass<"external-name-interop", "mlir::ModuleOp"> {
def ExternalNameConversion : Pass<"external-name-interop", "mlir::ModuleOp"> {
let summary = "Convert name for external interoperability";
let description = [{
Demangle FIR internal name and mangle them for external interoperability.
@ -165,7 +165,7 @@ def ExternalNameConversionPass : Pass<"external-name-interop", "mlir::ModuleOp">
let constructor = "::fir::createExternalNameConversionPass()";
}
def MemRefDataFlowOptPass : Pass<"fir-memref-dataflow-opt", "::mlir::func::FuncOp"> {
def MemRefDataFlowOpt : Pass<"fir-memref-dataflow-opt", "::mlir::func::FuncOp"> {
let summary =
"Perform store/load forwarding and potentially removing dead stores.";
let description = [{
@ -182,7 +182,7 @@ def MemRefDataFlowOptPass : Pass<"fir-memref-dataflow-opt", "::mlir::func::FuncO
// functions into the module, which is invalid if a finer grain mlir::Operation
// is used as the pass specification says to not touch things outside hte scope
// of the operation being processed.
def SimplifyIntrinsicsPass : Pass<"simplify-intrinsics", "mlir::ModuleOp"> {
def SimplifyIntrinsics : Pass<"simplify-intrinsics", "mlir::ModuleOp"> {
let summary = "Intrinsics simplification";
let description = [{
Qualifying intrinsics calls are replaced with calls to a specialized and
@ -198,7 +198,7 @@ def SimplifyIntrinsicsPass : Pass<"simplify-intrinsics", "mlir::ModuleOp"> {
];
}
def MemoryAllocationOptPass : Pass<"memory-allocation-opt", "mlir::func::FuncOp"> {
def MemoryAllocationOpt : Pass<"memory-allocation-opt", "mlir::func::FuncOp"> {
let summary = "Convert stack to heap allocations and vice versa.";
let description = [{
Convert stack allocations to heap allocations and vice versa based on
@ -216,7 +216,7 @@ def MemoryAllocationOptPass : Pass<"memory-allocation-opt", "mlir::func::FuncOp"
let constructor = "::fir::createMemoryAllocationPass()";
}
def SimplifyRegionLitePass : Pass<"simplify-region-lite", "mlir::ModuleOp"> {
def SimplifyRegionLite : Pass<"simplify-region-lite", "mlir::ModuleOp"> {
let summary = "Region simplification";
let description = [{
Run region DCE and erase unreachable blocks in regions.
@ -224,7 +224,7 @@ def SimplifyRegionLitePass : Pass<"simplify-region-lite", "mlir::ModuleOp"> {
let constructor = "::fir::createSimplifyRegionLitePass()";
}
def AlgebraicSimplificationPass : Pass<"flang-algebraic-simplification"> {
def AlgebraicSimplification : Pass<"flang-algebraic-simplification"> {
let summary = "";
let description = [{
Run algebraic simplifications for Math/Complex/etc. dialect operations.

2
flang/include/flang/Tools/CLOptions.inc
View File

@ -183,7 +183,7 @@ inline void createDefaultFIROptimizerPassPipeline(
// convert control flow to CFG form
fir::addCfgConversionPass(pm);
pm.addPass(mlir::createConvertSCFToControlFlowPass());
pm.addPass(mlir::createConvertSCFToCFPass());
pm.addPass(mlir::createCanonicalizerPass(config));
pm.addPass(fir::createSimplifyRegionLitePass());

12
flang/lib/Optimizer/CodeGen/BoxedProcedure.cpp
View File

@ -6,10 +6,10 @@
//
//===----------------------------------------------------------------------===//
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "PassDetail.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/LowLevelIntrinsics.h"
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
@ -19,11 +19,6 @@
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/DialectConversion.h"
namespace fir {
#define GEN_PASS_DEF_BOXEDPROCEDUREPASS
#include "flang/Optimizer/CodeGen/CGPasses.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-procedure-pointer"
using namespace fir;
@ -174,8 +169,7 @@ private:
/// the frame pointer during execution. In LLVM IR, the frame pointer is
/// designated with the `nest` attribute. The thunk's address will then be used
/// as the call target instead of the original function's address directly.
class BoxedProcedurePass
: public fir::impl::BoxedProcedurePassBase<BoxedProcedurePass> {
class BoxedProcedurePass : public BoxedProcedurePassBase<BoxedProcedurePass> {
public:
BoxedProcedurePass() { options = {true}; }
BoxedProcedurePass(bool useThunks) { options = {useThunks}; }

15
flang/lib/Optimizer/CodeGen/CodeGen.cpp
View File

@ -11,8 +11,8 @@
//===----------------------------------------------------------------------===//
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "CGOps.h"
#include "PassDetail.h"
#include "flang/ISO_Fortran_binding.h"
#include "flang/Optimizer/Dialect/FIRAttr.h"
#include "flang/Optimizer/Dialect/FIROps.h"
@ -26,19 +26,12 @@
#include "mlir/Conversion/MathToLLVM/MathToLLVM.h"
#include "mlir/Conversion/MathToLibm/MathToLibm.h"
#include "mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Matchers.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Target/LLVMIR/ModuleTranslation.h"
#include "llvm/ADT/ArrayRef.h"
namespace fir {
#define GEN_PASS_DEF_FIRTOLLVMLOWERINGPASS
#include "flang/Optimizer/CodeGen/CGPasses.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-codegen"
// fir::LLVMTypeConverter for converting to LLVM IR dialect types.
@ -3050,8 +3043,7 @@ struct IsPresentOpConversion : public FIROpConversion<fir::IsPresentOp> {
auto ptr = adaptor.getOperands()[0];
if (isPresent.getVal().getType().isa<fir::BoxCharType>()) {
[[maybe_unused]] auto structTy =
ptr.getType().cast<mlir::LLVM::LLVMStructType>();
[[maybe_unused]] auto structTy = ptr.getType().cast<mlir::LLVM::LLVMStructType>();
assert(!structTy.isOpaque() && !structTy.getBody().empty());
ptr = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, ptr, 0);
@ -3288,8 +3280,7 @@ namespace {
///
/// This pass lowers all FIR dialect operations to LLVM IR dialect. An
/// MLIR pass is used to lower residual Std dialect to LLVM IR dialect.
class FIRToLLVMLowering
: public fir::impl::FIRToLLVMLoweringPassBase<FIRToLLVMLowering> {
class FIRToLLVMLowering : public fir::FIRToLLVMLoweringBase<FIRToLLVMLowering> {
public:
FIRToLLVMLowering() = default;
FIRToLLVMLowering(fir::FIRToLLVMPassOptions options) : options{options} {}

12
flang/lib/Optimizer/CodeGen/PreCGRewrite.cpp
View File

@ -10,9 +10,9 @@
//
//===----------------------------------------------------------------------===//
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "CGOps.h"
#include "PassDetail.h"
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
@ -21,11 +21,6 @@
#include "llvm/ADT/STLExtras.h"
#include "llvm/Support/Debug.h"
namespace fir {
#define GEN_PASS_DEF_CODEGENREWRITEPASS
#include "flang/Optimizer/CodeGen/CGPasses.h.inc"
} // namespace fir
//===----------------------------------------------------------------------===//
// Codegen rewrite: rewriting of subgraphs of ops
//===----------------------------------------------------------------------===//
@ -263,8 +258,7 @@ public:
}
};
class CodeGenRewrite
: public fir::impl::CodeGenRewritePassBase<CodeGenRewrite> {
class CodeGenRewrite : public fir::CodeGenRewriteBase<CodeGenRewrite> {
public:
void runOn(mlir::Operation *op, mlir::Region &region) {
auto &context = getContext();

11
flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
View File

@ -14,12 +14,12 @@
//
//===----------------------------------------------------------------------===//
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "PassDetail.h"
#include "Target.h"
#include "flang/Optimizer/Builder/Character.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/Todo.h"
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIROpsSupport.h"
@ -30,11 +30,6 @@
#include "llvm/ADT/TypeSwitch.h"
#include "llvm/Support/Debug.h"
namespace fir {
#define GEN_PASS_DEF_TARGETREWRITEPASS
#include "flang/Optimizer/CodeGen/CGPasses.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-target-rewrite"
namespace {
@ -71,7 +66,7 @@ struct FixupTy {
/// generation that traverses the FIR and modifies types and operations to a
/// form that is appropriate for the specific target. LLVM IR has specific
/// idioms that are used for distinct target processor and ABI combinations.
class TargetRewrite : public fir::impl::TargetRewritePassBase<TargetRewrite> {
class TargetRewrite : public fir::TargetRewriteBase<TargetRewrite> {
public:
TargetRewrite(const fir::TargetRewriteOptions &options) {
noCharacterConversion = options.noCharacterConversion;

16
flang/lib/Optimizer/Transforms/AbstractResult.cpp
View File

@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Builder/Todo.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
@ -18,12 +19,6 @@
#include "mlir/Transforms/Passes.h"
#include "llvm/ADT/TypeSwitch.h"
namespace fir {
#define GEN_PASS_DEF_ABSTRACTRESULTONFUNCOPTPASS
#define GEN_PASS_DEF_ABSTRACTRESULTONGLOBALOPTPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-abstract-result-opt"
namespace fir {
@ -252,8 +247,8 @@ public:
};
class AbstractResultOnFuncOpt
: public AbstractResultOptTemplate<
AbstractResultOnFuncOpt, fir::impl::AbstractResultOnFuncOptPassBase> {
: public AbstractResultOptTemplate<AbstractResultOnFuncOpt,
fir::AbstractResultOnFuncOptBase> {
public:
void runOnSpecificOperation(mlir::func::FuncOp func, bool shouldBoxResult,
mlir::RewritePatternSet &patterns,
@ -297,9 +292,8 @@ inline static bool containsFunctionTypeWithAbstractResult(mlir::Type type) {
}
class AbstractResultOnGlobalOpt
: public AbstractResultOptTemplate<
AbstractResultOnGlobalOpt,
fir::impl::AbstractResultOnGlobalOptPassBase> {
: public AbstractResultOptTemplate<AbstractResultOnGlobalOpt,
fir::AbstractResultOnGlobalOptBase> {
public:
void runOnSpecificOperation(fir::GlobalOp global, bool,
mlir::RewritePatternSet &,

8
flang/lib/Optimizer/Transforms/AffineDemotion.cpp
View File

@ -16,6 +16,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
@ -35,11 +36,6 @@
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/Debug.h"
namespace fir {
#define GEN_PASS_DEF_AFFINEDIALECTDEMOTIONPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-affine-demotion"
using namespace fir;
@ -141,7 +137,7 @@ public:
};
class AffineDialectDemotion
: public fir::impl::AffineDialectDemotionPassBase<AffineDialectDemotion> {
: public AffineDialectDemotionBase<AffineDialectDemotion> {
public:
void runOnOperation() override {
auto *context = &getContext();

8
flang/lib/Optimizer/Transforms/AffinePromotion.cpp
View File

@ -15,6 +15,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
@ -30,11 +31,6 @@
#include "llvm/ADT/Optional.h"
#include "llvm/Support/Debug.h"
namespace fir {
#define GEN_PASS_DEF_AFFINEDIALECTPROMOTIONPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-affine-promotion"
using namespace fir;
@ -583,7 +579,7 @@ public:
/// Promote fir.do_loop and fir.if to affine.for and affine.if, in the cases
/// where such a promotion is possible.
class AffineDialectPromotion
: public fir::impl::AffineDialectPromotionPassBase<AffineDialectPromotion> {
: public AffineDialectPromotionBase<AffineDialectPromotion> {
public:
void runOnOperation() override {

10
flang/lib/Optimizer/Transforms/AlgebraicSimplification.cpp
View File

@ -11,22 +11,16 @@
// the parameters of the patterns for Fortran programs.
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Dialect/Math/IR/Math.h"
#include "mlir/Dialect/Math/Transforms/Passes.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
namespace fir {
#define GEN_PASS_DEF_ALGEBRAICSIMPLIFICATIONPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
using namespace mlir;
namespace {
struct AlgebraicSimplification
: public fir::impl::AlgebraicSimplificationPassBase<
AlgebraicSimplification> {
: public fir::AlgebraicSimplificationBase<AlgebraicSimplification> {
AlgebraicSimplification(const GreedyRewriteConfig &rewriteConfig) {
config = rewriteConfig;
}

10
flang/lib/Optimizer/Transforms/AnnotateConstant.cpp
View File

@ -6,24 +6,18 @@
//
//===----------------------------------------------------------------------===//
// #include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/IR/BuiltinAttributes.h"
namespace fir {
#define GEN_PASS_DEF_ANNOTATECONSTANTOPERANDSPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-annotate-constant"
using namespace fir;
namespace {
struct AnnotateConstantOperands
: impl::AnnotateConstantOperandsPassBase<AnnotateConstantOperands> {
: AnnotateConstantOperandsBase<AnnotateConstantOperands> {
void runOnOperation() override {
auto *context = &getContext();
mlir::Dialect *firDialect = context->getLoadedDialect("fir");

8
flang/lib/Optimizer/Transforms/ArrayValueCopy.cpp
View File

@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Builder/Array.h"
#include "flang/Optimizer/Builder/BoxValue.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
@ -21,11 +22,6 @@
#include "mlir/Transforms/DialectConversion.h"
#include "llvm/Support/Debug.h"
namespace fir {
#define GEN_PASS_DEF_ARRAYVALUECOPYPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-array-value-copy"
using namespace fir;
@ -1330,7 +1326,7 @@ public:
};
class ArrayValueCopyConverter
: public fir::impl::ArrayValueCopyPassBase<ArrayValueCopyConverter> {
: public ArrayValueCopyBase<ArrayValueCopyConverter> {
public:
void runOnOperation() override {
auto func = getOperation();

9
flang/lib/Optimizer/Transforms/CharacterConversion.cpp
View File

@ -6,24 +6,19 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
#include "flang/Optimizer/Support/FIRContext.h"
#include "flang/Optimizer/Support/KindMapping.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/IR/Diagnostics.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/DialectConversion.h"
#include "llvm/Support/Debug.h"
namespace fir {
#define GEN_PASS_DEF_CHARACTERCONVERSIONPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-character-conversion"
namespace {
@ -100,7 +95,7 @@ public:
/// Rewrite the `fir.char_convert` op into a loop. This pass must be run only on
/// fir::CharConvertOp.
class CharacterConversion
: public fir::impl::CharacterConversionPassBase<CharacterConversion> {
: public fir::CharacterConversionBase<CharacterConversion> {
public:
void runOnOperation() override {
CharacterConversionOptions clOpts{useRuntimeCalls.getValue()};

10
flang/lib/Optimizer/Transforms/ExternalNameConversion.cpp
View File

@ -6,7 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Support/InternalNames.h"
#include "flang/Optimizer/Transforms/Passes.h"
@ -17,11 +17,6 @@
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/DialectConversion.h"
namespace fir {
#define GEN_PASS_DEF_EXTERNALNAMECONVERSIONPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
using namespace mlir;
//===----------------------------------------------------------------------===//
@ -122,8 +117,7 @@ public:
};
class ExternalNameConversionPass
: public fir::impl::ExternalNameConversionPassBase<
ExternalNameConversionPass> {
: public fir::ExternalNameConversionBase<ExternalNameConversionPass> {
public:
mlir::ModuleOp getModule() { return getOperation(); }
void runOnOperation() override;

9
flang/lib/Optimizer/Transforms/MemRefDataFlowOpt.cpp
View File

@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
@ -18,11 +19,6 @@
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallVector.h"
namespace fir {
#define GEN_PASS_DEF_MEMREFDATAFLOWOPTPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "fir-memref-dataflow-opt"
using namespace mlir;
@ -98,8 +94,7 @@ private:
mlir::DominanceInfo *domInfo;
};
class MemDataFlowOpt
: public fir::impl::MemRefDataFlowOptPassBase<MemDataFlowOpt> {
class MemDataFlowOpt : public fir::MemRefDataFlowOptBase<MemDataFlowOpt> {
public:
void runOnOperation() override {
mlir::func::FuncOp f = getOperation();

8
flang/lib/Optimizer/Transforms/MemoryAllocation.cpp
View File

@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
@ -17,11 +18,6 @@
#include "mlir/Transforms/Passes.h"
#include "llvm/ADT/TypeSwitch.h"
namespace fir {
#define GEN_PASS_DEF_MEMORYALLOCATIONOPTPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-memory-allocation-opt"
// Number of elements in an array does not determine where it is allocated.
@ -155,7 +151,7 @@ private:
/// 2. If a stack allocation is an array with a runtime evaluated size make
/// it a heap allocation.
class MemoryAllocationOpt
: public fir::impl::MemoryAllocationOptPassBase<MemoryAllocationOpt> {
: public fir::MemoryAllocationOptBase<MemoryAllocationOpt> {
public:
MemoryAllocationOpt() {
// Set options with default values. (See Passes.td.) Note that the

29
flang/lib/Optimizer/Transforms/PassDetail.h Normal file
View File

@ -0,0 +1,29 @@
//===- PassDetail.h - Optimizer Transforms Pass class details ---*- C++ -*-===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
#ifndef FORTRAN_OPTMIZER_TRANSFORMS_PASSDETAIL_H
#define FORTRAN_OPTMIZER_TRANSFORMS_PASSDETAIL_H
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/Math/IR/Math.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassRegistry.h"
namespace fir {
#define GEN_PASS_CLASSES
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#endif // FORTRAN_OPTMIZER_TRANSFORMS_PASSDETAIL_H

8
flang/lib/Optimizer/Transforms/RewriteLoop.cpp
View File

@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Transforms/Passes.h"
@ -16,11 +17,6 @@
#include "mlir/Transforms/DialectConversion.h"
#include "llvm/Support/CommandLine.h"
namespace fir {
#define GEN_PASS_DEF_CFGCONVERSIONPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
using namespace fir;
using namespace mlir;
@ -301,7 +297,7 @@ public:
};
/// Convert FIR structured control flow ops to CFG ops.
class CfgConversion : public fir::impl::CFGConversionPassBase<CfgConversion> {
class CfgConversion : public CFGConversionBase<CfgConversion> {
public:
void runOnOperation() override {
auto *context = &getContext();

9
flang/lib/Optimizer/Transforms/SimplifyIntrinsics.cpp
View File

@ -22,6 +22,7 @@
/// and small in size.
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Builder/BoxValue.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/Todo.h"
@ -29,7 +30,6 @@
#include "flang/Optimizer/Dialect/FIRType.h"
#include "flang/Optimizer/Support/FIRContext.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/IR/Matchers.h"
#include "mlir/IR/TypeUtilities.h"
#include "mlir/Pass/Pass.h"
@ -40,17 +40,12 @@
#include "llvm/Support/Debug.h"
#include "llvm/Support/raw_ostream.h"
namespace fir {
#define GEN_PASS_DEF_SIMPLIFYINTRINSICSPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
#define DEBUG_TYPE "flang-simplify-intrinsics"
namespace {
class SimplifyIntrinsicsPass
: public fir::impl::SimplifyIntrinsicsPassBase<SimplifyIntrinsicsPass> {
: public fir::SimplifyIntrinsicsBase<SimplifyIntrinsicsPass> {
using FunctionTypeGeneratorTy =
std::function<mlir::FunctionType(fir::FirOpBuilder &)>;
using FunctionBodyGeneratorTy =

8
flang/lib/Optimizer/Transforms/SimplifyRegionLite.cpp
View File

@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "PassDetail.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/IR/PatternMatch.h"
@ -14,15 +15,10 @@
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
#include "mlir/Transforms/RegionUtils.h"
namespace fir {
#define GEN_PASS_DEF_SIMPLIFYREGIONLITEPASS
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
namespace {
class SimplifyRegionLitePass
: public fir::impl::SimplifyRegionLitePassBase<SimplifyRegionLitePass> {
: public fir::SimplifyRegionLiteBase<SimplifyRegionLitePass> {
public:
void runOnOperation() override;
};

5
mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h
View File

@ -17,13 +17,12 @@ class LLVMTypeConverter;
class RewritePatternSet;
class Pass;
#define GEN_PASS_DECL_CONVERTAMDGPUTOROCDLPASS
#include "mlir/Conversion/Passes.h.inc"
void populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns,
amdgpu::Chipset chipset);
std::unique_ptr<Pass> createConvertAMDGPUToROCDLPass();
} // namespace mlir
#endif // MLIR_CONVERSION_AMDGPUTOROCDL_AMDGPUTOROCDL_H_

8
mlir/include/mlir/Conversion/AffineToStandard/AffineToStandard.h
View File

@ -23,9 +23,6 @@ class ValueRange;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTAFFINETOSTANDARDPASS
#include "mlir/Conversion/Passes.h.inc"
/// Collect a set of patterns to convert from the Affine dialect to the Standard
/// dialect, in particular convert structured affine control flow into CFG
/// branch-based control flow.
@ -43,6 +40,11 @@ Value lowerAffineLowerBound(AffineForOp op, OpBuilder &builder);
/// standard arithmetic operations.
Value lowerAffineUpperBound(AffineForOp op, OpBuilder &builder);
/// Lowers affine control flow operations (ForStmt, IfStmt and AffineApplyOp)
/// to equivalent lower-level constructs (flow of basic blocks and arithmetic
/// primitives).
std::unique_ptr<Pass> createLowerAffinePass();
} // namespace mlir
#endif // MLIR_CONVERSION_AFFINETOSTANDARD_AFFINETOSTANDARD_H

4
mlir/include/mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h
View File

@ -17,13 +17,11 @@ class LLVMTypeConverter;
class RewritePatternSet;
class Pass;
#define GEN_PASS_DECL_CONVERTARITHMETICTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
namespace arith {
void populateArithmeticToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);
std::unique_ptr<Pass> createConvertArithmeticToLLVMPass();
} // namespace arith
} // namespace mlir

7
mlir/include/mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h
View File

@ -18,17 +18,12 @@ class SPIRVTypeConverter;
class RewritePatternSet;
class Pass;
#define GEN_PASS_DECL_CONVERTARITHMETICTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
namespace arith {
void populateArithmeticToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
} // namespace arith
std::unique_ptr<OperationPass<>> createConvertArithmeticToSPIRVPass();
} // namespace arith
} // namespace mlir
#endif // MLIR_CONVERSION_ARITHMETICTOSPIRV_ARITHMETICTOSPIRV_H

13
mlir/include/mlir/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.h
View File

@ -9,19 +9,18 @@
#ifndef MLIR_CONVERSION_ARMNEON2DTOINTR_ARMNEON2DTOINTR_H_
#define MLIR_CONVERSION_ARMNEON2DTOINTR_ARMNEON2DTOINTR_H_
#include <memory>
#include "mlir/Pass/Pass.h"
namespace mlir {
class Pass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTARMNEON2DTOINTRPASS
#include "mlir/Conversion/Passes.h.inc"
/// Populates patterns for the lowering of Arm NEON 2D ops to intrinsics.
/// See createConvertArmNeon2dToIntrPass.
void populateConvertArmNeon2dToIntrPatterns(RewritePatternSet &patterns);
/// Creates a pass to lower Arm NEON 2D ops to intrinsics, i.e.
/// equivalent ops operating on flattened 1D vectors and mapping more
/// directly to the corresponding Arm NEON instruction.
std::unique_ptr<Pass> createConvertArmNeon2dToIntrPass();
} // namespace mlir
#endif // MLIR_CONVERSION_ARMNEON2DTOINTR_ARMNEON2DTOINTR_H_

7
mlir/include/mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h
View File

@ -21,8 +21,8 @@ class MLIRContext;
class TypeConverter;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTASYNCTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Create a pass to convert Async operations to the LLVM dialect.
std::unique_ptr<OperationPass<ModuleOp>> createConvertAsyncToLLVMPass();
/// Populates patterns for async structural type conversions.
///
@ -36,9 +36,6 @@ void populateAsyncStructuralTypeConversionsAndLegality(
TypeConverter &typeConverter, RewritePatternSet &patterns,
ConversionTarget &target);
/// Create a pass to convert Async operations to the LLVM dialect.
std::unique_ptr<OperationPass<ModuleOp>> createConvertAsyncToLLVMPass();
} // namespace mlir
#endif // MLIR_CONVERSION_ASYNCTOLLVM_ASYNCTOLLVM_H

4
mlir/include/mlir/Conversion/BufferizationToMemRef/BufferizationToMemRef.h
View File

@ -15,14 +15,12 @@ namespace mlir {
class Pass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTBUFFERIZATIONTOMEMREFPASS
#include "mlir/Conversion/Passes.h.inc"
/// Collect a set of patterns to convert memory-related operations from the
/// Bufferization dialect to the MemRef dialect.
void populateBufferizationToMemRefConversionPatterns(
RewritePatternSet &patterns);
std::unique_ptr<Pass> createBufferizationToMemRefPass();
} // namespace mlir
#endif // MLIR_CONVERSION_BUFFERIZATIONTOMEMREF_BUFFERIZATIONTOMEMREF_H

6
mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h
View File

@ -15,9 +15,6 @@ class LLVMTypeConverter;
class Pass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTCOMPLEXTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
class ComplexStructBuilder : public StructBuilder {
public:
/// Construct a helper for the given complex number value.
@ -41,6 +38,9 @@ public:
void populateComplexToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);
/// Create a pass to convert Complex operations to the LLVMIR dialect.
std::unique_ptr<Pass> createConvertComplexToLLVMPass();
} // namespace mlir
#endif // MLIR_CONVERSION_COMPLEXTOLLVM_COMPLEXTOLLVM_H_

3
mlir/include/mlir/Conversion/ComplexToLibm/ComplexToLibm.h
View File

@ -14,9 +14,6 @@ namespace mlir {
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTCOMPLEXTOLIBMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Populate the given list with patterns that convert from Complex to Libm
/// calls.
void populateComplexToLibmConversionPatterns(RewritePatternSet &patterns,

6
mlir/include/mlir/Conversion/ComplexToStandard/ComplexToStandard.h
View File

@ -14,12 +14,12 @@ namespace mlir {
class RewritePatternSet;
class Pass;
#define GEN_PASS_DECL_CONVERTCOMPLEXTOSTANDARDPASS
#include "mlir/Conversion/Passes.h.inc"
/// Populate the given list with patterns that convert from Complex to Standard.
void populateComplexToStandardConversionPatterns(RewritePatternSet &patterns);
/// Create a pass to convert Complex operations to the Standard dialect.
std::unique_ptr<Pass> createConvertComplexToStandardPass();
} // namespace mlir
#endif // MLIR_CONVERSION_COMPLEXTOSTANDARD_COMPLEXTOSTANDARD_H_

5
mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h
View File

@ -20,9 +20,6 @@ class LLVMTypeConverter;
class RewritePatternSet;
class Pass;
#define GEN_PASS_DECL_CONVERTCONTROLFLOWTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
namespace cf {
/// Collect the patterns to convert from the ControlFlow dialect to LLVM. The
/// conversion patterns capture the LLVMTypeConverter by reference meaning the
@ -30,6 +27,8 @@ namespace cf {
void populateControlFlowToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);
/// Creates a pass to convert the ControlFlow dialect into the LLVMIR dialect.
std::unique_ptr<Pass> createConvertControlFlowToLLVMPass();
} // namespace cf
} // namespace mlir

3
mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.h
View File

@ -18,9 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_CONVERTCONTROLFLOWTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert ControlFlow ops to SPIR-V ops.
std::unique_ptr<OperationPass<>> createConvertControlFlowToSPIRVPass();

4
mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h
View File

@ -10,7 +10,6 @@
#define MLIR_CONVERSION_FUNCTOLLVM_CONVERTFUNCTOLLVMPASS_H_
#include <memory>
#include <string>
namespace mlir {
class LowerToLLVMOptions;
@ -19,9 +18,6 @@ template <typename T>
class OperationPass;
class Pass;
#define GEN_PASS_DECL_CONVERTFUNCTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert the Func dialect into the LLVMIR dialect.
std::unique_ptr<OperationPass<ModuleOp>> createConvertFuncToLLVMPass();
std::unique_ptr<OperationPass<ModuleOp>>

3
mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRVPass.h
View File

@ -18,9 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_CONVERTFUNCTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert Func ops to SPIR-V ops.
std::unique_ptr<OperationPass<>> createConvertFuncToSPIRVPass();

19
mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h
View File

@ -38,22 +38,12 @@ namespace LLVM {
class LLVMDialect;
} // namespace LLVM
#define GEN_PASS_DECL_CONVERTGPUTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
using OwnedBlob = std::unique_ptr<std::vector<char>>;
using BlobGenerator =
std::function<OwnedBlob(const std::string &, Location, StringRef)>;
using LoweringCallback = std::function<std::unique_ptr<llvm::Module>(
Operation *, llvm::LLVMContext &, StringRef)>;
/// Collect a set of patterns to convert from the GPU dialect to LLVM and
/// populate converter for gpu types.
void populateGpuToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns,
StringRef gpuBinaryAnnotation = {},
bool kernelBarePtrCallConv = false);
/// Creates a pass to convert a GPU operations into a sequence of GPU runtime
/// calls.
///
@ -61,7 +51,14 @@ void populateGpuToLLVMConversionPatterns(LLVMTypeConverter &converter,
/// instead uses a small wrapper library that exports a stable and conveniently
/// typed ABI on top of GPU runtimes such as CUDA or ROCm (HIP).
std::unique_ptr<OperationPass<ModuleOp>>
createConvertGpuToLLVMPass(bool kernelBarePtrCallConv = false);
createGpuToLLVMConversionPass(bool kernelBarePtrCallConv = false);
/// Collect a set of patterns to convert from the GPU dialect to LLVM and
/// populate converter for gpu types.
void populateGpuToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns,
StringRef gpuBinaryAnnotation = {},
bool kernelBarePtrCallConv = false);
} // namespace mlir

5
mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h
View File

@ -25,9 +25,6 @@ class GPUModuleOp;
class MMAMatrixType;
} // namespace gpu
#define GEN_PASS_DECL_CONVERTGPUTONVVMPASS
#include "mlir/Conversion/Passes.h.inc"
LLVM::LLVMStructType convertMMAToLLVMType(gpu::MMAMatrixType type);
/// Configure target to convert from the GPU dialect to NVVM.
@ -44,7 +41,7 @@ void populateGpuWMMAToNVVMConversionPatterns(LLVMTypeConverter &converter,
/// Creates a pass that lowers GPU dialect operations to NVVM counterparts. The
/// index bitwidth used for the lowering of the device side index computations
/// is configurable.
std::unique_ptr<OperationPass<gpu::GPUModuleOp>> createConvertGpuToNVVMPass(
std::unique_ptr<OperationPass<gpu::GPUModuleOp>> createLowerGpuOpsToNVVMOpsPass(
unsigned indexBitwidth = kDeriveIndexBitwidthFromDataLayout);
} // namespace mlir

6
mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h
View File

@ -24,9 +24,6 @@ namespace gpu {
class GPUModuleOp;
} // namespace gpu
#define GEN_PASS_DECL_CONVERTGPUTOROCDLPASS
#include "mlir/Conversion/Passes.h.inc"
/// Collect a set of patterns to convert from the GPU dialect to ROCDL.
/// If `runtime` is Unknown, gpu.printf will not be lowered
/// The resulting pattern set should be run over a gpu.module op
@ -40,7 +37,8 @@ void configureGpuToROCDLConversionLegality(ConversionTarget &target);
/// Creates a pass that lowers GPU dialect operations to ROCDL counterparts. The
/// index bitwidth used for the lowering of the device side index computations
/// is configurable.
std::unique_ptr<OperationPass<gpu::GPUModuleOp>> createConvertGpuToROCDLPass(
std::unique_ptr<OperationPass<gpu::GPUModuleOp>>
createLowerGpuOpsToROCDLOpsPass(
const std::string &chipset = "gfx900",
unsigned indexBitwidth = kDeriveIndexBitwidthFromDataLayout,
bool useBarePtrCallConv = false,

3
mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h
View File

@ -21,9 +21,6 @@ class ModuleOp;
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTGPUTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert GPU kernel ops to corresponding SPIR-V ops. For a
/// gpu.func to be converted, it should have a spv.entry_point_abi attribute.
/// If `mapMemorySpace` is true, performs MemRef memory space to SPIR-V mapping

4
mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h
View File

@ -24,10 +24,6 @@ class ModuleOp;
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTVULKANLAUNCHFUNCTOVULKANCALLSPASS
#define GEN_PASS_DECL_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNCPASS
#include "mlir/Conversion/Passes.h.inc"
std::unique_ptr<OperationPass<ModuleOp>>
createConvertVulkanLaunchFuncToVulkanCallsPass();

3
mlir/include/mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h
View File

@ -18,9 +18,6 @@ template <typename T>
class OperationPass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTLINALGTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Populate the given list with patterns that convert from Linalg to LLVM.
void populateLinalgToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);

9
mlir/include/mlir/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.h
View File

@ -13,18 +13,13 @@
#ifndef MLIR_CONVERSION_LINALGTOSPIRV_LINALGTOSPIRVPASS_H
#define MLIR_CONVERSION_LINALGTOSPIRV_LINALGTOSPIRVPASS_H
#include <memory>
#include "mlir/Pass/Pass.h"
namespace mlir {
class ModuleOp;
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTLINALGTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates and returns a pass to convert Linalg ops to SPIR-V ops.
std::unique_ptr<OperationPass<ModuleOp>> createConvertLinalgToSPIRVPass();
std::unique_ptr<OperationPass<ModuleOp>> createLinalgToSPIRVPass();
} // namespace mlir

3
mlir/include/mlir/Conversion/LinalgToStandard/LinalgToStandard.h
View File

@ -17,9 +17,6 @@ class ModuleOp;
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTLINALGTOSTANDARDPASS
#include "mlir/Conversion/Passes.h.inc"
namespace linalg {
//===----------------------------------------------------------------------===//

6
mlir/include/mlir/Conversion/MathToFuncs/MathToFuncs.h
View File

@ -14,9 +14,9 @@
namespace mlir {
class Pass;
#define GEN_PASS_DECL_CONVERTMATHTOFUNCSPASS
#include "mlir/Conversion/Passes.h.inc"
// Pass to convert some Math operations into calls of functions
// containing software implementation of these operations.
std::unique_ptr<Pass> createConvertMathToFuncsPass();
} // namespace mlir
#endif // MLIR_CONVERSION_MATHTOFUNCS_MATHTOFUNCS_H

4
mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h
View File

@ -17,12 +17,10 @@ class LLVMTypeConverter;
class RewritePatternSet;
class Pass;
#define GEN_PASS_DECL_CONVERTMATHTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
void populateMathToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);
std::unique_ptr<Pass> createConvertMathToLLVMPass();
} // namespace mlir
#endif // MLIR_CONVERSION_MATHTOLLVM_MATHTOLLVM_H

3
mlir/include/mlir/Conversion/MathToLibm/MathToLibm.h
View File

@ -14,9 +14,6 @@ namespace mlir {
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTMATHTOLIBMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Populate the given list with patterns that convert from Math to Libm calls.
/// If log1pBenefit is present, use it instead of benefit for the Log1p op.
void populateMathToLibmConversionPatterns(

3
mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRVPass.h
View File

@ -18,9 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_CONVERTMATHTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert Math ops to SPIR-V ops.
std::unique_ptr<OperationPass<>> createConvertMathToSPIRVPass();

4
mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h
View File

@ -16,14 +16,12 @@ class Pass;
class LLVMTypeConverter;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTMEMREFTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Collect a set of patterns to convert memory-related operations from the
/// MemRef dialect to the LLVM dialect.
void populateMemRefToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);
std::unique_ptr<Pass> createMemRefToLLVMPass();
} // namespace mlir
#endif // MLIR_CONVERSION_MEMREFTOLLVM_MEMREFTOLLVM_H

4
mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.h
View File

@ -19,10 +19,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_MAPMEMREFSTORAGECLASSPASS
#define GEN_PASS_DECL_CONVERTMEMREFTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to map numeric MemRef memory spaces to symbolic SPIR-V
/// storage classes. The mapping is read from the command-line option.
std::unique_ptr<OperationPass<>> createMapMemRefStorageClassPass();

5
mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h
View File

@ -16,12 +16,11 @@ class LLVMTypeConverter;
class RewritePatternSet;
class Pass;
#define GEN_PASS_DECL_CONVERTNVGPUTONVVMPASS
#include "mlir/Conversion/Passes.h.inc"
void populateNVGPUToNVVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);
std::unique_ptr<Pass> createConvertNVGPUToNVVMPass();
} // namespace mlir
#endif // MLIR_CONVERSION_NVGPUTONVVM_NVGPUTONVVMPASS_H_

3
mlir/include/mlir/Conversion/OpenACCToLLVM/ConvertOpenACCToLLVM.h
View File

@ -18,9 +18,6 @@ template <typename T>
class OperationPass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTOPENACCTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
static constexpr unsigned kPtrBasePosInDataDescriptor = 0;
static constexpr unsigned kPtrPosInDataDescriptor = 1;
static constexpr unsigned kSizePosInDataDescriptor = 2;

3
mlir/include/mlir/Conversion/OpenACCToSCF/ConvertOpenACCToSCF.h
View File

@ -16,9 +16,6 @@ template <typename T>
class OperationPass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTOPENACCTOSCFPASS
#include "mlir/Conversion/Passes.h.inc"
/// Collect the patterns to convert from the OpenACC dialect to OpenACC with
/// SCF dialect.
void populateOpenACCToSCFConversionPatterns(RewritePatternSet &patterns);

3
mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h
View File

@ -19,9 +19,6 @@ template <typename T>
class OperationPass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTOOPENMPTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Configure dynamic conversion legality of regionless operations from OpenMP
/// to LLVM.
void configureOpenMPToLLVMConversionLegality(ConversionTarget &target,

3
mlir/include/mlir/Conversion/PDLToPDLInterp/PDLToPDLInterp.h
View File

@ -20,9 +20,6 @@ class ModuleOp;
template <typename OpT>
class OperationPass;
#define GEN_PASS_DECL_CONVERTPDLTOPDLINTERPPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates and returns a pass to convert PDL ops to PDL interpreter ops.
std::unique_ptr<OperationPass<ModuleOp>> createPDLToPDLInterpPass();

160
mlir/include/mlir/Conversion/Passes.td
View File

@ -15,7 +15,7 @@ include "mlir/Pass/PassBase.td"
// AffineToStandard
//===----------------------------------------------------------------------===//
def ConvertAffineToStandardPass : Pass<"lower-affine"> {
def ConvertAffineToStandard : Pass<"lower-affine"> {
let summary = "Lower Affine operations to a combination of Standard and SCF "
"operations";
let description = [{
@ -65,6 +65,7 @@ def ConvertAffineToStandardPass : Pass<"lower-affine"> {
if they do not depend on the loop iterator value or on the result of
`affine.apply`.
}];
let constructor = "mlir::createLowerAffinePass()";
let dependentDialects = [
"memref::MemRefDialect",
"scf::SCFDialect",
@ -76,11 +77,12 @@ def ConvertAffineToStandardPass : Pass<"lower-affine"> {
// AMDGPUToROCDL
//===----------------------------------------------------------------------===//
def ConvertAMDGPUToROCDLPass : Pass<"convert-amdgpu-to-rocdl"> {
def ConvertAMDGPUToROCDL : Pass<"convert-amdgpu-to-rocdl"> {
let summary = "Convert AMDGPU dialect to ROCDL dialect";
let description = [{
This pass converts supported AMDGPU ops to ROCDL dialect intrinsics.
}];
let constructor = "mlir::createConvertAMDGPUToROCDLPass()";
let dependentDialects = [
"LLVM::LLVMDialect",
"ROCDL::ROCDLDialect",
@ -94,11 +96,12 @@ def ConvertAMDGPUToROCDLPass : Pass<"convert-amdgpu-to-rocdl"> {
// ArithmeticToLLVM
//===----------------------------------------------------------------------===//
def ConvertArithmeticToLLVMPass : Pass<"convert-arith-to-llvm"> {
def ConvertArithmeticToLLVM : Pass<"convert-arith-to-llvm"> {
let summary = "Convert Arithmetic dialect to LLVM dialect";
let description = [{
This pass converts supported Arithmetic ops to LLVM dialect instructions.
}];
let constructor = "mlir::arith::createConvertArithmeticToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
let options = [
Option<"indexBitwidth", "index-bitwidth", "unsigned",
@ -111,9 +114,9 @@ def ConvertArithmeticToLLVMPass : Pass<"convert-arith-to-llvm"> {
// ArithmeticToSPIRV
//===----------------------------------------------------------------------===//
def ConvertArithmeticToSPIRVPass : Pass<"convert-arith-to-spirv"> {
def ConvertArithmeticToSPIRV : Pass<"convert-arith-to-spirv"> {
let summary = "Convert Arithmetic dialect to SPIR-V dialect";
let constructor = "mlir::createConvertArithmeticToSPIRVPass()";
let constructor = "mlir::arith::createConvertArithmeticToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
let options = [
Option<"emulateNon32BitScalarTypes", "emulate-non-32-bit-scalar-types",
@ -127,13 +130,9 @@ def ConvertArithmeticToSPIRVPass : Pass<"convert-arith-to-spirv"> {
// ArmNeon2dToIntr
//===----------------------------------------------------------------------===//
def ConvertArmNeon2dToIntrPass : Pass<"arm-neon-2d-to-intr"> {
def ConvertArmNeon2dToIntr : Pass<"arm-neon-2d-to-intr"> {
let summary = "Convert Arm NEON structured ops to intrinsics";
let description = [{
Lower Arm NEON 2D ops to intrinsics, i.e. equivalent ops operating on
flattened 1D vectors and mapping more directly to the corresponding Arm
NEON instruction.
}];
let constructor = "mlir::createConvertArmNeon2dToIntrPass()";
let dependentDialects = ["arm_neon::ArmNeonDialect", "vector::VectorDialect"];
}
@ -141,7 +140,7 @@ def ConvertArmNeon2dToIntrPass : Pass<"arm-neon-2d-to-intr"> {
// AsyncToLLVM
//===----------------------------------------------------------------------===//
def ConvertAsyncToLLVMPass : Pass<"convert-async-to-llvm", "ModuleOp"> {
def ConvertAsyncToLLVM : Pass<"convert-async-to-llvm", "ModuleOp"> {
let summary = "Convert the operations from the async dialect into the LLVM "
"dialect";
let description = [{
@ -159,7 +158,7 @@ def ConvertAsyncToLLVMPass : Pass<"convert-async-to-llvm", "ModuleOp"> {
// BufferizationToMemRef
//===----------------------------------------------------------------------===//
def ConvertBufferizationToMemRefPass : Pass<"convert-bufferization-to-memref"> {
def ConvertBufferizationToMemRef : Pass<"convert-bufferization-to-memref"> {
let summary = "Convert operations from the Bufferization dialect to the "
"MemRef dialect";
let description = [{
@ -186,6 +185,7 @@ def ConvertBufferizationToMemRefPass : Pass<"convert-bufferization-to-memref"> {
and hence does not resolve any memory leaks.
}];
let constructor = "mlir::createBufferizationToMemRefPass()";
let dependentDialects = ["arith::ArithmeticDialect", "memref::MemRefDialect"];
}
@ -193,8 +193,9 @@ def ConvertBufferizationToMemRefPass : Pass<"convert-bufferization-to-memref"> {
// ComplexToLLVM
//===----------------------------------------------------------------------===//
def ConvertComplexToLLVMPass : Pass<"convert-complex-to-llvm"> {
def ConvertComplexToLLVM : Pass<"convert-complex-to-llvm"> {
let summary = "Convert Complex dialect to LLVM dialect";
let constructor = "mlir::createConvertComplexToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
}
@ -202,7 +203,7 @@ def ConvertComplexToLLVMPass : Pass<"convert-complex-to-llvm"> {
// ComplexToLibm
//===----------------------------------------------------------------------===//
def ConvertComplexToLibmPass : Pass<"convert-complex-to-libm", "ModuleOp"> {
def ConvertComplexToLibm : Pass<"convert-complex-to-libm", "ModuleOp"> {
let summary = "Convert Complex dialect to libm calls";
let description = [{
This pass converts supported Complex ops to libm calls.
@ -217,8 +218,9 @@ def ConvertComplexToLibmPass : Pass<"convert-complex-to-libm", "ModuleOp"> {
// ComplexToStandard
//===----------------------------------------------------------------------===//
def ConvertComplexToStandardPass : Pass<"convert-complex-to-standard"> {
def ConvertComplexToStandard : Pass<"convert-complex-to-standard"> {
let summary = "Convert Complex dialect to standard dialect";
let constructor = "mlir::createConvertComplexToStandardPass()";
let dependentDialects = ["math::MathDialect"];
}
@ -226,7 +228,7 @@ def ConvertComplexToStandardPass : Pass<"convert-complex-to-standard"> {
// ControlFlowToLLVM
//===----------------------------------------------------------------------===//
def ConvertControlFlowToLLVMPass : Pass<"convert-cf-to-llvm", "ModuleOp"> {
def ConvertControlFlowToLLVM : Pass<"convert-cf-to-llvm", "ModuleOp"> {
let summary = "Convert ControlFlow operations to the LLVM dialect";
let description = [{
Convert ControlFlow operations into LLVM IR dialect operations.
@ -235,6 +237,7 @@ def ConvertControlFlowToLLVMPass : Pass<"convert-cf-to-llvm", "ModuleOp"> {
IR dialect operations, the pass will fail. Any LLVM IR operations or types
already present in the IR will be kept as is.
}];
let constructor = "mlir::cf::createConvertControlFlowToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
let options = [
Option<"indexBitwidth", "index-bitwidth", "unsigned",
@ -247,7 +250,7 @@ def ConvertControlFlowToLLVMPass : Pass<"convert-cf-to-llvm", "ModuleOp"> {
// ControlFlowToSPIRV
//===----------------------------------------------------------------------===//
def ConvertControlFlowToSPIRVPass : Pass<"convert-cf-to-spirv"> {
def ConvertControlFlowToSPIRV : Pass<"convert-cf-to-spirv"> {
let summary = "Convert ControlFlow dialect to SPIR-V dialect";
let constructor = "mlir::createConvertControlFlowToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
@ -263,7 +266,7 @@ def ConvertControlFlowToSPIRVPass : Pass<"convert-cf-to-spirv"> {
// FuncToLLVM
//===----------------------------------------------------------------------===//
def ConvertFuncToLLVMPass : Pass<"convert-func-to-llvm", "ModuleOp"> {
def ConvertFuncToLLVM : Pass<"convert-func-to-llvm", "ModuleOp"> {
let summary = "Convert from the Func dialect to the LLVM dialect";
let description = [{
Convert Func dialect operations into the LLVM IR dialect operations.
@ -308,7 +311,7 @@ def ConvertFuncToLLVMPass : Pass<"convert-func-to-llvm", "ModuleOp"> {
// FuncToSPIRV
//===----------------------------------------------------------------------===//
def ConvertFuncToSPIRVPass : Pass<"convert-func-to-spirv"> {
def ConvertFuncToSPIRV : Pass<"convert-func-to-spirv"> {
let summary = "Convert Func dialect to SPIR-V dialect";
let constructor = "mlir::createConvertFuncToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
@ -324,13 +327,13 @@ def ConvertFuncToSPIRVPass : Pass<"convert-func-to-spirv"> {
// GPUCommon
//===----------------------------------------------------------------------===//
def ConvertGpuToLLVMPass : Pass<"gpu-to-llvm", "ModuleOp"> {
def GpuToLLVMConversionPass : Pass<"gpu-to-llvm", "ModuleOp"> {
let summary = "Convert GPU dialect to LLVM dialect with GPU runtime calls";
let constructor = "mlir::createConvertGpuToLLVMPass()";
let constructor = "mlir::createGpuToLLVMConversionPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
}
def LowerHostCodeToLLVMPass : Pass<"lower-host-to-llvm", "ModuleOp"> {
def LowerHostCodeToLLVM : Pass<"lower-host-to-llvm", "ModuleOp"> {
let summary = "Lowers the host module code and `gpu.launch_func` to LLVM";
let constructor = "mlir::createLowerHostCodeToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
@ -340,9 +343,9 @@ def LowerHostCodeToLLVMPass : Pass<"lower-host-to-llvm", "ModuleOp"> {
// GPUToNVVM
//===----------------------------------------------------------------------===//
def ConvertGpuToNVVMPass : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {
def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {
let summary = "Generate NVVM operations for gpu operations";
let constructor = "mlir::createConvertGpuToNVVMPass()";
let constructor = "mlir::createLowerGpuOpsToNVVMOpsPass()";
let dependentDialects = [
"cf::ControlFlowDialect",
"memref::MemRefDialect",
@ -359,9 +362,9 @@ def ConvertGpuToNVVMPass : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {
// GPUToROCDL
//===----------------------------------------------------------------------===//
def ConvertGpuToROCDLPass : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
let summary = "Generate ROCDL operations for gpu operations";
let constructor = "mlir::createConvertGpuToROCDLPass()";
let constructor = "mlir::createLowerGpuOpsToROCDLOpsPass()";
let dependentDialects = ["ROCDL::ROCDLDialect"];
let options = [
Option<"chipset", "chipset", "std::string",
@ -389,7 +392,7 @@ def ConvertGpuToROCDLPass : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
// GPUToSPIRV
//===----------------------------------------------------------------------===//
def ConvertGPUToSPIRVPass : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
let summary = "Convert GPU dialect to SPIR-V dialect";
let description = [{
This pass converts supported GPU device ops to SPIR-V ops. It does not
@ -411,7 +414,7 @@ def ConvertGPUToSPIRVPass : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
// GPUToVulkan
//===----------------------------------------------------------------------===//
def ConvertGpuLaunchFuncToVulkanLaunchFuncPass
def ConvertGpuLaunchFuncToVulkanLaunchFunc
: Pass<"convert-gpu-launch-to-vulkan-launch", "ModuleOp"> {
let summary = "Convert gpu.launch_func to vulkanLaunch external call";
let description = [{
@ -421,7 +424,7 @@ def ConvertGpuLaunchFuncToVulkanLaunchFuncPass
let dependentDialects = ["spirv::SPIRVDialect"];
}
def ConvertVulkanLaunchFuncToVulkanCallsPass
def ConvertVulkanLaunchFuncToVulkanCalls
: Pass<"launch-func-to-vulkan", "ModuleOp"> {
let summary = "Convert vulkanLaunch external call to Vulkan runtime external "
"calls";
@ -436,7 +439,7 @@ def ConvertVulkanLaunchFuncToVulkanCallsPass
// LinalgToLLVM
//===----------------------------------------------------------------------===//
def ConvertLinalgToLLVMPass : Pass<"convert-linalg-to-llvm", "ModuleOp"> {
def ConvertLinalgToLLVM : Pass<"convert-linalg-to-llvm", "ModuleOp"> {
let summary = "Convert the operations from the linalg dialect into the LLVM "
"dialect";
let constructor = "mlir::createConvertLinalgToLLVMPass()";
@ -447,7 +450,7 @@ def ConvertLinalgToLLVMPass : Pass<"convert-linalg-to-llvm", "ModuleOp"> {
// LinalgToStandard
//===----------------------------------------------------------------------===//
def ConvertLinalgToStandardPass : Pass<"convert-linalg-to-std", "ModuleOp"> {
def ConvertLinalgToStandard : Pass<"convert-linalg-to-std", "ModuleOp"> {
let summary = "Convert the operations from the linalg dialect into the "
"Standard dialect";
let constructor = "mlir::createConvertLinalgToStandardPass()";
@ -458,13 +461,13 @@ def ConvertLinalgToStandardPass : Pass<"convert-linalg-to-std", "ModuleOp"> {
// LinalgToSPIRV
//===----------------------------------------------------------------------===//
def ConvertLinalgToSPIRVPass : Pass<"convert-linalg-to-spirv", "ModuleOp"> {
def ConvertLinalgToSPIRV : Pass<"convert-linalg-to-spirv", "ModuleOp"> {
let summary = "Convert Linalg dialect to SPIR-V dialect";
let description = [{
This pass converts supported Linalg ops to SPIR-V ops. It's quite
experimental and are expected to migrate to other proper conversions.
}];
let constructor = "mlir::createConvertLinalgToSPIRVPass()";
let constructor = "mlir::createLinalgToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
}
@ -472,7 +475,7 @@ def ConvertLinalgToSPIRVPass : Pass<"convert-linalg-to-spirv", "ModuleOp"> {
// MathToLibm
//===----------------------------------------------------------------------===//
def ConvertMathToLibmPass : Pass<"convert-math-to-libm", "ModuleOp"> {
def ConvertMathToLibm : Pass<"convert-math-to-libm", "ModuleOp"> {
let summary = "Convert Math dialect to libm calls";
let description = [{
This pass converts supported Math ops to libm calls.
@ -489,11 +492,12 @@ def ConvertMathToLibmPass : Pass<"convert-math-to-libm", "ModuleOp"> {
// MathToLLVM
//===----------------------------------------------------------------------===//
def ConvertMathToLLVMPass : Pass<"convert-math-to-llvm"> {
def ConvertMathToLLVM : Pass<"convert-math-to-llvm"> {
let summary = "Convert Math dialect to LLVM dialect";
let description = [{
This pass converts supported Math ops to LLVM dialect intrinsics.
}];
let constructor = "mlir::createConvertMathToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
}
@ -501,7 +505,7 @@ def ConvertMathToLLVMPass : Pass<"convert-math-to-llvm"> {
// MathToSPIRV
//===----------------------------------------------------------------------===//
def ConvertMathToSPIRVPass : Pass<"convert-math-to-spirv"> {
def ConvertMathToSPIRV : Pass<"convert-math-to-spirv"> {
let summary = "Convert Math dialect to SPIR-V dialect";
let constructor = "mlir::createConvertMathToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
@ -511,13 +515,14 @@ def ConvertMathToSPIRVPass : Pass<"convert-math-to-spirv"> {
// MathToFuncs
//===----------------------------------------------------------------------===//
def ConvertMathToFuncsPass : Pass<"convert-math-to-funcs", "ModuleOp"> {
def ConvertMathToFuncs : Pass<"convert-math-to-funcs", "ModuleOp"> {
let summary = "Convert Math operations to calls of outlined implementations.";
let description = [{
This pass converts supported Math ops to calls of compiler generated
functions implementing these operations in software.
The LLVM dialect is used for LinkonceODR linkage of the generated functions.
}];
let constructor = "mlir::createConvertMathToFuncsPass()";
let dependentDialects = [
"arith::ArithmeticDialect",
"cf::ControlFlowDialect",
@ -531,9 +536,10 @@ def ConvertMathToFuncsPass : Pass<"convert-math-to-funcs", "ModuleOp"> {
// MemRefToLLVM
//===----------------------------------------------------------------------===//
def ConvertMemRefToLLVMPass : Pass<"convert-memref-to-llvm", "ModuleOp"> {
def ConvertMemRefToLLVM : Pass<"convert-memref-to-llvm", "ModuleOp"> {
let summary = "Convert operations from the MemRef dialect to the LLVM "
"dialect";
let constructor = "mlir::createMemRefToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
let options = [
Option<"useAlignedAlloc", "use-aligned-alloc", "bool", /*default=*/"false",
@ -553,7 +559,7 @@ def ConvertMemRefToLLVMPass : Pass<"convert-memref-to-llvm", "ModuleOp"> {
// MemRefToSPIRV
//===----------------------------------------------------------------------===//
def MapMemRefStorageClassPass : Pass<"map-memref-spirv-storage-class"> {
def MapMemRefStorageClass : Pass<"map-memref-spirv-storage-class"> {
let summary = "Map numeric MemRef memory spaces to SPIR-V storage classes";
let constructor = "mlir::createMapMemRefStorageClassPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
@ -563,7 +569,7 @@ def MapMemRefStorageClassPass : Pass<"map-memref-spirv-storage-class"> {
];
}
def ConvertMemRefToSPIRVPass : Pass<"convert-memref-to-spirv"> {
def ConvertMemRefToSPIRV : Pass<"convert-memref-to-spirv"> {
let summary = "Convert MemRef dialect to SPIR-V dialect";
let constructor = "mlir::createConvertMemRefToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
@ -578,11 +584,12 @@ def ConvertMemRefToSPIRVPass : Pass<"convert-memref-to-spirv"> {
// NVGPUToNVVM
//===----------------------------------------------------------------------===//
def ConvertNVGPUToNVVMPass : Pass<"convert-nvgpu-to-nvvm"> {
def ConvertNVGPUToNVVM : Pass<"convert-nvgpu-to-nvvm"> {
let summary = "Convert NVGPU dialect to NVVM dialect";
let description = [{
This pass converts supported NVGPU ops to NVVM dialect intrinsics.
}];
let constructor = "mlir::createConvertNVGPUToNVVMPass()";
let dependentDialects = [
"NVVM::NVVMDialect",
];
@ -593,7 +600,7 @@ def ConvertNVGPUToNVVMPass : Pass<"convert-nvgpu-to-nvvm"> {
// OpenACCToSCF
//===----------------------------------------------------------------------===//
def ConvertOpenACCToSCFPass : Pass<"convert-openacc-to-scf", "ModuleOp"> {
def ConvertOpenACCToSCF : Pass<"convert-openacc-to-scf", "ModuleOp"> {
let summary = "Convert the OpenACC ops to OpenACC with SCF dialect";
let constructor = "mlir::createConvertOpenACCToSCFPass()";
let dependentDialects = ["scf::SCFDialect", "acc::OpenACCDialect"];
@ -603,7 +610,7 @@ def ConvertOpenACCToSCFPass : Pass<"convert-openacc-to-scf", "ModuleOp"> {
// OpenACCToLLVM
//===----------------------------------------------------------------------===//
def ConvertOpenACCToLLVMPass : Pass<"convert-openacc-to-llvm", "ModuleOp"> {
def ConvertOpenACCToLLVM : Pass<"convert-openacc-to-llvm", "ModuleOp"> {
let summary = "Convert the OpenACC ops to LLVM dialect";
let constructor = "mlir::createConvertOpenACCToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
@ -613,7 +620,7 @@ def ConvertOpenACCToLLVMPass : Pass<"convert-openacc-to-llvm", "ModuleOp"> {
// OpenMPToLLVM
//===----------------------------------------------------------------------===//
def ConvertOpenMPToLLVMPass : Pass<"convert-openmp-to-llvm", "ModuleOp"> {
def ConvertOpenMPToLLVM : Pass<"convert-openmp-to-llvm", "ModuleOp"> {
let summary = "Convert the OpenMP ops to OpenMP ops with LLVM dialect";
let constructor = "mlir::createConvertOpenMPToLLVMPass()";
let dependentDialects = ["LLVM::LLVMDialect"];
@ -623,7 +630,7 @@ def ConvertOpenMPToLLVMPass : Pass<"convert-openmp-to-llvm", "ModuleOp"> {
// PDLToPDLInterp
//===----------------------------------------------------------------------===//
def ConvertPDLToPDLInterpPass : Pass<"convert-pdl-to-pdl-interp", "ModuleOp"> {
def ConvertPDLToPDLInterp : Pass<"convert-pdl-to-pdl-interp", "ModuleOp"> {
let summary = "Convert PDL ops to PDL interpreter ops";
let constructor = "mlir::createPDLToPDLInterpPass()";
let dependentDialects = ["pdl_interp::PDLInterpDialect"];
@ -633,7 +640,7 @@ def ConvertPDLToPDLInterpPass : Pass<"convert-pdl-to-pdl-interp", "ModuleOp"> {
// ReconcileUnrealizedCasts
//===----------------------------------------------------------------------===//
def ReconcileUnrealizedCastsPass : Pass<"reconcile-unrealized-casts"> {
def ReconcileUnrealizedCasts : Pass<"reconcile-unrealized-casts"> {
let summary = "Simplify and eliminate unrealized conversion casts";
let description = [{
Eliminate `unrealized_conversion_cast` operations, commonly introduced by
@ -652,15 +659,17 @@ def ReconcileUnrealizedCastsPass : Pass<"reconcile-unrealized-casts"> {
and the producer operation is converted by another pass, each of which
produces an unrealized cast. This pass can be used to clean up the IR.
}];
let constructor = "mlir::createReconcileUnrealizedCastsPass()";
}
//===----------------------------------------------------------------------===//
// SCFToControlFlow
//===----------------------------------------------------------------------===//
def ConvertSCFToControlFlowPass : Pass<"convert-scf-to-cf"> {
def SCFToControlFlow : Pass<"convert-scf-to-cf"> {
let summary = "Convert SCF dialect to ControlFlow dialect, replacing structured"
" control flow with a CFG";
let constructor = "mlir::createConvertSCFToCFPass()";
let dependentDialects = ["cf::ControlFlowDialect"];
}
@ -668,7 +677,7 @@ def ConvertSCFToControlFlowPass : Pass<"convert-scf-to-cf"> {
// SCFToOpenMP
//===----------------------------------------------------------------------===//
def ConvertSCFToOpenMPPass : Pass<"convert-scf-to-openmp", "ModuleOp"> {
def ConvertSCFToOpenMP : Pass<"convert-scf-to-openmp", "ModuleOp"> {
let summary = "Convert SCF parallel loop to OpenMP parallel + workshare "
"constructs.";
let constructor = "mlir::createConvertSCFToOpenMPPass()";
@ -680,7 +689,7 @@ def ConvertSCFToOpenMPPass : Pass<"convert-scf-to-openmp", "ModuleOp"> {
// SCFToSPIRV
//===----------------------------------------------------------------------===//
def ConvertSCFToSPIRVPass : Pass<"convert-scf-to-spirv"> {
def SCFToSPIRV : Pass<"convert-scf-to-spirv"> {
let summary = "Convert SCF dialect to SPIR-V dialect.";
let description = [{
This pass converts SCF ops into SPIR-V structured control flow ops.
@ -697,7 +706,7 @@ def ConvertSCFToSPIRVPass : Pass<"convert-scf-to-spirv"> {
// SCFToGPU
//===----------------------------------------------------------------------===//
def ConvertAffineForToGPUPass
def ConvertAffineForToGPU
: InterfacePass<"convert-affine-for-to-gpu", "FunctionOpInterface"> {
let summary = "Convert top-level AffineFor Ops to GPU kernels";
let constructor = "mlir::createAffineForToGPUPass()";
@ -710,14 +719,9 @@ def ConvertAffineForToGPUPass
];
}
def ConvertParallelLoopToGpuPass : Pass<"convert-parallel-loops-to-gpu"> {
def ConvertParallelLoopToGpu : Pass<"convert-parallel-loops-to-gpu"> {
let summary = "Convert mapped scf.parallel ops to gpu launch operations";
let description = [{
Convert `scf.parallel` operations into a `gpu.launch` operation. The
mapping of loop dimensions to launch dimensions is derived from mapping
attributes. See `ParallelToGpuLaunchLowering::matchAndRewrite` for a
description of the used attributes.
}];
let constructor = "mlir::createParallelLoopToGpuPass()";
let dependentDialects = ["AffineDialect", "gpu::GPUDialect"];
}
@ -725,7 +729,7 @@ def ConvertParallelLoopToGpuPass : Pass<"convert-parallel-loops-to-gpu"> {
// ShapeToStandard
//===----------------------------------------------------------------------===//
def ConvertShapeToStandardPass : Pass<"convert-shape-to-std", "ModuleOp"> {
def ConvertShapeToStandard : Pass<"convert-shape-to-std", "ModuleOp"> {
let summary = "Convert operations from the shape dialect into the standard "
"dialect";
let constructor = "mlir::createConvertShapeToStandardPass()";
@ -734,7 +738,7 @@ def ConvertShapeToStandardPass : Pass<"convert-shape-to-std", "ModuleOp"> {
];
}
def ConvertShapeConstraintsPass : Pass<"convert-shape-constraints"> {
def ConvertShapeConstraints : Pass<"convert-shape-constraints"> {
let summary = "Convert shape constraint operations to the standard dialect";
let description = [{
This pass eliminates shape constraints from the program, converting them to
@ -745,6 +749,7 @@ def ConvertShapeConstraintsPass : Pass<"convert-shape-constraints"> {
can happen at a different part of the program than general shape
computation lowering.
}];
let constructor = "mlir::createConvertShapeConstraintsPass()";
let dependentDialects = ["cf::ControlFlowDialect", "scf::SCFDialect"];
}
@ -752,7 +757,7 @@ def ConvertShapeConstraintsPass : Pass<"convert-shape-constraints"> {
// SPIRVToLLVM
//===----------------------------------------------------------------------===//
def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
def ConvertSPIRVToLLVM : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
let summary = "Convert SPIR-V dialect to LLVM dialect";
let description = [{
See https://mlir.llvm.org/docs/SPIRVToLLVMDialectConversion/
@ -766,7 +771,7 @@ def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
// TensorToLinalg
//===----------------------------------------------------------------------===//
def ConvertTensorToLinalgPass : Pass<"convert-tensor-to-linalg", "ModuleOp"> {
def ConvertTensorToLinalg : Pass<"convert-tensor-to-linalg", "ModuleOp"> {
let summary = "Convert some Tensor dialect ops to Linalg dialect";
let constructor = "mlir::createConvertTensorToLinalgPass()";
let dependentDialects = [
@ -775,11 +780,12 @@ def ConvertTensorToLinalgPass : Pass<"convert-tensor-to-linalg", "ModuleOp"> {
];
}
//===----------------------------------------------------------------------===//
// TensorToSPIRV
//===----------------------------------------------------------------------===//
def ConvertTensorToSPIRVPass : Pass<"convert-tensor-to-spirv"> {
def ConvertTensorToSPIRV : Pass<"convert-tensor-to-spirv"> {
let summary = "Convert Tensor dialect to SPIR-V dialect";
let constructor = "mlir::createConvertTensorToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
@ -795,7 +801,7 @@ def ConvertTensorToSPIRVPass : Pass<"convert-tensor-to-spirv"> {
// TosaToArith
//===----------------------------------------------------------------------===//
def ConvertTosaToArithPass : Pass<"tosa-to-arith"> {
def TosaToArith : Pass<"tosa-to-arith"> {
let summary = "Lower TOSA to the Arith dialect";
let dependentDialects = [
"arith::ArithmeticDialect",
@ -814,52 +820,60 @@ def ConvertTosaToArithPass : Pass<"tosa-to-arith"> {
"bool", /*default=*/"false",
"Whether to prioritze lowering to 32-bit operations">
];
let constructor = "tosa::createTosaToArith()";
}
//===----------------------------------------------------------------------===//
// TosaToLinalg
//===----------------------------------------------------------------------===//
def ConvertTosaToLinalgPass
def TosaToLinalg
: InterfacePass<"tosa-to-linalg", "FunctionOpInterface"> {
let summary = "Lower TOSA to LinAlg on tensors";
let description = [{
Pass that converts TOSA operations to the equivalent operations using the
tensor operations in LinAlg.
}];
let constructor = "tosa::createTosaToLinalg()";
}
//===----------------------------------------------------------------------===//
// TosaToLinalgNamed
//===----------------------------------------------------------------------===//
def ConvertTosaToLinalgNamedPass
def TosaToLinalgNamed
: InterfacePass<"tosa-to-linalg-named", "FunctionOpInterface"> {
let summary = "Lower TOSA to LinAlg named operations";
let description = [{
Pass that converts TOSA operations to the equivalent operations using the
Linalg named operations.
}];
let constructor = "tosa::createTosaToLinalgNamed()";
}
//===----------------------------------------------------------------------===//
// TosaToSCF
//===----------------------------------------------------------------------===//
def ConvertTosaToSCFPass : Pass<"tosa-to-scf"> {
def TosaToSCF : Pass<"tosa-to-scf"> {
let summary = "Lower TOSA to the SCF dialect";
let dependentDialects = ["tensor::TensorDialect, scf::SCFDialect"];
let description = [{
Pass that converts TOSA's control flow operations to the equivalent SCF
operations.
}];
let constructor = "tosa::createTosaToSCF()";
}
//===----------------------------------------------------------------------===//
// TosaToTensor
//===----------------------------------------------------------------------===//
def ConvertTosaToTensorPass : Pass<"tosa-to-tensor"> {
def TosaToTensor : Pass<"tosa-to-tensor"> {
let summary = "Lower TOSA to the Tensor dialect";
let dependentDialects = [
"tensor::TensorDialect",
@ -868,13 +882,15 @@ def ConvertTosaToTensorPass : Pass<"tosa-to-tensor"> {
Pass that converts TOSA operations to the equivalent operations using the
operations in the Tensor dialect.
}];
let constructor = "tosa::createTosaToTensor()";
}
//===----------------------------------------------------------------------===//
// VectorToGPU
//===----------------------------------------------------------------------===//
def ConvertVectorToGPUPass : Pass<"convert-vector-to-gpu"> {
def ConvertVectorToGPU : Pass<"convert-vector-to-gpu"> {
let summary = "Lower the operations from the vector dialect into the GPU "
"dialect";
let constructor = "mlir::createConvertVectorToGPUPass()";
@ -893,7 +909,7 @@ def ConvertVectorToGPUPass : Pass<"convert-vector-to-gpu"> {
// VectorToSCF
//===----------------------------------------------------------------------===//
def ConvertVectorToSCFPass : Pass<"convert-vector-to-scf"> {
def ConvertVectorToSCF : Pass<"convert-vector-to-scf"> {
let summary = "Lower the operations from the vector dialect into the SCF "
"dialect";
let constructor = "mlir::createConvertVectorToSCFPass()";
@ -919,7 +935,7 @@ def ConvertVectorToSCFPass : Pass<"convert-vector-to-scf"> {
// VectorToLLVM
//===----------------------------------------------------------------------===//
def ConvertVectorToLLVMPass : Pass<"convert-vector-to-llvm", "ModuleOp"> {
def ConvertVectorToLLVM : Pass<"convert-vector-to-llvm", "ModuleOp"> {
let summary = "Lower the operations from the vector dialect into the LLVM "
"dialect";
let description = [{
@ -966,7 +982,7 @@ def ConvertVectorToLLVMPass : Pass<"convert-vector-to-llvm", "ModuleOp"> {
// VectorToSPIRV
//===----------------------------------------------------------------------===//
def ConvertVectorToSPIRVPass : Pass<"convert-vector-to-spirv"> {
def ConvertVectorToSPIRV : Pass<"convert-vector-to-spirv"> {
let summary = "Convert Vector dialect to SPIR-V dialect";
let constructor = "mlir::createConvertVectorToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];

6
mlir/include/mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h
View File

@ -15,13 +15,13 @@ namespace mlir {
class Pass;
class RewritePatternSet;
#define GEN_PASS_DECL_RECONCILEUNREALIZEDCASTSPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass that eliminates noop `unrealized_conversion_cast` operation
/// sequences.
std::unique_ptr<Pass> createReconcileUnrealizedCastsPass();
/// Populates `patterns` with rewrite patterns that eliminate noop
/// `unrealized_conversion_cast` operation sequences.
void populateReconcileUnrealizedCastsPatterns(RewritePatternSet &patterns);
} // namespace mlir
#endif // MLIR_CONVERSION_RECONCILEUNREALIZEDCASTS_RECONCILEUNREALIZEDCASTS_H_

7
mlir/include/mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h
View File

@ -15,13 +15,14 @@ namespace mlir {
class Pass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTSCFTOCONTROLFLOWPASS
#include "mlir/Conversion/Passes.h.inc"
/// Collect a set of patterns to convert SCF operations to CFG branch-based
/// operations within the ControlFlow dialect.
void populateSCFToControlFlowConversionPatterns(RewritePatternSet &patterns);
/// Creates a pass to convert SCF operations to CFG branch-based operation in
/// the ControlFlow dialect.
std::unique_ptr<Pass> createConvertSCFToCFPass();
} // namespace mlir
#endif // MLIR_CONVERSION_SCFTOCONTROLFLOW_SCFTOCONTROLFLOW_H_

10
mlir/include/mlir/Conversion/SCFToGPU/SCFToGPUPass.h
View File

@ -18,10 +18,6 @@ template <typename T>
class InterfacePass;
class Pass;
#define GEN_PASS_DECL_CONVERTAFFINEFORTOGPUPASS
#define GEN_PASS_DECL_CONVERTPARALLELLOOPTOGPUPASS
#include "mlir/Conversion/Passes.h.inc"
/// Create a pass that converts loop nests into GPU kernels. It considers
/// top-level affine.for operations as roots of loop nests and converts them to
/// the gpu.launch operations if possible.
@ -34,6 +30,12 @@ std::unique_ptr<InterfacePass<FunctionOpInterface>>
createAffineForToGPUPass(unsigned numBlockDims, unsigned numThreadDims);
std::unique_ptr<InterfacePass<FunctionOpInterface>> createAffineForToGPUPass();
/// Creates a pass that converts scf.parallel operations into a gpu.launch
/// operation. The mapping of loop dimensions to launch dimensions is derived
/// from mapping attributes. See ParallelToGpuLaunchLowering::matchAndRewrite
/// for a description of the used attributes.
std::unique_ptr<Pass> createParallelLoopToGpuPass();
} // namespace mlir
#endif // MLIR_CONVERSION_SCFTOGPU_SCFTOGPUPASS_H_

3
mlir/include/mlir/Conversion/SCFToOpenMP/SCFToOpenMP.h
View File

@ -16,9 +16,6 @@ class ModuleOp;
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTSCFTOOPENMPPASS
#include "mlir/Conversion/Passes.h.inc"
std::unique_ptr<OperationPass<ModuleOp>> createConvertSCFToOpenMPPass();
} // namespace mlir

3
mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRVPass.h
View File

@ -18,9 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_CONVERTSCFTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert SCF ops into SPIR-V ops.
std::unique_ptr<OperationPass<>> createConvertSCFToSPIRVPass();

4
mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h
View File

@ -20,10 +20,6 @@ class ModuleOp;
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_LOWERHOSTCODETOLLVMPASS
#define GEN_PASS_DECL_CONVERTSPIRVTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to emulate `gpu.launch_func` call in LLVM dialect and lower
/// the host module code to LLVM.
///

8
mlir/include/mlir/Conversion/ShapeToStandard/ShapeToStandard.h
View File

@ -19,16 +19,14 @@ template <typename T>
class OperationPass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTSHAPECONSTRAINTSPASS
#define GEN_PASS_DECL_CONVERTSHAPETOSTANDARDPASS
#include "mlir/Conversion/Passes.h.inc"
void populateShapeToStandardConversionPatterns(RewritePatternSet &patterns);
std::unique_ptr<OperationPass<ModuleOp>> createConvertShapeToStandardPass();
void populateConvertShapeConstraintsConversionPatterns(
RewritePatternSet &patterns);
std::unique_ptr<OperationPass<ModuleOp>> createConvertShapeToStandardPass();
std::unique_ptr<Pass> createConvertShapeConstraintsPass();
} // namespace mlir

3
mlir/include/mlir/Conversion/TensorToLinalg/TensorToLinalgPass.h
View File

@ -18,9 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_CONVERTTENSORTOLINALGPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert Tensor ops to Linalg ops.
std::unique_ptr<OperationPass<ModuleOp>> createConvertTensorToLinalgPass();

3
mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRVPass.h
View File

@ -18,9 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_CONVERTTENSORTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert Tensor ops to SPIR-V ops.
std::unique_ptr<OperationPass<>> createConvertTensorToSPIRVPass();

6
mlir/include/mlir/Conversion/TosaToArith/TosaToArith.h
View File

@ -16,12 +16,10 @@
#include "mlir/Pass/Pass.h"
namespace mlir {
#define GEN_PASS_DECL_CONVERTTOSATOARITHPASS
#include "mlir/Conversion/Passes.h.inc"
namespace tosa {
std::unique_ptr<Pass> createTosaToArith();
void populateTosaToArithConversionPatterns(RewritePatternSet *patterns);
void populateTosaRescaleToArithConversionPatterns(RewritePatternSet *patterns,

8
mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h
View File

@ -17,13 +17,11 @@
#include "mlir/Pass/Pass.h"
namespace mlir {
#define GEN_PASS_DECL_CONVERTTOSATOLINALGPASS
#define GEN_PASS_DECL_CONVERTTOSATOLINALGNAMEDPASS
#include "mlir/Conversion/Passes.h.inc"
namespace tosa {
std::unique_ptr<Pass> createTosaToLinalg();
std::unique_ptr<Pass> createTosaToLinalgNamed();
/// Populates passes to convert from TOSA to Linalg on buffers. At the end of
/// the pass, the function will only contain linalg ops or standard ops if the
/// pipeline succeeds. The option to disable decompositions is available for

6
mlir/include/mlir/Conversion/TosaToSCF/TosaToSCF.h
View File

@ -16,12 +16,10 @@
#include "mlir/Pass/Pass.h"
namespace mlir {
#define GEN_PASS_DECL_CONVERTTOSATOSCFPASS
#include "mlir/Conversion/Passes.h.inc"
namespace tosa {
std::unique_ptr<Pass> createTosaToSCF();
void populateTosaToSCFConversionPatterns(RewritePatternSet *patterns);
/// Populates passes to convert from TOSA to SCF.

6
mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h
View File

@ -16,12 +16,10 @@
#include "mlir/Pass/Pass.h"
namespace mlir {
#define GEN_PASS_DECL_CONVERTTOSATOTENSORPASS
#include "mlir/Conversion/Passes.h.inc"
namespace tosa {
std::unique_ptr<Pass> createTosaToTensor();
void populateTosaToTensorConversionPatterns(RewritePatternSet *patterns);
} // namespace tosa

3
mlir/include/mlir/Conversion/VectorToGPU/VectorToGPU.h
View File

@ -16,9 +16,6 @@ class MLIRContext;
class Pass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTVECTORTOGPUPASS
#include "mlir/Conversion/Passes.h.inc"
/// Patterns to transform vector ops into a canonical form to convert to MMA
/// matrix operations. If `useNvGpu` is true, then the patterns will populated
/// will prepare for conversion to `nvgpu` mma operations rather than the `gpu`

3
mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h
View File

@ -16,9 +16,6 @@ class ModuleOp;
template <typename T>
class OperationPass;
#define GEN_PASS_DECL_CONVERTVECTORTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
/// Options to control Vector to LLVM lowering.
///
/// This should kept in sync with VectorToLLVM options defined for the

3
mlir/include/mlir/Conversion/VectorToSCF/VectorToSCF.h
View File

@ -16,9 +16,6 @@ class MLIRContext;
class Pass;
class RewritePatternSet;
#define GEN_PASS_DECL_CONVERTVECTORTOSCFPASS
#include "mlir/Conversion/Passes.h.inc"
/// When lowering an N-d vector transfer op to an (N-1)-d vector transfer op,
/// a temporary buffer is created through which individual (N-1)-d vector are
/// staged. This pattern can be applied multiple time, until the transfer op

3
mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRVPass.h
View File

@ -18,9 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_CONVERTVECTORTOSPIRVPASS
#include "mlir/Conversion/Passes.h.inc"
/// Creates a pass to convert Vector Ops to SPIR-V ops.
std::unique_ptr<OperationPass<>> createConvertVectorToSPIRVPass();

16
mlir/include/mlir/Dialect/Affine/Passes.h
View File

@ -28,22 +28,6 @@ class AffineForOp;
/// producer-consumer and sibling fusion.
enum FusionMode { Greedy, ProducerConsumer, Sibling };
#define GEN_PASS_DECL_AFFINEDATACOPYGENERATIONPASS
#define GEN_PASS_DECL_AFFINELOOPFUSIONPASS
#define GEN_PASS_DECL_AFFINELOOPINVARIANTCODEMOTIONPASS
#define GEN_PASS_DECL_AFFINELOOPTILINGPASS
#define GEN_PASS_DECL_AFFINELOOPUNROLLPASS
#define GEN_PASS_DECL_AFFINELOOPUNROLLANDJAMPASS
#define GEN_PASS_DECL_AFFINEPIPELINEDATATRANSFERPASS
#define GEN_PASS_DECL_AFFINESCALARREPLACEMENTPASS
#define GEN_PASS_DECL_AFFINEVECTORIZEPASS
#define GEN_PASS_DECL_AFFINEPARALLELIZEPASS
#define GEN_PASS_DECL_AFFINELOOPNORMALIZEPASS
#define GEN_PASS_DECL_LOOPCOALESCINGPASS
#define GEN_PASS_DECL_SIMPLIFYAFFINESTRUCTURESPASS
#define GEN_PASS_DECL_AFFINEEXPANDINDEXOPSPASS
#include "mlir/Dialect/Affine/Passes.h.inc"
/// Creates a simplification pass for affine structures (maps and sets). In
/// addition, this pass also normalizes memrefs to have the trivial (identity)
/// layout map.

28
mlir/include/mlir/Dialect/Affine/Passes.td
View File

@ -15,7 +15,7 @@
include "mlir/Pass/PassBase.td"
def AffineDataCopyGenerationPass : Pass<"affine-data-copy-generate", "func::FuncOp"> {
def AffineDataCopyGeneration : Pass<"affine-data-copy-generate", "func::FuncOp"> {
let summary = "Generate explicit copying for affine memory operations";
let constructor = "mlir::createAffineDataCopyGenerationPass()";
let dependentDialects = ["memref::MemRefDialect"];
@ -43,7 +43,7 @@ def AffineDataCopyGenerationPass : Pass<"affine-data-copy-generate", "func::Func
];
}
def AffineLoopFusionPass : Pass<"affine-loop-fusion", "func::FuncOp"> {
def AffineLoopFusion : Pass<"affine-loop-fusion", "func::FuncOp"> {
let summary = "Fuse affine loop nests";
let description = [{
This pass performs fusion of loop nests using a slicing-based approach. It
@ -175,13 +175,13 @@ def AffineLoopFusionPass : Pass<"affine-loop-fusion", "func::FuncOp"> {
let dependentDialects = ["memref::MemRefDialect"];
}
def AffineLoopInvariantCodeMotionPass
def AffineLoopInvariantCodeMotion
: Pass<"affine-loop-invariant-code-motion", "func::FuncOp"> {
let summary = "Hoist loop invariant instructions outside of affine loops";
let constructor = "mlir::createAffineLoopInvariantCodeMotionPass()";
}
def AffineLoopTilingPass : Pass<"affine-loop-tile", "func::FuncOp"> {
def AffineLoopTiling : Pass<"affine-loop-tile", "func::FuncOp"> {
let summary = "Tile affine loop nests";
let constructor = "mlir::createLoopTilingPass()";
let options = [
@ -197,7 +197,7 @@ def AffineLoopTilingPass : Pass<"affine-loop-tile", "func::FuncOp"> {
];
}
def AffineLoopUnrollPass : Pass<"affine-loop-unroll", "func::FuncOp"> {
def AffineLoopUnroll : Pass<"affine-loop-unroll", "func::FuncOp"> {
let summary = "Unroll affine loops";
let constructor = "mlir::createLoopUnrollPass()";
let options = [
@ -217,7 +217,7 @@ def AffineLoopUnrollPass : Pass<"affine-loop-unroll", "func::FuncOp"> {
];
}
def AffineLoopUnrollAndJamPass : Pass<"affine-loop-unroll-jam", "func::FuncOp"> {
def AffineLoopUnrollAndJam : Pass<"affine-loop-unroll-jam", "func::FuncOp"> {
let summary = "Unroll and jam affine loops";
let constructor = "mlir::createLoopUnrollAndJamPass()";
let options = [
@ -227,7 +227,7 @@ def AffineLoopUnrollAndJamPass : Pass<"affine-loop-unroll-jam", "func::FuncOp">
];
}
def AffinePipelineDataTransferPass
def AffinePipelineDataTransfer
: Pass<"affine-pipeline-data-transfer", "func::FuncOp"> {
let summary = "Pipeline non-blocking data transfers between explicitly "
"managed levels of the memory hierarchy";
@ -296,7 +296,7 @@ def AffinePipelineDataTransferPass
let constructor = "mlir::createPipelineDataTransferPass()";
}
def AffineScalarReplacementPass : Pass<"affine-scalrep", "func::FuncOp"> {
def AffineScalarReplacement : Pass<"affine-scalrep", "func::FuncOp"> {
let summary = "Replace affine memref acceses by scalars by forwarding stores "
"to loads and eliminating redundant loads";
let description = [{
@ -342,7 +342,7 @@ def AffineScalarReplacementPass : Pass<"affine-scalrep", "func::FuncOp"> {
let constructor = "mlir::createAffineScalarReplacementPass()";
}
def AffineVectorizePass : Pass<"affine-super-vectorize", "func::FuncOp"> {
def AffineVectorize : Pass<"affine-super-vectorize", "func::FuncOp"> {
let summary = "Vectorize to a target independent n-D vector abstraction";
let constructor = "mlir::createSuperVectorizePass()";
let dependentDialects = ["vector::VectorDialect"];
@ -366,7 +366,7 @@ def AffineVectorizePass : Pass<"affine-super-vectorize", "func::FuncOp"> {
];
}
def AffineParallelizePass : Pass<"affine-parallelize", "func::FuncOp"> {
def AffineParallelize : Pass<"affine-parallelize", "func::FuncOp"> {
let summary = "Convert affine.for ops into 1-D affine.parallel";
let constructor = "mlir::createAffineParallelizePass()";
let options = [
@ -379,25 +379,25 @@ def AffineParallelizePass : Pass<"affine-parallelize", "func::FuncOp"> {
];
}
def AffineLoopNormalizePass : Pass<"affine-loop-normalize", "func::FuncOp"> {
def AffineLoopNormalize : Pass<"affine-loop-normalize", "func::FuncOp"> {
let summary = "Apply normalization transformations to affine loop-like ops";
let constructor = "mlir::createAffineLoopNormalizePass()";
}
def LoopCoalescingPass : Pass<"affine-loop-coalescing", "func::FuncOp"> {
def LoopCoalescing : Pass<"affine-loop-coalescing", "func::FuncOp"> {
let summary = "Coalesce nested loops with independent bounds into a single "
"loop";
let constructor = "mlir::createLoopCoalescingPass()";
let dependentDialects = ["arith::ArithmeticDialect"];
}
def SimplifyAffineStructuresPass : Pass<"affine-simplify-structures", "func::FuncOp"> {
def SimplifyAffineStructures : Pass<"affine-simplify-structures", "func::FuncOp"> {
let summary = "Simplify affine expressions in maps/sets and normalize "
"memrefs";
let constructor = "mlir::createSimplifyAffineStructuresPass()";
}
def AffineExpandIndexOpsPass : Pass<"affine-expand-index-ops"> {
def AffineExpandIndexOps : Pass<"affine-expand-index-ops"> {
let summary = "Lower affine operations operating on indices into more fundamental operations";
let constructor = "mlir::createAffineExpandIndexOpsPass()";
}

5
mlir/include/mlir/Dialect/Arithmetic/Transforms/Passes.h
View File

@ -14,11 +14,6 @@
namespace mlir {
namespace arith {
#define GEN_PASS_DECL_ARITHMETICBUFFERIZEPASS
#define GEN_PASS_DECL_ARITHMETICEXPANDOPSPASS
#define GEN_PASS_DECL_ARITHMETICUNSIGNEDWHENEQUIVALENTPASS
#include "mlir/Dialect/Arithmetic/Transforms/Passes.h.inc"
/// Create a pass to bufferize Arithmetic ops.
std::unique_ptr<Pass> createArithmeticBufferizePass();

6
mlir/include/mlir/Dialect/Arithmetic/Transforms/Passes.td
View File

@ -11,7 +11,7 @@
include "mlir/Pass/PassBase.td"
def ArithmeticBufferizePass : Pass<"arith-bufferize", "ModuleOp"> {
def ArithmeticBufferize : Pass<"arith-bufferize", "ModuleOp"> {
let summary = "Bufferize Arithmetic dialect ops.";
let description = [{
This pass bufferizes arith dialect ops.
@ -28,12 +28,12 @@ def ArithmeticBufferizePass : Pass<"arith-bufferize", "ModuleOp"> {
];
}
def ArithmeticExpandOpsPass : Pass<"arith-expand"> {
def ArithmeticExpandOps : Pass<"arith-expand"> {
let summary = "Legalize Arithmetic ops to be convertible to LLVM.";
let constructor = "mlir::arith::createArithmeticExpandOpsPass()";
}
def ArithmeticUnsignedWhenEquivalentPass : Pass<"arith-unsigned-when-equivalent"> {
def ArithmeticUnsignedWhenEquivalent : Pass<"arith-unsigned-when-equivalent"> {
let summary = "Replace signed ops with unsigned ones where they are proven equivalent";
let description = [{
Replace signed ops with their unsigned equivalents when integer range analysis

7
mlir/include/mlir/Dialect/Async/Passes.h
View File

@ -18,13 +18,6 @@
namespace mlir {
class ModuleOp;
#define GEN_PASS_DECL_ASYNCPARALLELFORPASS
#define GEN_PASS_DECL_ASYNCTOASYNCRUNTIMEPASS
#define GEN_PASS_DECL_ASYNCRUNTIMEREFCOUNTINGPASS
#define GEN_PASS_DECL_ASYNCRUNTIMEREFCOUNTINGOPTPASS
#define GEN_PASS_DECL_ASYNCRUNTIMEPOLICYBASEDREFCOUNTINGPASS
#include "mlir/Dialect/Async/Passes.h.inc"
std::unique_ptr<Pass> createAsyncParallelForPass();
std::unique_ptr<Pass> createAsyncParallelForPass(bool asyncDispatch,

10
mlir/include/mlir/Dialect/Async/Passes.td
View File

@ -11,7 +11,7 @@
include "mlir/Pass/PassBase.td"
def AsyncParallelForPass : Pass<"async-parallel-for", "ModuleOp"> {
def AsyncParallelFor : Pass<"async-parallel-for", "ModuleOp"> {
let summary = "Convert scf.parallel operations to multiple async compute ops "
"executed concurrently for non-overlapping iteration ranges";
let constructor = "mlir::createAsyncParallelForPass()";
@ -40,7 +40,7 @@ def AsyncParallelForPass : Pass<"async-parallel-for", "ModuleOp"> {
];
}
def AsyncToAsyncRuntimePass : Pass<"async-to-async-runtime", "ModuleOp"> {
def AsyncToAsyncRuntime : Pass<"async-to-async-runtime", "ModuleOp"> {
let summary = "Lower high level async operations (e.g. async.execute) to the"
"explicit async.runtime and async.coro operations";
let constructor = "mlir::createAsyncToAsyncRuntimePass()";
@ -54,7 +54,7 @@ def AsyncToAsyncRuntimePass : Pass<"async-to-async-runtime", "ModuleOp"> {
let dependentDialects = ["async::AsyncDialect"];
}
def AsyncRuntimeRefCountingPass : Pass<"async-runtime-ref-counting"> {
def AsyncRuntimeRefCounting : Pass<"async-runtime-ref-counting"> {
let summary = "Automatic reference counting for Async runtime operations";
let description = [{
This pass works at the async runtime abtraction level, after all
@ -71,14 +71,14 @@ def AsyncRuntimeRefCountingPass : Pass<"async-runtime-ref-counting"> {
let dependentDialects = ["async::AsyncDialect"];
}
def AsyncRuntimeRefCountingOptPass : Pass<"async-runtime-ref-counting-opt"> {
def AsyncRuntimeRefCountingOpt : Pass<"async-runtime-ref-counting-opt"> {
let summary = "Optimize automatic reference counting operations for the"
"Async runtime by removing redundant operations";
let constructor = "mlir::createAsyncRuntimeRefCountingOptPass()";
let dependentDialects = ["async::AsyncDialect"];
}
def AsyncRuntimePolicyBasedRefCountingPass
def AsyncRuntimePolicyBasedRefCounting
: Pass<"async-runtime-policy-based-ref-counting"> {
let summary = "Policy based reference counting for Async runtime operations";
let description = [{

13
mlir/include/mlir/Dialect/Bufferization/Transforms/Passes.h
View File

@ -17,19 +17,6 @@ struct OneShotBufferizationOptions;
// Passes
//===----------------------------------------------------------------------===//
#define GEN_PASS_DECL_BUFFERDEALLOCATIONPASS
#define GEN_PASS_DECL_BUFFERHOISTINGPASS
#define GEN_PASS_DECL_BUFFERLOOPHOISTINGPASS
#define GEN_PASS_DECL_BUFFERRESULTSTOOUTPARAMSPASS
#define GEN_PASS_DECL_FINALIZINGBUFFERIZEPASS
#define GEN_PASS_DECL_BUFFERIZATIONBUFFERIZEPASS
#define GEN_PASS_DECL_DROPEQUIVALENTBUFFERRESULTSPASS
#define GEN_PASS_DECL_ONESHOTBUFFERIZEPASS
#define GEN_PASS_DECL_PROMOTEBUFFERSTOSTACKPASS
#define GEN_PASS_DECL_TENSORCOPYINSERTIONPASS
#define GEN_PASS_DECL_ALLOCTENSORELIMINATIONPASS
#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc"
/// Creates an instance of the BufferDeallocation pass to free all allocated
/// buffers.
std::unique_ptr<Pass> createBufferDeallocationPass();

22
mlir/include/mlir/Dialect/Bufferization/Transforms/Passes.td
View File

@ -11,7 +11,7 @@
include "mlir/Pass/PassBase.td"
def BufferDeallocationPass : Pass<"buffer-deallocation", "func::FuncOp"> {
def BufferDeallocation : Pass<"buffer-deallocation", "func::FuncOp"> {
let summary = "Adds all required dealloc operations for all allocations in "
"the input program";
let description = [{
@ -88,7 +88,7 @@ def BufferDeallocationPass : Pass<"buffer-deallocation", "func::FuncOp"> {
let constructor = "mlir::bufferization::createBufferDeallocationPass()";
}
def BufferHoistingPass : Pass<"buffer-hoisting", "func::FuncOp"> {
def BufferHoisting : Pass<"buffer-hoisting", "func::FuncOp"> {
let summary = "Optimizes placement of allocation operations by moving them "
"into common dominators and out of nested regions";
let description = [{
@ -98,7 +98,7 @@ def BufferHoistingPass : Pass<"buffer-hoisting", "func::FuncOp"> {
let constructor = "mlir::bufferization::createBufferHoistingPass()";
}
def BufferLoopHoistingPass : Pass<"buffer-loop-hoisting", "func::FuncOp"> {
def BufferLoopHoisting : Pass<"buffer-loop-hoisting", "func::FuncOp"> {
let summary = "Optimizes placement of allocation operations by moving them "
"out of loop nests";
let description = [{
@ -108,7 +108,7 @@ def BufferLoopHoistingPass : Pass<"buffer-loop-hoisting", "func::FuncOp"> {
let constructor = "mlir::bufferization::createBufferLoopHoistingPass()";
}
def BufferResultsToOutParamsPass : Pass<"buffer-results-to-out-params", "ModuleOp"> {
def BufferResultsToOutParams : Pass<"buffer-results-to-out-params", "ModuleOp"> {
let summary = "Converts memref-typed function results to out-params";
let description = [{
Some calling conventions prefer to pass output memrefs as "out params". The
@ -133,7 +133,7 @@ def BufferResultsToOutParamsPass : Pass<"buffer-results-to-out-params", "ModuleO
let dependentDialects = ["memref::MemRefDialect"];
}
def FinalizingBufferizePass : Pass<"finalizing-bufferize", "func::FuncOp"> {
def FinalizingBufferize : Pass<"finalizing-bufferize", "func::FuncOp"> {
let summary = "Finalize a partial bufferization";
let description = [{
A bufferize pass that finalizes a partial bufferization by removing
@ -149,12 +149,12 @@ def FinalizingBufferizePass : Pass<"finalizing-bufferize", "func::FuncOp"> {
let constructor = "mlir::bufferization::createFinalizingBufferizePass()";
}
def BufferizationBufferizePass : Pass<"bufferization-bufferize", "func::FuncOp"> {
def BufferizationBufferize : Pass<"bufferization-bufferize", "func::FuncOp"> {
let summary = "Bufferize the `bufferization` dialect";
let constructor = "mlir::bufferization::createBufferizationBufferizePass()";
}
def DropEquivalentBufferResultsPass : Pass<"drop-equivalent-buffer-results", "ModuleOp"> {
def DropEquivalentBufferResults : Pass<"drop-equivalent-buffer-results", "ModuleOp"> {
let summary = "Remove MemRef return values that are equivalent to a bbArg";
let description = [{
This pass removes MemRef return values from functions if they are equivalent
@ -168,7 +168,7 @@ def DropEquivalentBufferResultsPass : Pass<"drop-equivalent-buffer-results", "Mo
let dependentDialects = ["memref::MemRefDialect"];
}
def OneShotBufferizePass : Pass<"one-shot-bufferize", "ModuleOp"> {
def OneShotBufferize : Pass<"one-shot-bufferize", "ModuleOp"> {
let summary = "One-Shot Bufferize";
let description = [{
This pass bufferizes all ops that implement `BufferizableOpInterface`. It
@ -303,7 +303,7 @@ def OneShotBufferizePass : Pass<"one-shot-bufferize", "ModuleOp"> {
let constructor = "mlir::bufferization::createOneShotBufferizePass()";
}
def PromoteBuffersToStackPass : Pass<"promote-buffers-to-stack", "func::FuncOp"> {
def PromoteBuffersToStack : Pass<"promote-buffers-to-stack", "func::FuncOp"> {
let summary = "Promotes heap-based allocations to automatically managed "
"stack-based allocations";
let description = [{
@ -324,7 +324,7 @@ def PromoteBuffersToStackPass : Pass<"promote-buffers-to-stack", "func::FuncOp">
];
}
def TensorCopyInsertionPass : Pass<"tensor-copy-insertion"> {
def TensorCopyInsertion : Pass<"tensor-copy-insertion"> {
let summary = "Make all tensor IR inplaceable by inserting copies";
let description = [{
This pass runs One-Shot Analysis and inserts copies for all OpOperands that
@ -355,7 +355,7 @@ def TensorCopyInsertionPass : Pass<"tensor-copy-insertion"> {
let constructor = "mlir::bufferization::createTensorCopyInsertionPass()";
}
def AllocTensorEliminationPass : Pass<"eliminate-alloc-tensors"> {
def AllocTensorElimination : Pass<"eliminate-alloc-tensors"> {
let summary = "Try to eliminate all alloc_tensor ops.";
let description = [{
This pass tries to eliminate all insert_slice op-anchored alloc_tensor ops.

4
mlir/include/mlir/Dialect/Func/Transforms/Passes.h
View File

@ -25,10 +25,6 @@ class BufferizeTypeConverter;
class RewritePatternSet;
namespace func {
#define GEN_PASS_DECL_FUNCBUFFERIZEPASS
#include "mlir/Dialect/Func/Transforms/Passes.h.inc"
/// Creates an instance of func bufferization pass.
std::unique_ptr<Pass> createFuncBufferizePass();

2
mlir/include/mlir/Dialect/Func/Transforms/Passes.td
View File

@ -11,7 +11,7 @@
include "mlir/Pass/PassBase.td"
def FuncBufferizePass : Pass<"func-bufferize", "ModuleOp"> {
def FuncBufferize : Pass<"func-bufferize", "ModuleOp"> {
let summary = "Bufferize func/call/return ops";
let description = [{
A bufferize pass that bufferizes func.func and func.call ops.

6
mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
View File

@ -27,12 +27,6 @@ namespace func {
class FuncOp;
} // namespace func
#define GEN_PASS_DECL_GPULAUNCHSINKINDEXCOMPUTATIONSPASS
#define GEN_PASS_DECL_GPUKERNELOUTLININGPASS
#define GEN_PASS_DECL_GPUASYNCREGIONPASS
#define GEN_PASS_DECL_GPUMAPPARALLELLOOPSPASS
#include "mlir/Dialect/GPU/Transforms/Passes.h.inc"
/// Pass that moves ops which are likely an index computation into gpu.launch
/// body.
std::unique_ptr<Pass> createGpuLauchSinkIndexComputationsPass();

4
mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
View File

@ -11,13 +11,13 @@
include "mlir/Pass/PassBase.td"
def GpuLaunchSinkIndexComputationsPass : Pass<"gpu-launch-sink-index-computations"> {
def GpuLaunchSinkIndexComputations : Pass<"gpu-launch-sink-index-computations"> {
let summary = "Sink index computations into gpu.launch body";
let constructor = "mlir::createGpuLauchSinkIndexComputationsPass()";
let dependentDialects = ["mlir::gpu::GPUDialect"];
}
def GpuKernelOutliningPass : Pass<"gpu-kernel-outlining", "ModuleOp"> {
def GpuKernelOutlining : Pass<"gpu-kernel-outlining", "ModuleOp"> {
let summary = "Outline gpu.launch bodies to kernel functions";
let constructor = "mlir::createGpuKernelOutliningPass()";
let dependentDialects = ["mlir::DLTIDialect"];

3
mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h
View File

@ -17,9 +17,6 @@ class Pass;
namespace LLVM {
#define GEN_PASS_DECL_LLVMLEGALIZEFOREXPORTPASS
#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
/// Make argument-taking successors of each block distinct. PHI nodes in LLVM
/// IR use the predecessor ID to identify which value to take. They do not
/// support different values coming from the same predecessor. If a block has

3
mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h
View File

@ -16,9 +16,6 @@ class Pass;
namespace NVVM {
#define GEN_PASS_DECL_NVVMOPTIMIZEFORTARGETPASS
#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
/// Creates a pass that optimizes LLVM IR for the NVVM target.
std::unique_ptr<Pass> createOptimizeForTargetPass();

6
mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td
View File

@ -11,12 +11,12 @@
include "mlir/Pass/PassBase.td"
def LLVMLegalizeForExportPass : Pass<"llvm-legalize-for-export"> {
def LLVMLegalizeForExport : Pass<"llvm-legalize-for-export"> {
let summary = "Legalize LLVM dialect to be convertible to LLVM IR";
let constructor = "::mlir::LLVM::createLegalizeForExportPass()";
}
def LLVMRequestCWrappersPass
def LLVMRequestCWrappers
: Pass<"llvm-request-c-wrappers", "::mlir::func::FuncOp"> {
let summary = "Request C wrapper emission for all functions";
let description = [{
@ -29,7 +29,7 @@ def LLVMRequestCWrappersPass
let constructor = "::mlir::LLVM::createRequestCWrappersPass()";
}
def NVVMOptimizeForTargetPass : Pass<"llvm-optimize-for-nvvm-target"> {
def NVVMOptimizeForTarget : Pass<"llvm-optimize-for-nvvm-target"> {
let summary = "Optimize NVVM IR";
let constructor = "::mlir::NVVM::createOptimizeForTargetPass()";
}

4
mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h
View File

@ -15,10 +15,6 @@ namespace mlir {
class Pass;
namespace LLVM {
#define GEN_PASS_DECL_LLVMREQUESTCWRAPPERSPASS
#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
std::unique_ptr<Pass> createRequestCWrappersPass();
} // namespace LLVM
} // namespace mlir

24
mlir/include/mlir/Dialect/Linalg/Passes.h
View File

@ -26,30 +26,6 @@ namespace bufferization {
struct OneShotBufferizationOptions;
} // namespace bufferization
#define GEN_PASS_DECL_CONVERTELEMENTWISETOLINALGPASS
#define GEN_PASS_DECL_LINALGINITTENSORTOALLOCTENSORPASS
#define GEN_PASS_DECL_LINALGFOLDUNITEXTENTDIMSPASS
#define GEN_PASS_DECL_LINALGELEMENTWISEOPFUSIONPASS
#define GEN_PASS_DECL_LINALGNAMEDOPCONVERSIONPASS
#define GEN_PASS_DECL_LINALGINLINESCALAROPERANDSPASS
#define GEN_PASS_DECL_LINALGLOWERTOAFFINELOOPSPASS
#define GEN_PASS_DECL_LINALGLOWERTOLOOPSPASS
#define GEN_PASS_DECL_LINALGLOWERTOPARALLELLOOPSPASS
#define GEN_PASS_DECL_LINALGBUFFERIZEPASS
#define GEN_PASS_DECL_LINALGTILINGPASS
#define GEN_PASS_DECL_LINALGGENERALIZATIONPASS
#define GEN_PASS_DECL_LINALGDETENSORIZEPASS
#define GEN_PASS_DECL_LINALGSTRATEGYTILEANDFUSEPASS
#define GEN_PASS_DECL_LINALGSTRATEGYTILEPASS
#define GEN_PASS_DECL_LINALGSTRATEGYPADPASS
#define GEN_PASS_DECL_LINALGSTRATEGYDECOMPOSEPASS
#define GEN_PASS_DECL_LINALGSTRATEGYPEELPASS
#define GEN_PASS_DECL_LINALGSTRATEGYVECTORIZEPASS
#define GEN_PASS_DECL_LINALGSTRATEGYENABLEPASS
#define GEN_PASS_DECL_LINALGSTRATEGYLOWERVECTORSPASS
#define GEN_PASS_DECL_LINALGSTRATEGYREMOVEMARKERSPASS
#include "mlir/Dialect/Linalg/Passes.h.inc"
std::unique_ptr<Pass> createConvertElementwiseToLinalgPass();
std::unique_ptr<Pass> createLinalgFoldUnitExtentDimsPass();

26
mlir/include/mlir/Dialect/Linalg/Passes.td
View File

@ -11,7 +11,7 @@
include "mlir/Pass/PassBase.td"
def ConvertElementwiseToLinalgPass : Pass<"convert-elementwise-to-linalg", ""> {
def ConvertElementwiseToLinalg : Pass<"convert-elementwise-to-linalg", ""> {
let summary = "Convert ElementwiseMappable ops to linalg";
let description = [{
Convert ops with the `ElementwiseMappable` trait to linalg parallel loops.
@ -24,7 +24,7 @@ def ConvertElementwiseToLinalgPass : Pass<"convert-elementwise-to-linalg", ""> {
let dependentDialects = ["linalg::LinalgDialect", "memref::MemRefDialect"];
}
def LinalgInitTensorToAllocTensorPass : Pass<"linalg-init-tensor-to-alloc-tensor"> {
def LinalgInitTensorToAllocTensor : Pass<"linalg-init-tensor-to-alloc-tensor"> {
let summary = "Replace all init_tensor ops by alloc_tensor ops.";
let description = [{
init_tensor ops return a tensor of unspecified contents who's only purpose
@ -34,7 +34,7 @@ def LinalgInitTensorToAllocTensorPass : Pass<"linalg-init-tensor-to-alloc-tensor
let constructor = "mlir::createLinalgInitTensorToAllocTensorPass()";
}
def LinalgFoldUnitExtentDimsPass : Pass<"linalg-fold-unit-extent-dims", ""> {
def LinalgFoldUnitExtentDims : Pass<"linalg-fold-unit-extent-dims", ""> {
let summary = "Remove unit-extent dimension in Linalg ops on tensors";
let constructor = "mlir::createLinalgFoldUnitExtentDimsPass()";
let options = [
@ -48,7 +48,7 @@ def LinalgFoldUnitExtentDimsPass : Pass<"linalg-fold-unit-extent-dims", ""> {
];
}
def LinalgElementwiseOpFusionPass : Pass<"linalg-fuse-elementwise-ops"> {
def LinalgElementwiseOpFusion : Pass<"linalg-fuse-elementwise-ops"> {
let summary = "Fuse elementwise operations on tensors";
let constructor = "mlir::createLinalgElementwiseOpFusionPass()";
let dependentDialects = [
@ -56,13 +56,13 @@ def LinalgElementwiseOpFusionPass : Pass<"linalg-fuse-elementwise-ops"> {
];
}
def LinalgNamedOpConversionPass : Pass<"linalg-named-op-conversion"> {
def LinalgNamedOpConversion: Pass<"linalg-named-op-conversion"> {
let summary = "Convert from one named linalg op to another.";
let constructor = "mlir::createLinalgNamedOpConversionPass()";
let dependentDialects = ["linalg::LinalgDialect", "tensor::TensorDialect"];
}
def LinalgInlineScalarOperandsPass : Pass<"linalg-inline-scalar-operands", "func::FuncOp"> {
def LinalgInlineScalarOperands : Pass<"linalg-inline-scalar-operands", "func::FuncOp"> {
let summary = "Inline scalar operands into linalg generic ops";
let constructor = "mlir::createLinalgInlineScalarOperandsPass()";
let dependentDialects = [
@ -70,7 +70,7 @@ def LinalgInlineScalarOperandsPass : Pass<"linalg-inline-scalar-operands", "func
];
}
def LinalgLowerToAffineLoopsPass : Pass<"convert-linalg-to-affine-loops", "func::FuncOp"> {
def LinalgLowerToAffineLoops : Pass<"convert-linalg-to-affine-loops", "func::FuncOp"> {
let summary = "Lower the operations from the linalg dialect into affine "
"loops";
let constructor = "mlir::createConvertLinalgToAffineLoopsPass()";
@ -78,7 +78,7 @@ def LinalgLowerToAffineLoopsPass : Pass<"convert-linalg-to-affine-loops", "func:
"AffineDialect", "linalg::LinalgDialect", "memref::MemRefDialect"];
}
def LinalgLowerToLoopsPass : Pass<"convert-linalg-to-loops", "func::FuncOp"> {
def LinalgLowerToLoops : Pass<"convert-linalg-to-loops", "func::FuncOp"> {
let summary = "Lower the operations from the linalg dialect into loops";
let constructor = "mlir::createConvertLinalgToLoopsPass()";
let dependentDialects = [
@ -88,7 +88,7 @@ def LinalgLowerToLoopsPass : Pass<"convert-linalg-to-loops", "func::FuncOp"> {
];
}
def LinalgLowerToParallelLoopsPass
def LinalgLowerToParallelLoops
: Pass<"convert-linalg-to-parallel-loops", "func::FuncOp"> {
let summary = "Lower the operations from the linalg dialect into parallel "
"loops";
@ -101,7 +101,7 @@ def LinalgLowerToParallelLoopsPass
];
}
def LinalgBufferizePass : Pass<"linalg-bufferize", "func::FuncOp"> {
def LinalgBufferize : Pass<"linalg-bufferize", "func::FuncOp"> {
let summary = "Bufferize the linalg dialect";
let constructor = "mlir::createLinalgBufferizePass()";
let dependentDialects = [
@ -112,7 +112,7 @@ def LinalgBufferizePass : Pass<"linalg-bufferize", "func::FuncOp"> {
];
}
def LinalgTilingPass : Pass<"linalg-tile", "func::FuncOp"> {
def LinalgTiling : Pass<"linalg-tile", "func::FuncOp"> {
let summary = "Tile operations in the linalg dialect";
let constructor = "mlir::createLinalgTilingPass()";
let dependentDialects = [
@ -128,13 +128,13 @@ def LinalgTilingPass : Pass<"linalg-tile", "func::FuncOp"> {
];
}
def LinalgGeneralizationPass : Pass<"linalg-generalize-named-ops", "func::FuncOp"> {
def LinalgGeneralization : Pass<"linalg-generalize-named-ops", "func::FuncOp"> {
let summary = "Convert named ops into generic ops";
let constructor = "mlir::createLinalgGeneralizationPass()";
let dependentDialects = ["linalg::LinalgDialect"];
}
def LinalgDetensorizePass : Pass<"linalg-detensorize", ""> {
def LinalgDetensorize : Pass<"linalg-detensorize", ""> {
let summary = "Detensorize linalg ops";
let constructor = "mlir::createLinalgDetensorizePass()";
let dependentDialects = [];

7
mlir/include/mlir/Dialect/MemRef/Transforms/Passes.h
View File

@ -86,13 +86,6 @@ LogicalResult multiBuffer(memref::AllocOp allocOp, unsigned multiplier);
// Passes
//===----------------------------------------------------------------------===//
#define GEN_PASS_DECL_EXPANDOPSPASS
#define GEN_PASS_DECL_FOLDMEMREFALIASOPSPASS
#define GEN_PASS_DECL_NORMALIZEMEMREFSPASS
#define GEN_PASS_DECL_RESOLVERANKEDSHAPETYPERESULTDIMSPASS
#define GEN_PASS_DECL_RESOLVESHAPEDTYPERESULTDIMSPASS
#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc"
/// Creates an instance of the ExpandOps pass that legalizes memref dialect ops
/// to be convertible to LLVM. For example, `memref.reshape` gets converted to
/// `memref_reinterpret_cast`.

10
mlir/include/mlir/Dialect/MemRef/Transforms/Passes.td
View File

@ -11,12 +11,12 @@
include "mlir/Pass/PassBase.td"
def ExpandOpsPass : Pass<"memref-expand"> {
def ExpandOps : Pass<"memref-expand"> {
let summary = "Legalize memref operations to be convertible to LLVM.";
let constructor = "mlir::memref::createExpandOpsPass()";
}
def FoldMemRefAliasOpsPass : Pass<"fold-memref-alias-ops"> {
def FoldMemRefAliasOps : Pass<"fold-memref-alias-ops"> {
let summary = "Fold memref alias ops into consumer load/store ops";
let description = [{
The pass folds loading/storing from/to memref aliasing ops to loading/storing
@ -28,7 +28,7 @@ def FoldMemRefAliasOpsPass : Pass<"fold-memref-alias-ops"> {
];
}
def NormalizeMemRefsPass : Pass<"normalize-memrefs", "ModuleOp"> {
def NormalizeMemRefs : Pass<"normalize-memrefs", "ModuleOp"> {
let summary = "Normalize memrefs";
let description = [{
This pass transforms memref types with a non-trivial
@ -144,7 +144,7 @@ def NormalizeMemRefsPass : Pass<"normalize-memrefs", "ModuleOp"> {
let dependentDialects = ["AffineDialect"];
}
def ResolveRankedShapeTypeResultDimsPass :
def ResolveRankedShapeTypeResultDims :
Pass<"resolve-ranked-shaped-type-result-dims"> {
let summary = "Resolve memref.dim of result values of ranked shape type";
let description = [{
@ -159,7 +159,7 @@ def ResolveRankedShapeTypeResultDimsPass :
];
}
def ResolveShapedTypeResultDimsPass : Pass<"resolve-shaped-type-result-dims"> {
def ResolveShapedTypeResultDims : Pass<"resolve-shaped-type-result-dims"> {
let summary = "Resolve memref.dim of result values";
let description = [{
The pass resolves memref.dim of result of operations that

3
mlir/include/mlir/Dialect/NVGPU/Passes.h
View File

@ -17,9 +17,6 @@
namespace mlir {
namespace nvgpu {
#define GEN_PASS_DECL_OPTIMIZESHAREDMEMORYPASS
#include "mlir/Dialect/NVGPU/Passes.h.inc"
/// Create a pass to optimize shared memory reads and writes.
std::unique_ptr<Pass> createOptimizeSharedMemoryPass();

2
mlir/include/mlir/Dialect/NVGPU/Passes.td
View File

@ -11,7 +11,7 @@
include "mlir/Pass/PassBase.td"
def OptimizeSharedMemoryPass : Pass<"nvgpu-optimize-shared-memory"> {
def OptimizeSharedMemory : Pass<"nvgpu-optimize-shared-memory"> {
let summary = "Optimizes accesses to shard memory memrefs in order to reduce bank conflicts.";
let constructor = "mlir::nvgpu::createOptimizeSharedMemoryPass()";
let dependentDialects = [

52
mlir/include/mlir/Dialect/SCF/Transforms/Passes.h
View File

@ -17,17 +17,47 @@
namespace mlir {
#define GEN_PASS_DECL_SCFBUFFERIZEPASS
#define GEN_PASS_DECL_SCFFORLOOPCANONICALIZATIONPASS
#define GEN_PASS_DECL_SCFFORLOOPPEELINGPASS
#define GEN_PASS_DECL_SCFFORLOOPSPECIALIZATIONPASS
#define GEN_PASS_DECL_SCFPARALLELLOOPFUSIONPASS
#define GEN_PASS_DECL_SCFPARALLELLOOPCOLLAPSINGPASS
#define GEN_PASS_DECL_SCFPARALLELLOOPSPECIALIZATIONPASS
#define GEN_PASS_DECL_SCFPARALLELLOOPTILINGPASS
#define GEN_PASS_DECL_SCFFORLOOPRANGEFOLDINGPASS
#define GEN_PASS_DECL_SCFFORTOWHILELOOPPASS
#include "mlir/Dialect/SCF/Transforms/Passes.h.inc"
/// Creates a pass that bufferizes the SCF dialect.
std::unique_ptr<Pass> createSCFBufferizePass();
/// Creates a pass that specializes for loop for unrolling and
/// vectorization.
std::unique_ptr<Pass> createForLoopSpecializationPass();
/// Creates a pass that peels for loops at their upper bounds for
/// better vectorization.
std::unique_ptr<Pass> createForLoopPeelingPass();
/// Creates a pass that canonicalizes affine.min and affine.max operations
/// inside of scf.for loops with known lower and upper bounds.
std::unique_ptr<Pass> createSCFForLoopCanonicalizationPass();
/// Creates a pass that transforms a single ParallelLoop over N induction
/// variables into another ParallelLoop over less than N induction variables.
std::unique_ptr<Pass> createParallelLoopCollapsingPass();
/// Creates a loop fusion pass which fuses parallel loops.
std::unique_ptr<Pass> createParallelLoopFusionPass();
/// Creates a pass that specializes parallel loop for unrolling and
/// vectorization.
std::unique_ptr<Pass> createParallelLoopSpecializationPass();
/// Creates a pass which tiles innermost parallel loops.
/// If noMinMaxBounds, the upper bound of the inner loop will
/// be a same value among different outter loop iterations, and
/// an additional inbound check will be emitted inside the internal
/// loops.
std::unique_ptr<Pass>
createParallelLoopTilingPass(llvm::ArrayRef<int64_t> tileSize = {},
bool noMinMaxBounds = false);
/// Creates a pass which folds arith ops on induction variable into
/// loop range.
std::unique_ptr<Pass> createForLoopRangeFoldingPass();
// Creates a pass which lowers for loops into while loops.
std::unique_ptr<Pass> createForToWhileLoopPass();
//===----------------------------------------------------------------------===//
// Registration

53
mlir/include/mlir/Dialect/SCF/Transforms/Passes.td
View File

@ -11,30 +11,26 @@
include "mlir/Pass/PassBase.td"
def SCFBufferizePass : Pass<"scf-bufferize"> {
def SCFBufferize : Pass<"scf-bufferize"> {
let summary = "Bufferize the scf dialect.";
let constructor = "mlir::createSCFBufferizePass()";
let dependentDialects = ["bufferization::BufferizationDialect",
"memref::MemRefDialect"];
}
// Note: Making these canonicalization patterns would require a dependency
// of the SCF dialect on the Affine/Tensor/MemRef dialects or vice versa.
def SCFForLoopCanonicalizationPass
def SCFForLoopCanonicalization
: Pass<"scf-for-loop-canonicalization"> {
let summary = "Canonicalize operations within scf.for loop bodies";
let description = [{
Canonicalize `affine.min` and `affine.max` operations inside of `scf.for`
loops with known lower and upper bounds.
}];
let constructor = "mlir::createSCFForLoopCanonicalizationPass()";
let dependentDialects = ["AffineDialect", "tensor::TensorDialect",
"memref::MemRefDialect"];
}
def SCFForLoopPeelingPass : Pass<"scf-for-loop-peeling"> {
def SCFForLoopPeeling : Pass<"scf-for-loop-peeling"> {
let summary = "Peel `for` loops at their upper bounds.";
let description = [{
Peel for loops at their upper bounds for better vectorization.
}];
let constructor = "mlir::createForLoopPeelingPass()";
let options = [
Option<"skipPartial", "skip-partial", "bool",
/*default=*/"true",
@ -44,20 +40,19 @@ def SCFForLoopPeelingPass : Pass<"scf-for-loop-peeling"> {
let dependentDialects = ["AffineDialect"];
}
def SCFForLoopSpecializationPass : Pass<"scf-for-loop-specialization"> {
def SCFForLoopSpecialization : Pass<"scf-for-loop-specialization"> {
let summary = "Specialize `for` loops for vectorization";
let constructor = "mlir::createForLoopSpecializationPass()";
}
def SCFParallelLoopFusionPass : Pass<"scf-parallel-loop-fusion"> {
def SCFParallelLoopFusion : Pass<"scf-parallel-loop-fusion"> {
let summary = "Fuse adjacent parallel loops";
let constructor = "mlir::createParallelLoopFusionPass()";
}
def SCFParallelLoopCollapsingPass : Pass<"scf-parallel-loop-collapsing"> {
def SCFParallelLoopCollapsing : Pass<"scf-parallel-loop-collapsing"> {
let summary = "Collapse parallel loops to use less induction variables";
let description = [{
Transforms a single `ParallelLoop` over N induction variables into
another `ParallelLoop` over less than N induction variables.
}];
let constructor = "mlir::createParallelLoopCollapsingPass()";
let options = [
ListOption<"clCollapsedIndices0", "collapsed-indices-0", "unsigned",
"Which loop indices to combine 0th loop index">,
@ -68,22 +63,15 @@ def SCFParallelLoopCollapsingPass : Pass<"scf-parallel-loop-collapsing"> {
];
}
def SCFParallelLoopSpecializationPass
def SCFParallelLoopSpecialization
: Pass<"scf-parallel-loop-specialization"> {
let summary = "Specialize parallel loops for vectorization";
let description = [{
Specialize parallel loop for unrolling and vectorization.
}];
let constructor = "mlir::createParallelLoopSpecializationPass()";
}
def SCFParallelLoopTilingPass : Pass<"scf-parallel-loop-tiling"> {
def SCFParallelLoopTiling : Pass<"scf-parallel-loop-tiling"> {
let summary = "Tile parallel loops";
let description = [{
Tile innermost parallel loops.
If `noMinMaxBounds`, the upper bound of the inner loop will be a same
value among different outter loop iterations, and an additional inbound
check will be emitted inside the internal loops.
}];
let constructor = "mlir::createParallelLoopTilingPass()";
let options = [
ListOption<"tileSizes", "parallel-loop-tile-sizes", "int64_t",
"Factors to tile parallel loops by">,
@ -95,15 +83,14 @@ def SCFParallelLoopTilingPass : Pass<"scf-parallel-loop-tiling"> {
let dependentDialects = ["AffineDialect"];
}
def SCFForLoopRangeFoldingPass : Pass<"scf-for-loop-range-folding"> {
def SCFForLoopRangeFolding : Pass<"scf-for-loop-range-folding"> {
let summary = "Fold add/mul ops into loop range";
let description = [{
Fold arith ops on induction variable into loop range.
}];
let constructor = "mlir::createForLoopRangeFoldingPass()";
}
def SCFForToWhileLoopPass : Pass<"scf-for-to-while"> {
def SCFForToWhileLoop : Pass<"scf-for-to-while"> {
let summary = "Convert SCF for loops to SCF while loops";
let constructor = "mlir::createForToWhileLoopPass()";
let description = [{
This pass transforms SCF.ForOp operations to SCF.WhileOp. The For loop
condition is placed in the 'before' region of the while operation, and the

8
mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h
View File

@ -27,14 +27,6 @@ class ModuleOp;
// Passes
//===----------------------------------------------------------------------===//
#define GEN_PASS_DECL_SPIRVDECORATECOMPOSITETYPELAYOUTPASS
#define GEN_PASS_DECL_SPIRVCANONICALIZEGLPASS
#define GEN_PASS_DECL_SPIRVLOWERABIATTRIBUTESPASS
#define GEN_PASS_DECL_SPIRVREWRITEINSERTSPASS
#define GEN_PASS_DECL_SPIRVUNIFYALIASEDRESOURCEPASS
#define GEN_PASS_DECL_SPIRVUPDATEVCEPASS
#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc"
/// Creates a pass to run canoncalization patterns that involve GL ops.
/// These patterns cannot be run in default canonicalization because GL ops
/// aren't always available. So they should be involed specifically when needed.

10
mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td
View File

@ -11,18 +11,18 @@
include "mlir/Pass/PassBase.td"
def SPIRVDecorateCompositeTypeLayoutPass
: Pass<"decorate-spirv-composite-type-layout", "mlir::ModuleOp"> {
def SPIRVCompositeTypeLayout
: Pass<"decorate-spirv-composite-type-layout", "ModuleOp"> {
let summary = "Decorate SPIR-V composite type with layout info";
let constructor = "mlir::spirv::createDecorateSPIRVCompositeTypeLayoutPass()";
}
def SPIRVCanonicalizeGLPass : Pass<"spirv-canonicalize-gl", ""> {
def SPIRVCanonicalizeGL : Pass<"spirv-canonicalize-gl", ""> {
let summary = "Run canonicalization involving GLSL ops";
let constructor = "mlir::spirv::createCanonicalizeGLPass()";
}
def SPIRVLowerABIAttributesPass : Pass<"spirv-lower-abi-attrs", "spirv::ModuleOp"> {
def SPIRVLowerABIAttributes : Pass<"spirv-lower-abi-attrs", "spirv::ModuleOp"> {
let summary = "Decorate SPIR-V composite type with layout info";
let constructor = "mlir::spirv::createLowerABIAttributesPass()";
}
@ -40,7 +40,7 @@ def SPIRVUnifyAliasedResourcePass
let constructor = "mlir::spirv::createUnifyAliasedResourcePass()";
}
def SPIRVUpdateVCEPass : Pass<"spirv-update-vce", "spirv::ModuleOp"> {
def SPIRVUpdateVCE : Pass<"spirv-update-vce", "spirv::ModuleOp"> {
let summary = "Deduce and attach minimal (version, capabilities, extensions) "
"requirements to spv.module ops";
let constructor = "mlir::spirv::createUpdateVersionCapabilityExtensionPass()";

Some files were not shown because too many files have changed in this diff Show More