diff options
Diffstat (limited to 'mlir/lib')
172 files changed, 1320 insertions, 1073 deletions
diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp index 0e2894aae756..81b8f6e7a03a 100644 --- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp +++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp @@ -7,11 +7,17 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h" -#include "../PassDetail.h" + #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Dialect/AMDGPU/AMDGPUDialect.h" #include "mlir/Dialect/LLVMIR/ROCDLDialect.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTAMDGPUTOROCDL +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::amdgpu; @@ -269,7 +275,7 @@ struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> { }; struct ConvertAMDGPUToROCDLPass - : public ConvertAMDGPUToROCDLBase<ConvertAMDGPUToROCDLPass> { + : public impl::ConvertAMDGPUToROCDLBase<ConvertAMDGPUToROCDLPass> { ConvertAMDGPUToROCDLPass() = default; void runOnOperation() override { diff --git a/mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp b/mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp index 2c9a1eaa2bf6..a1e8e31d9a92 100644 --- a/mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp +++ b/mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp @@ -13,7 +13,6 @@ #include "mlir/Conversion/AffineToStandard/AffineToStandard.h" -#include "../PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -22,10 +21,14 @@ #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/IntegerSet.h" #include "mlir/IR/MLIRContext.h" -#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/Passes.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTAFFINETOSTANDARD +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::vector; @@ -544,7 +547,8 @@ void mlir::populateAffineToVectorConversionPatterns( } namespace { -class LowerAffinePass : public ConvertAffineToStandardBase<LowerAffinePass> { +class LowerAffinePass + : public impl::ConvertAffineToStandardBase<LowerAffinePass> { void runOnOperation() override { RewritePatternSet patterns(&getContext()); populateAffineToStdConversionPatterns(patterns); diff --git a/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp b/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp index df6df155f41d..cdfe8589c4ff 100644 --- a/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp +++ b/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp @@ -7,12 +7,18 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" -#include "../PassDetail.h" + #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/VectorPattern.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/TypeUtilities.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTARITHMETICTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -315,7 +321,7 @@ CmpFOpLowering::matchAndRewrite(arith::CmpFOp op, OpAdaptor adaptor, namespace { struct ConvertArithmeticToLLVMPass - : public ConvertArithmeticToLLVMBase<ConvertArithmeticToLLVMPass> { + : public impl::ConvertArithmeticToLLVMBase<ConvertArithmeticToLLVMPass> { ConvertArithmeticToLLVMPass() = default; void runOnOperation() override { diff --git a/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp b/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp index 6ca37c667dab..b65d13c13b94 100644 --- a/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp +++ b/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h" -#include "../PassDetail.h" + #include "../SPIRVCommon/Pattern.h" #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" @@ -20,6 +20,11 @@ #include "llvm/ADT/ArrayRef.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTARITHMETICTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "arith-to-spirv-pattern" using namespace mlir; @@ -855,8 +860,8 @@ AddICarryOpPattern::matchAndRewrite(arith::AddUICarryOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const { Type dstElemTy = adaptor.getLhs().getType(); Location loc = op->getLoc(); - Value result = rewriter.create<spirv::IAddCarryOp>( - loc, adaptor.getLhs(), adaptor.getRhs()); + Value result = rewriter.create<spirv::IAddCarryOp>(loc, adaptor.getLhs(), + adaptor.getRhs()); Value sumResult = rewriter.create<spirv::CompositeExtractOp>( loc, result, llvm::makeArrayRef(0)); @@ -948,7 +953,7 @@ void mlir::arith::populateArithmeticToSPIRVPatterns( namespace { struct ConvertArithmeticToSPIRVPass - : public ConvertArithmeticToSPIRVBase<ConvertArithmeticToSPIRVPass> { + : public impl::ConvertArithmeticToSPIRVBase<ConvertArithmeticToSPIRVPass> { void runOnOperation() override { Operation *op = getOperation(); auto targetAttr = spirv::lookupTargetEnvOrDefault(op); diff --git a/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp b/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp index a2f5641a0f5e..691cd23c6ed1 100644 --- a/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp +++ b/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.h" -#include "../PassDetail.h" + #include "mlir/Dialect/ArmNeon/ArmNeonDialect.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "mlir/IR/PatternMatch.h" @@ -15,6 +15,11 @@ #include "mlir/Pass/PassRegistry.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTARMNEON2DTOINTR +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::arm_neon; @@ -47,7 +52,7 @@ public: }; class ConvertArmNeon2dToIntr - : public ConvertArmNeon2dToIntrBase<ConvertArmNeon2dToIntr> { + : public impl::ConvertArmNeon2dToIntrBase<ConvertArmNeon2dToIntr> { void runOnOperation() override { auto *context = &getContext(); diff --git a/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp b/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp index 26154e4ed66d..8733459dc471 100644 --- a/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp +++ b/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp @@ -8,7 +8,6 @@ #include "mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h" -#include "../PassDetail.h" #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/TypeConverter.h" @@ -24,6 +23,11 @@ #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTASYNCTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "convert-async-to-llvm" using namespace mlir; @@ -986,7 +990,7 @@ public: namespace { struct ConvertAsyncToLLVMPass - : public ConvertAsyncToLLVMBase<ConvertAsyncToLLVMPass> { + : public impl::ConvertAsyncToLLVMBase<ConvertAsyncToLLVMPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp b/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp index d98bb2a23335..fcbd27f371ad 100644 --- a/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp +++ b/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp @@ -12,14 +12,20 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/BufferizationToMemRef/BufferizationToMemRef.h" -#include "../PassDetail.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/BuiltinTypes.h" +#include "mlir/Pass/Pass.h" #include "mlir/Support/LogicalResult.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTBUFFERIZATIONTOMEMREF +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -82,7 +88,7 @@ void mlir::populateBufferizationToMemRefConversionPatterns( namespace { struct BufferizationToMemRefPass - : public ConvertBufferizationToMemRefBase<BufferizationToMemRefPass> { + : public impl::ConvertBufferizationToMemRefBase<BufferizationToMemRefPass> { BufferizationToMemRefPass() = default; void runOnOperation() override { diff --git a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp index 820086e6d2ba..64e64c7935aa 100644 --- a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp +++ b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp @@ -8,12 +8,17 @@ #include "mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h" -#include "../PassDetail.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Complex/IR/Complex.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTCOMPLEXTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::LLVM; @@ -318,7 +323,7 @@ void mlir::populateComplexToLLVMConversionPatterns( namespace { struct ConvertComplexToLLVMPass - : public ConvertComplexToLLVMBase<ConvertComplexToLLVMPass> { + : public impl::ConvertComplexToLLVMBase<ConvertComplexToLLVMPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp b/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp index 558b1cc66a32..fa8db2c96517 100644 --- a/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp +++ b/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp @@ -8,10 +8,15 @@ #include "mlir/Conversion/ComplexToLibm/ComplexToLibm.h" -#include "../PassDetail.h" #include "mlir/Dialect/Complex/IR/Complex.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/PatternMatch.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTCOMPLEXTOLIBM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -117,7 +122,7 @@ void mlir::populateComplexToLibmConversionPatterns(RewritePatternSet &patterns, namespace { struct ConvertComplexToLibmPass - : public ConvertComplexToLibmBase<ConvertComplexToLibmPass> { + : public impl::ConvertComplexToLibmBase<ConvertComplexToLibmPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp b/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp index 064b0db08a41..d1af84c22c02 100644 --- a/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp +++ b/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp @@ -8,16 +8,20 @@ #include "mlir/Conversion/ComplexToStandard/ComplexToStandard.h" -#include <memory> -#include <type_traits> - -#include "../PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Complex/IR/Complex.h" #include "mlir/Dialect/Math/IR/Math.h" #include "mlir/IR/ImplicitLocOpBuilder.h" #include "mlir/IR/PatternMatch.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +#include <memory> +#include <type_traits> + +namespace mlir { +#define GEN_PASS_DEF_CONVERTCOMPLEXTOSTANDARD +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -1081,7 +1085,7 @@ void mlir::populateComplexToStandardConversionPatterns( namespace { struct ConvertComplexToStandardPass - : public ConvertComplexToStandardBase<ConvertComplexToStandardPass> { + : public impl::ConvertComplexToStandardBase<ConvertComplexToStandardPass> { void runOnOperation() override; }; diff --git a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp index 89012704541d..ea7fc8f5b7ad 100644 --- a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp +++ b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp @@ -12,7 +12,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h" -#include "../PassDetail.h" + #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Conversion/LLVMCommon/VectorPattern.h" @@ -21,10 +21,16 @@ #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/PatternMatch.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/StringRef.h" #include <functional> +namespace mlir { +#define GEN_PASS_DEF_CONVERTCONTROLFLOWTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; #define PASS_NAME "convert-cf-to-llvm" @@ -196,7 +202,7 @@ void mlir::cf::populateControlFlowToLLVMConversionPatterns( namespace { /// A pass converting MLIR operations into the LLVM IR dialect. struct ConvertControlFlowToLLVM - : public ConvertControlFlowToLLVMBase<ConvertControlFlowToLLVM> { + : public impl::ConvertControlFlowToLLVMBase<ConvertControlFlowToLLVM> { ConvertControlFlowToLLVM() = default; /// Run the dialect converter on the module. diff --git a/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp b/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp index 6cd237ea3e0e..69b2a874d7b0 100644 --- a/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp +++ b/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp @@ -11,17 +11,23 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTCONTROLFLOWTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { /// A pass converting MLIR ControlFlow operations into the SPIR-V dialect. class ConvertControlFlowToSPIRVPass - : public ConvertControlFlowToSPIRVBase<ConvertControlFlowToSPIRVPass> { + : public impl::ConvertControlFlowToSPIRVBase< + ConvertControlFlowToSPIRVPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp index 4b80bcb90cc3..1a77a1673ffa 100644 --- a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp +++ b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp @@ -11,12 +11,12 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" +#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h" + #include "mlir/Analysis/DataLayoutAnalysis.h" #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h" #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h" -#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Conversion/LLVMCommon/VectorPattern.h" @@ -47,6 +47,11 @@ #include <algorithm> #include <functional> +namespace mlir { +#define GEN_PASS_DEF_CONVERTFUNCTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; #define PASS_NAME "convert-func-to-llvm" @@ -699,7 +704,7 @@ void mlir::populateFuncToLLVMConversionPatterns(LLVMTypeConverter &converter, namespace { /// A pass converting Func operations into the LLVM IR dialect. struct ConvertFuncToLLVMPass - : public ConvertFuncToLLVMBase<ConvertFuncToLLVMPass> { + : public impl::ConvertFuncToLLVMBase<ConvertFuncToLLVMPass> { ConvertFuncToLLVMPass() = default; ConvertFuncToLLVMPass(bool useBarePtrCallConv, unsigned indexBitwidth, bool useAlignedAlloc, diff --git a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp index d2416feea763..303692ad8bcb 100644 --- a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp +++ b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp @@ -11,17 +11,22 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRVPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTFUNCTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { /// A pass converting MLIR Func operations into the SPIR-V dialect. class ConvertFuncToSPIRVPass - : public ConvertFuncToSPIRVBase<ConvertFuncToSPIRVPass> { + : public impl::ConvertFuncToSPIRVBase<ConvertFuncToSPIRVPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp b/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp index e4448ff1f310..378ee5340cd6 100644 --- a/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp +++ b/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp @@ -15,7 +15,6 @@ #include "mlir/Conversion/GPUCommon/GPUCommonPass.h" -#include "../PassDetail.h" #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h" #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h" @@ -38,6 +37,11 @@ #include "llvm/Support/Error.h" #include "llvm/Support/FormatVariadic.h" +namespace mlir { +#define GEN_PASS_DEF_GPUTOLLVMCONVERSIONPASS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; static constexpr const char *kGpuBinaryStorageSuffix = "_gpubin_cst"; @@ -45,7 +49,7 @@ static constexpr const char *kGpuBinaryStorageSuffix = "_gpubin_cst"; namespace { class GpuToLLVMConversionPass - : public GpuToLLVMConversionPassBase<GpuToLLVMConversionPass> { + : public impl::GpuToLLVMConversionPassBase<GpuToLLVMConversionPass> { public: GpuToLLVMConversionPass() = default; diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp index 2e89a3a23508..f75cb85de30b 100644 --- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp +++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp @@ -36,7 +36,11 @@ #include "../GPUCommon/GPUOpsLowering.h" #include "../GPUCommon/IndexIntrinsicsOpLowering.h" #include "../GPUCommon/OpToFuncCallLowering.h" -#include "../PassDetail.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTGPUOPSTONVVMOPS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -152,7 +156,7 @@ struct GPULaneIdOpToNVVM : ConvertOpToLLVMPattern<gpu::LaneIdOp> { /// This pass only handles device code and is not meant to be run on GPU host /// code. struct LowerGpuOpsToNVVMOpsPass - : public ConvertGpuOpsToNVVMOpsBase<LowerGpuOpsToNVVMOpsPass> { + : public impl::ConvertGpuOpsToNVVMOpsBase<LowerGpuOpsToNVVMOpsPass> { LowerGpuOpsToNVVMOpsPass() = default; LowerGpuOpsToNVVMOpsPass(unsigned indexBitwidth) { this->indexBitwidth = indexBitwidth; diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp index 2c7f70046519..43213e8d412e 100644 --- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp +++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp @@ -37,7 +37,11 @@ #include "../GPUCommon/GPUOpsLowering.h" #include "../GPUCommon/IndexIntrinsicsOpLowering.h" #include "../GPUCommon/OpToFuncCallLowering.h" -#include "../PassDetail.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTGPUOPSTOROCDLOPS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -62,7 +66,7 @@ namespace { // This pass only handles device code and is not meant to be run on GPU host // code. struct LowerGpuOpsToROCDLOpsPass - : public ConvertGpuOpsToROCDLOpsBase<LowerGpuOpsToROCDLOpsPass> { + : public impl::ConvertGpuOpsToROCDLOpsBase<LowerGpuOpsToROCDLOpsPass> { LowerGpuOpsToROCDLOpsPass() = default; LowerGpuOpsToROCDLOpsPass(const std::string &chipset, unsigned indexBitwidth, bool useBarePtrCallConv, diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp index 480c903f83ec..b880c65ff847 100644 --- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp +++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp @@ -13,7 +13,6 @@ #include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h" -#include "../PassDetail.h" #include "mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h" #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h" #include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h" @@ -23,6 +22,11 @@ #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTGPUTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -35,7 +39,7 @@ namespace { /// replace it). /// /// 2) Lower the body of the spirv::ModuleOp. -class GPUToSPIRVPass : public ConvertGPUToSPIRVBase<GPUToSPIRVPass> { +class GPUToSPIRVPass : public impl::ConvertGPUToSPIRVBase<GPUToSPIRVPass> { public: explicit GPUToSPIRVPass(bool mapMemorySpace) : mapMemorySpace(mapMemorySpace) {} diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp index 361dae64ceaf..8f084958e11f 100644 --- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp +++ b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp @@ -13,8 +13,8 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h" + #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" @@ -23,8 +23,14 @@ #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/BuiltinTypes.h" +#include "mlir/Pass/Pass.h" #include "mlir/Target/SPIRV/Serialization.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; static constexpr const char *kSPIRVBlobAttrName = "spirv_blob"; @@ -38,7 +44,7 @@ namespace { /// function and attaching binary data and entry point name as an attributes to /// created vulkan launch call op. class ConvertGpuLaunchFuncToVulkanLaunchFunc - : public ConvertGpuLaunchFuncToVulkanLaunchFuncBase< + : public impl::ConvertGpuLaunchFuncToVulkanLaunchFuncBase< ConvertGpuLaunchFuncToVulkanLaunchFunc> { public: void runOnOperation() override; diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp index 3242add0799d..f3afc67933bc 100644 --- a/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp +++ b/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp @@ -14,16 +14,21 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinOps.h" - +#include "mlir/Pass/Pass.h" #include "llvm/ADT/SmallString.h" #include "llvm/Support/FormatVariadic.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTVULKANLAUNCHFUNCTOVULKANCALLS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; static constexpr const char *kCInterfaceVulkanLaunch = @@ -53,7 +58,7 @@ namespace { /// * deinitVulkan -- deinitializes vulkan runtime /// class VulkanLaunchFuncToVulkanCallsPass - : public ConvertVulkanLaunchFuncToVulkanCallsBase< + : public impl::ConvertVulkanLaunchFuncToVulkanCallsBase< VulkanLaunchFuncToVulkanCallsPass> { private: void initializeCachedTypes() { diff --git a/mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp b/mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp index 04512f0bf61f..7bc3680c2eae 100644 --- a/mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp +++ b/mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp @@ -8,7 +8,6 @@ #include "mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h" -#include "../PassDetail.h" #include "mlir/Conversion/AffineToStandard/AffineToStandard.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" @@ -41,6 +40,11 @@ #include "llvm/Support/Allocator.h" #include "llvm/Support/ErrorHandling.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTLINALGTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::LLVM; using namespace mlir::linalg; @@ -74,7 +78,7 @@ void mlir::populateLinalgToLLVMConversionPatterns(LLVMTypeConverter &converter, namespace { struct ConvertLinalgToLLVMPass - : public ConvertLinalgToLLVMBase<ConvertLinalgToLLVMPass> { + : public impl::ConvertLinalgToLLVMBase<ConvertLinalgToLLVMPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.cpp b/mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.cpp index 1238cfe4bc87..5732413a26e3 100644 --- a/mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.cpp +++ b/mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.cpp @@ -7,17 +7,24 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/LinalgToSPIRV/LinalgToSPIRV.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTLINALGTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; namespace { /// A pass converting MLIR Linalg ops into SPIR-V ops. -class LinalgToSPIRVPass : public ConvertLinalgToSPIRVBase<LinalgToSPIRVPass> { +class LinalgToSPIRVPass + : public impl::ConvertLinalgToSPIRVBase<LinalgToSPIRVPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp b/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp index 2779d247ddf1..299a908f2cc8 100644 --- a/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp +++ b/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp @@ -8,7 +8,6 @@ #include "mlir/Conversion/LinalgToStandard/LinalgToStandard.h" -#include "../PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" @@ -16,6 +15,12 @@ #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTLINALGTOSTANDARD +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::linalg; @@ -121,7 +126,7 @@ void mlir::linalg::populateLinalgToStandardConversionPatterns( namespace { struct ConvertLinalgToStandardPass - : public ConvertLinalgToStandardBase<ConvertLinalgToStandardPass> { + : public impl::ConvertLinalgToStandardBase<ConvertLinalgToStandardPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp b/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp index 25ee8cec62d8..e64611b7a245 100644 --- a/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp +++ b/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/MathToFuncs/MathToFuncs.h" -#include "../PassDetail.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/Func/IR/FuncOps.h" @@ -18,10 +18,16 @@ #include "mlir/Dialect/Vector/Utils/VectorUtils.h" #include "mlir/IR/ImplicitLocOpBuilder.h" #include "mlir/IR/TypeUtilities.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/DenseMap.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTMATHTOFUNCS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -318,7 +324,7 @@ IPowIOpLowering::matchAndRewrite(math::IPowIOp op, namespace { struct ConvertMathToFuncsPass - : public ConvertMathToFuncsBase<ConvertMathToFuncsPass> { + : public impl::ConvertMathToFuncsBase<ConvertMathToFuncsPass> { ConvertMathToFuncsPass() = default; void runOnOperation() override; diff --git a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp index 8161cc5e419f..a9ce30efdd08 100644 --- a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp +++ b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp @@ -7,13 +7,19 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/MathToLLVM/MathToLLVM.h" -#include "../PassDetail.h" + #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Conversion/LLVMCommon/VectorPattern.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/Math/IR/Math.h" #include "mlir/IR/TypeUtilities.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTMATHTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -250,7 +256,7 @@ struct RsqrtOpLowering : public ConvertOpToLLVMPattern<math::RsqrtOp> { }; struct ConvertMathToLLVMPass - : public ConvertMathToLLVMBase<ConvertMathToLLVMPass> { + : public impl::ConvertMathToLLVMBase<ConvertMathToLLVMPass> { ConvertMathToLLVMPass() = default; void runOnOperation() override { diff --git a/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp b/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp index 6dcf4532dd82..561fdfcab586 100644 --- a/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp +++ b/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp @@ -8,7 +8,6 @@ #include "mlir/Conversion/MathToLibm/MathToLibm.h" -#include "../PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" @@ -18,6 +17,12 @@ #include "mlir/Dialect/Vector/Utils/VectorUtils.h" #include "mlir/IR/BuiltinDialect.h" #include "mlir/IR/PatternMatch.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTMATHTOLIBM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -193,7 +198,7 @@ void mlir::populateMathToLibmConversionPatterns( namespace { struct ConvertMathToLibmPass - : public ConvertMathToLibmBase<ConvertMathToLibmPass> { + : public impl::ConvertMathToLibmBase<ConvertMathToLibmPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp b/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp index 6ef71d9f27a1..6e9ab607fef8 100644 --- a/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp +++ b/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp @@ -11,17 +11,23 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/MathToSPIRV/MathToSPIRVPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/MathToSPIRV/MathToSPIRV.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTMATHTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; namespace { /// A pass converting MLIR Math operations into the SPIR-V dialect. class ConvertMathToSPIRVPass - : public ConvertMathToSPIRVBase<ConvertMathToSPIRVPass> { + : public impl::ConvertMathToSPIRVBase<ConvertMathToSPIRVPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp b/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp index 8b3348f039b2..caef7494110c 100644 --- a/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp +++ b/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" -#include "../PassDetail.h" + #include "mlir/Analysis/DataLayoutAnalysis.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" @@ -20,8 +20,14 @@ #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/BlockAndValueMapping.h" +#include "mlir/Pass/Pass.h" #include "llvm/ADT/SmallBitVector.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTMEMREFTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -2042,7 +2048,8 @@ void mlir::populateMemRefToLLVMConversionPatterns(LLVMTypeConverter &converter, } namespace { -struct MemRefToLLVMPass : public ConvertMemRefToLLVMBase<MemRefToLLVMPass> { +struct MemRefToLLVMPass + : public impl::ConvertMemRefToLLVMBase<MemRefToLLVMPass> { MemRefToLLVMPass() = default; void runOnOperation() override { diff --git a/mlir/lib/Conversion/MemRefToSPIRV/MapMemRefStorageClassPass.cpp b/mlir/lib/Conversion/MemRefToSPIRV/MapMemRefStorageClassPass.cpp index e63f0594cf23..006c69f9d0f6 100644 --- a/mlir/lib/Conversion/MemRefToSPIRV/MapMemRefStorageClassPass.cpp +++ b/mlir/lib/Conversion/MemRefToSPIRV/MapMemRefStorageClassPass.cpp @@ -11,9 +11,9 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" -#include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h" #include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.h" + +#include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h" #include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h" @@ -23,6 +23,11 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_MAPMEMREFSTORAGECLASS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "mlir-map-memref-storage-class" using namespace mlir; @@ -284,7 +289,7 @@ void spirv::populateMemorySpaceToStorageClassPatterns( namespace { class MapMemRefStorageClassPass final - : public MapMemRefStorageClassBase<MapMemRefStorageClassPass> { + : public impl::MapMemRefStorageClassBase<MapMemRefStorageClassPass> { public: explicit MapMemRefStorageClassPass() { memorySpaceMap = spirv::mapMemorySpaceToVulkanStorageClass; diff --git a/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp b/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp index 44fc17bde62f..13b2b45d34c6 100644 --- a/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp +++ b/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp @@ -11,17 +11,22 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTMEMREFTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { /// A pass converting MLIR MemRef operations into the SPIR-V dialect. class ConvertMemRefToSPIRVPass - : public ConvertMemRefToSPIRVBase<ConvertMemRefToSPIRVPass> { + : public impl::ConvertMemRefToSPIRVBase<ConvertMemRefToSPIRVPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp index ff515bc69d67..9f1d19ddd6dc 100644 --- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp +++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp @@ -7,12 +7,18 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h" -#include "../PassDetail.h" + #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/NVVMDialect.h" #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTNVGPUTONVVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -327,7 +333,7 @@ struct MmaSyncOptoNVVM : public ConvertOpToLLVMPattern<nvgpu::MmaSyncOp> { }; struct ConvertNVGPUToNVVMPass - : public ConvertNVGPUToNVVMBase<ConvertNVGPUToNVVMPass> { + : public impl::ConvertNVGPUToNVVMBase<ConvertNVGPUToNVVMPass> { ConvertNVGPUToNVVMPass() = default; void runOnOperation() override { diff --git a/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp b/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp index 43a49a59fca3..525d12b7b8d7 100644 --- a/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp +++ b/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp @@ -6,12 +6,18 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" -#include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Conversion/OpenACCToLLVM/ConvertOpenACCToLLVM.h" + +#include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/OpenACC/OpenACC.h" #include "mlir/IR/Builders.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTOPENACCTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -148,7 +154,7 @@ void mlir::populateOpenACCToLLVMConversionPatterns( namespace { struct ConvertOpenACCToLLVMPass - : public ConvertOpenACCToLLVMBase<ConvertOpenACCToLLVMPass> { + : public impl::ConvertOpenACCToLLVMBase<ConvertOpenACCToLLVMPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp b/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp index ddc4517383cb..96ed682e2b7d 100644 --- a/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp +++ b/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp @@ -6,13 +6,19 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/OpenACCToSCF/ConvertOpenACCToSCF.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/OpenACC/OpenACC.h" #include "mlir/Dialect/SCF/IR/SCF.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTOPENACCTOSCF +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; //===----------------------------------------------------------------------===// @@ -56,7 +62,7 @@ void mlir::populateOpenACCToSCFConversionPatterns(RewritePatternSet &patterns) { namespace { struct ConvertOpenACCToSCFPass - : public ConvertOpenACCToSCFBase<ConvertOpenACCToSCFPass> { + : public impl::ConvertOpenACCToSCFBase<ConvertOpenACCToSCFPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp b/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp index 455246511533..7194fa695c5f 100644 --- a/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp +++ b/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp @@ -8,7 +8,6 @@ #include "mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h" -#include "../PassDetail.h" #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h" #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h" @@ -18,6 +17,12 @@ #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/OpenMP/OpenMPDialect.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTOPENMPTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -132,7 +137,7 @@ void mlir::populateOpenMPToLLVMConversionPatterns(LLVMTypeConverter &converter, namespace { struct ConvertOpenMPToLLVMPass - : public ConvertOpenMPToLLVMBase<ConvertOpenMPToLLVMPass> { + : public impl::ConvertOpenMPToLLVMBase<ConvertOpenMPToLLVMPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp b/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp index c24620a6729b..42ea3c9dc395 100644 --- a/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp +++ b/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/PDLToPDLInterp/PDLToPDLInterp.h" -#include "../PassDetail.h" + #include "PredicateTree.h" #include "mlir/Dialect/PDL/IR/PDL.h" #include "mlir/Dialect/PDL/IR/PDLTypes.h" @@ -20,6 +20,11 @@ #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTPDLTOPDLINTERP +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::pdl_to_pdl_interp; @@ -909,7 +914,7 @@ void PatternLowering::generateOperationResultTypeRewriter( namespace { struct PDLToPDLInterpPass - : public ConvertPDLToPDLInterpBase<PDLToPDLInterpPass> { + : public impl::ConvertPDLToPDLInterpBase<PDLToPDLInterpPass> { void runOnOperation() final; }; } // namespace diff --git a/mlir/lib/Conversion/PassDetail.h b/mlir/lib/Conversion/PassDetail.h deleted file mode 100644 index 530e156024fd..000000000000 --- a/mlir/lib/Conversion/PassDetail.h +++ /dev/null @@ -1,115 +0,0 @@ -//===- PassDetail.h - Conversion 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 CONVERSION_PASSDETAIL_H_ -#define CONVERSION_PASSDETAIL_H_ - -#include "mlir/Pass/Pass.h" - -#include "mlir/Conversion/GPUToROCDL/Runtimes.h" -#include "mlir/IR/BuiltinOps.h" -#include "mlir/IR/FunctionInterfaces.h" - -namespace mlir { -class AffineDialect; - -// Forward declaration from Dialect.h -template <typename ConcreteDialect> -void registerDialect(DialectRegistry ®istry); - -namespace acc { -class OpenACCDialect; -} // namespace acc - -namespace arith { -class ArithmeticDialect; -} // namespace arith - -namespace cf { -class ControlFlowDialect; -} // namespace cf - -namespace complex { -class ComplexDialect; -} // namespace complex - -namespace gpu { -class GPUDialect; -class GPUModuleOp; -} // namespace gpu - -namespace func { -class FuncDialect; -} // namespace func - -namespace linalg { -class LinalgDialect; -} // namespace linalg - -namespace LLVM { -class LLVMDialect; -} // namespace LLVM - -namespace nvgpu { -class NVGPUDialect; -} - -namespace NVVM { -class NVVMDialect; -} // namespace NVVM - -namespace math { -class MathDialect; -} // namespace math - -namespace memref { -class MemRefDialect; -} // namespace memref - -namespace omp { -class OpenMPDialect; -} // namespace omp - -namespace pdl_interp { -class PDLInterpDialect; -} // namespace pdl_interp - -namespace ROCDL { -class ROCDLDialect; -} // namespace ROCDL - -namespace scf { -class SCFDialect; -} // namespace scf - -namespace spirv { -class SPIRVDialect; -} // namespace spirv - -namespace tensor { -class TensorDialect; -} // namespace tensor - -namespace tosa { -class TosaDialect; -} // namespace tosa - -namespace vector { -class VectorDialect; -} // namespace vector - -namespace arm_neon { -class ArmNeonDialect; -} // namespace arm_neon - -#define GEN_PASS_CLASSES -#include "mlir/Conversion/Passes.h.inc" - -} // namespace mlir - -#endif // CONVERSION_PASSDETAIL_H_ diff --git a/mlir/lib/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.cpp b/mlir/lib/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.cpp index 773e9b267c4f..6e4b6774e300 100644 --- a/mlir/lib/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.cpp +++ b/mlir/lib/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.cpp @@ -7,12 +7,17 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h" -#include "../PassDetail.h" + #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_RECONCILEUNREALIZEDCASTS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -103,7 +108,7 @@ struct UnrealizedConversionCastPassthrough /// Pass to simplify and eliminate unrealized conversion casts. struct ReconcileUnrealizedCasts - : public ReconcileUnrealizedCastsBase<ReconcileUnrealizedCasts> { + : public impl::ReconcileUnrealizedCastsBase<ReconcileUnrealizedCasts> { ReconcileUnrealizedCasts() = default; void runOnOperation() override { diff --git a/mlir/lib/Conversion/SCFToControlFlow/SCFToControlFlow.cpp b/mlir/lib/Conversion/SCFToControlFlow/SCFToControlFlow.cpp index 8bdf68354daa..489db87507cf 100644 --- a/mlir/lib/Conversion/SCFToControlFlow/SCFToControlFlow.cpp +++ b/mlir/lib/Conversion/SCFToControlFlow/SCFToControlFlow.cpp @@ -12,7 +12,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" -#include "../PassDetail.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/SCF/IR/SCF.h" @@ -24,13 +24,18 @@ #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/Passes.h" +namespace mlir { +#define GEN_PASS_DEF_SCFTOCONTROLFLOW +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::scf; namespace { struct SCFToControlFlowPass - : public SCFToControlFlowBase<SCFToControlFlowPass> { + : public impl::SCFToControlFlowBase<SCFToControlFlowPass> { void runOnOperation() override; }; diff --git a/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp b/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp index 42260b5987eb..47c266cd15b1 100644 --- a/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp +++ b/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp @@ -7,18 +7,24 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/SCFToGPU/SCFToGPUPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/SCFToGPU/SCFToGPU.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Complex/IR/Complex.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/SCF/IR/SCF.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" - #include "llvm/ADT/ArrayRef.h" #include "llvm/Support/CommandLine.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTAFFINEFORTOGPU +#define GEN_PASS_DEF_CONVERTPARALLELLOOPTOGPU +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::scf; @@ -26,7 +32,7 @@ namespace { // A pass that traverses top-level loops in the function and converts them to // GPU launch operations. Nested launches are not allowed, so this does not // walk the function recursively to avoid considering nested loops. -struct ForLoopMapper : public ConvertAffineForToGPUBase<ForLoopMapper> { +struct ForLoopMapper : public impl::ConvertAffineForToGPUBase<ForLoopMapper> { ForLoopMapper() = default; ForLoopMapper(unsigned numBlockDims, unsigned numThreadDims) { this->numBlockDims = numBlockDims; @@ -46,7 +52,7 @@ struct ForLoopMapper : public ConvertAffineForToGPUBase<ForLoopMapper> { }; struct ParallelLoopToGpuPass - : public ConvertParallelLoopToGpuBase<ParallelLoopToGpuPass> { + : public impl::ConvertParallelLoopToGpuBase<ParallelLoopToGpuPass> { void runOnOperation() override { RewritePatternSet patterns(&getContext()); populateParallelLoopToGPUPatterns(patterns); diff --git a/mlir/lib/Conversion/SCFToOpenMP/SCFToOpenMP.cpp b/mlir/lib/Conversion/SCFToOpenMP/SCFToOpenMP.cpp index e963f02de528..b534c27fae6b 100644 --- a/mlir/lib/Conversion/SCFToOpenMP/SCFToOpenMP.cpp +++ b/mlir/lib/Conversion/SCFToOpenMP/SCFToOpenMP.cpp @@ -12,7 +12,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/SCFToOpenMP/SCFToOpenMP.h" -#include "../PassDetail.h" + #include "mlir/Analysis/SliceAnalysis.h" #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" @@ -22,8 +22,14 @@ #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/IR/ImplicitLocOpBuilder.h" #include "mlir/IR/SymbolTable.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTSCFTOOPENMP +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; /// Matches a block containing a "simple" reduction. The expected shape of the @@ -443,7 +449,7 @@ static LogicalResult applyPatterns(ModuleOp module) { } /// A pass converting SCF operations to OpenMP operations. -struct SCFToOpenMPPass : public ConvertSCFToOpenMPBase<SCFToOpenMPPass> { +struct SCFToOpenMPPass : public impl::ConvertSCFToOpenMPBase<SCFToOpenMPPass> { /// Pass entry point. void runOnOperation() override { if (failed(applyPatterns(getOperation()))) diff --git a/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRVPass.cpp b/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRVPass.cpp index 1b22fadc4a34..86a8f23b9dd7 100644 --- a/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRVPass.cpp +++ b/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRVPass.cpp @@ -12,7 +12,6 @@ #include "mlir/Conversion/SCFToSPIRV/SCFToSPIRVPass.h" -#include "../PassDetail.h" #include "mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h" #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h" #include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h" @@ -21,10 +20,15 @@ #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +namespace mlir { +#define GEN_PASS_DEF_SCFTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { -struct SCFToSPIRVPass : public SCFToSPIRVBase<SCFToSPIRVPass> { +struct SCFToSPIRVPass : public impl::SCFToSPIRVBase<SCFToSPIRVPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp index f383f9a3995e..c09ad0eb6172 100644 --- a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp +++ b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp @@ -11,7 +11,8 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" +#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h" + #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h" #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" @@ -19,19 +20,23 @@ #include "mlir/Conversion/LLVMCommon/TypeConverter.h" #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" #include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h" -#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/SymbolTable.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" - #include "llvm/ADT/DenseMap.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Support/FormatVariadic.h" +namespace mlir { +#define GEN_PASS_DEF_LOWERHOSTCODETOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; static constexpr const char kSPIRVModule[] = "__spv__"; @@ -271,7 +276,7 @@ class GPULaunchLowering : public ConvertOpToLLVMPattern<gpu::LaunchFuncOp> { }; class LowerHostCodeToLLVM - : public LowerHostCodeToLLVMBase<LowerHostCodeToLLVM> { + : public impl::LowerHostCodeToLLVMBase<LowerHostCodeToLLVM> { public: void runOnOperation() override { ModuleOp module = getOperation(); diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp index 4c0f55815744..28236ce7f3e2 100644 --- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp +++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp @@ -11,18 +11,24 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/LLVMCommon/TypeConverter.h" #include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +#define GEN_PASS_DEF_CONVERTSPIRVTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir using namespace mlir; namespace { /// A pass converting MLIR SPIR-V operations into LLVM dialect. class ConvertSPIRVToLLVMPass - : public ConvertSPIRVToLLVMBase<ConvertSPIRVToLLVMPass> { + : public impl::ConvertSPIRVToLLVMBase<ConvertSPIRVToLLVMPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/ShapeToStandard/ConvertShapeConstraints.cpp b/mlir/lib/Conversion/ShapeToStandard/ConvertShapeConstraints.cpp index b3422767e1af..7df1407da6f9 100644 --- a/mlir/lib/Conversion/ShapeToStandard/ConvertShapeConstraints.cpp +++ b/mlir/lib/Conversion/ShapeToStandard/ConvertShapeConstraints.cpp @@ -8,7 +8,6 @@ #include "mlir/Conversion/ShapeToStandard/ShapeToStandard.h" -#include "../PassDetail.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Shape/IR/Shape.h" @@ -18,7 +17,13 @@ #include "mlir/Pass/PassRegistry.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTSHAPECONSTRAINTS +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; + namespace { #include "ShapeToStandard.cpp.inc" } // namespace @@ -49,7 +54,7 @@ namespace { // is emitted, witnesses are satisfied, so they are replace with // `shape.const_witness true`. class ConvertShapeConstraints - : public ConvertShapeConstraintsBase<ConvertShapeConstraints> { + : public impl::ConvertShapeConstraintsBase<ConvertShapeConstraints> { void runOnOperation() override { auto *func = getOperation(); auto *context = &getContext(); diff --git a/mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp b/mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp index 3a3a336f3ee6..dea878b07cc5 100644 --- a/mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp +++ b/mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp @@ -8,7 +8,6 @@ #include "mlir/Conversion/ShapeToStandard/ShapeToStandard.h" -#include "../PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/SCF/IR/SCF.h" @@ -16,9 +15,15 @@ #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/ImplicitLocOpBuilder.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/STLExtras.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTSHAPETOSTANDARD +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::shape; using namespace mlir::scf; @@ -680,7 +685,7 @@ namespace { namespace { /// Conversion pass. class ConvertShapeToStandardPass - : public ConvertShapeToStandardBase<ConvertShapeToStandardPass> { + : public impl::ConvertShapeToStandardBase<ConvertShapeToStandardPass> { void runOnOperation() override; }; diff --git a/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp b/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp index 8be029e8c5c2..3db90d791079 100644 --- a/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp +++ b/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp @@ -11,18 +11,23 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/TensorToLinalg/TensorToLinalgPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/TensorToLinalg/TensorToLinalg.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTTENSORTOLINALG +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { /// A pass converting MLIR Tensor operations into the Linalg dialect. class ConvertTensorToLinalgPass - : public ConvertTensorToLinalgBase<ConvertTensorToLinalgPass> { + : public impl::ConvertTensorToLinalgBase<ConvertTensorToLinalgPass> { void runOnOperation() override { auto &context = getContext(); ConversionTarget target(context); diff --git a/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRVPass.cpp b/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRVPass.cpp index 3a8ccc8a7b85..4c0dfb62939f 100644 --- a/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRVPass.cpp +++ b/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRVPass.cpp @@ -11,19 +11,24 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/TensorToSPIRV/TensorToSPIRVPass.h" -#include "../PassDetail.h" + #include "mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h" #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h" #include "mlir/Conversion/TensorToSPIRV/TensorToSPIRV.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTTENSORTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { /// A pass converting MLIR Tensor operations into the SPIR-V dialect. class ConvertTensorToSPIRVPass - : public ConvertTensorToSPIRVBase<ConvertTensorToSPIRVPass> { + : public impl::ConvertTensorToSPIRVBase<ConvertTensorToSPIRVPass> { void runOnOperation() override { MLIRContext *context = &getContext(); Operation *op = getOperation(); diff --git a/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp b/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp index a4e4de6ef7a8..788e2846d451 100644 --- a/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp +++ b/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp @@ -10,22 +10,26 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/TosaToArith/TosaToArith.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Pass/PassManager.h" #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_TOSATOARITH +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace tosa; namespace { -struct TosaToArith : public TosaToArithBase<TosaToArith> { +struct TosaToArith : public impl::TosaToArithBase<TosaToArith> { public: void runOnOperation() override { RewritePatternSet patterns(&getContext()); diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp index 0743f5702ee7..51f2ebca546b 100644 --- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp +++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp @@ -10,15 +10,15 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/TosaToLinalg/TosaToLinalg.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Math/IR/Math.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/Dialect/Tosa/Utils/QuantUtils.h" #include "mlir/IR/PatternMatch.h" @@ -26,10 +26,16 @@ #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_TOSATOLINALGNAMED +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { -struct TosaToLinalgNamed : public TosaToLinalgNamedBase<TosaToLinalgNamed> { +struct TosaToLinalgNamed + : public impl::TosaToLinalgNamedBase<TosaToLinalgNamed> { public: void getDependentDialects(DialectRegistry ®istry) const override { registry diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgPass.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgPass.cpp index f8f4a8f01b19..bcaad5589071 100644 --- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgPass.cpp +++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgPass.cpp @@ -10,15 +10,15 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/TosaToLinalg/TosaToLinalg.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Math/IR/Math.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/Dialect/Tosa/Utils/QuantUtils.h" #include "mlir/IR/PatternMatch.h" @@ -27,10 +27,15 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/Passes.h" +namespace mlir { +#define GEN_PASS_DEF_TOSATOLINALG +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { -struct TosaToLinalg : public TosaToLinalgBase<TosaToLinalg> { +struct TosaToLinalg : public impl::TosaToLinalgBase<TosaToLinalg> { public: void getDependentDialects(DialectRegistry ®istry) const override { registry diff --git a/mlir/lib/Conversion/TosaToSCF/TosaToSCFPass.cpp b/mlir/lib/Conversion/TosaToSCF/TosaToSCFPass.cpp index c86216a0c43a..759b730556d7 100644 --- a/mlir/lib/Conversion/TosaToSCF/TosaToSCFPass.cpp +++ b/mlir/lib/Conversion/TosaToSCF/TosaToSCFPass.cpp @@ -10,23 +10,28 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/TosaToSCF/TosaToSCF.h" + +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Pass/PassManager.h" #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_TOSATOSCF +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace tosa; namespace { -struct TosaToSCF : public TosaToSCFBase<TosaToSCF> { +struct TosaToSCF : public impl::TosaToSCFBase<TosaToSCF> { public: void runOnOperation() override { RewritePatternSet patterns(&getContext()); diff --git a/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp b/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp index 08d5c7d50640..cc129f99c340 100644 --- a/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp +++ b/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp @@ -10,23 +10,27 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/TosaToTensor/TosaToTensor.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Pass/PassManager.h" #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_TOSATOTENSOR +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace tosa; namespace { -struct TosaToTensor : public TosaToTensorBase<TosaToTensor> { +struct TosaToTensor : public impl::TosaToTensorBase<TosaToTensor> { public: void runOnOperation() override { RewritePatternSet patterns(&getContext()); diff --git a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp index 26758825bb01..a1708dfd3b56 100644 --- a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp +++ b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp @@ -15,7 +15,6 @@ #include "NvGpuSupport.h" #include "mlir/Conversion/VectorToGPU/VectorToGPU.h" -#include "../PassDetail.h" #include "mlir/Analysis/SliceAnalysis.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" @@ -31,6 +30,11 @@ #include "mlir/Transforms/Passes.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTVECTORTOGPU +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; /// For a vector TransferOpType `xferOp`, an empty `indices` vector, and an @@ -882,7 +886,7 @@ LogicalResult mlir::convertVectorToNVVMCompatibleMMASync(Operation *rootOp) { namespace { struct ConvertVectorToGPUPass - : public ConvertVectorToGPUBase<ConvertVectorToGPUPass> { + : public impl::ConvertVectorToGPUBase<ConvertVectorToGPUPass> { explicit ConvertVectorToGPUPass(bool useNvGpu_) { useNvGpu.setValue(useNvGpu_); diff --git a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.cpp b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.cpp index 3493e2c75165..d74a3cfbaf88 100644 --- a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.cpp +++ b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.cpp @@ -8,8 +8,6 @@ #include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h" -#include "../PassDetail.h" - #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/TypeConverter.h" #include "mlir/Dialect/AMX/AMXDialect.h" @@ -24,14 +22,20 @@ #include "mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h" #include "mlir/Dialect/X86Vector/Transforms.h" #include "mlir/Dialect/X86Vector/X86VectorDialect.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTVECTORTOLLVM +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::vector; namespace { struct LowerVectorToLLVMPass - : public ConvertVectorToLLVMBase<LowerVectorToLLVMPass> { + : public impl::ConvertVectorToLLVMBase<LowerVectorToLLVMPass> { LowerVectorToLLVMPass(const LowerVectorToLLVMOptions &options) { this->reassociateFPReductions = options.reassociateFPReductions; this->force32BitVectorIndices = options.force32BitVectorIndices; diff --git a/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp b/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp index 6bcceb6ed08a..15b0cd628cb4 100644 --- a/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp +++ b/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp @@ -14,7 +14,6 @@ #include "mlir/Conversion/VectorToSCF/VectorToSCF.h" -#include "../PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -26,6 +25,11 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/Passes.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTVECTORTOSCF +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; using vector::TransferReadOp; using vector::TransferWriteOp; @@ -1281,7 +1285,7 @@ void mlir::populateVectorToSCFConversionPatterns( namespace { struct ConvertVectorToSCFPass - : public ConvertVectorToSCFBase<ConvertVectorToSCFPass> { + : public impl::ConvertVectorToSCFBase<ConvertVectorToSCFPass> { ConvertVectorToSCFPass() = default; ConvertVectorToSCFPass(const VectorTransferToSCFOptions &options) { this->fullUnroll = options.unroll; diff --git a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp index 04bc41ebb0aa..483619ba708d 100644 --- a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp +++ b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp @@ -12,7 +12,6 @@ #include "mlir/Conversion/VectorToSPIRV/VectorToSPIRV.h" -#include "../PassDetail.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h" diff --git a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp index d3585cad4897..1ab62428d0ad 100644 --- a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp +++ b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp @@ -12,18 +12,22 @@ #include "mlir/Conversion/VectorToSPIRV/VectorToSPIRVPass.h" -#include "../PassDetail.h" #include "mlir/Conversion/VectorToSPIRV/VectorToSPIRV.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTVECTORTOSPIRV +#include "mlir/Conversion/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { struct ConvertVectorToSPIRVPass - : public ConvertVectorToSPIRVBase<ConvertVectorToSPIRVPass> { + : public impl::ConvertVectorToSPIRVBase<ConvertVectorToSPIRVPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp index 303498b8e2aa..815f002b97e4 100644 --- a/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp @@ -19,12 +19,13 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/Utils.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" -#include "mlir/Dialect/Affine/Passes.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/MapVector.h" @@ -32,6 +33,11 @@ #include "llvm/Support/Debug.h" #include <algorithm> +namespace mlir { +#define GEN_PASS_DEF_AFFINEDATACOPYGENERATION +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "affine-data-copy-generate" using namespace mlir; @@ -50,7 +56,7 @@ namespace { // TODO: We currently can't generate copies correctly when stores // are strided. Check for strided stores. struct AffineDataCopyGeneration - : public AffineDataCopyGenerationBase<AffineDataCopyGeneration> { + : public impl::AffineDataCopyGenerationBase<AffineDataCopyGeneration> { AffineDataCopyGeneration() = default; explicit AffineDataCopyGeneration(unsigned slowMemorySpace, unsigned fastMemorySpace, diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineExpandIndexOps.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineExpandIndexOps.cpp index c162aa2f2d05..0e30b8b22cc1 100644 --- a/mlir/lib/Dialect/Affine/Transforms/AffineExpandIndexOps.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineExpandIndexOps.cpp @@ -10,13 +10,17 @@ // fundamental operations. //===----------------------------------------------------------------------===// -#include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Passes.h" -#include "PassDetail.h" +#include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_AFFINEEXPANDINDEXOPS +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -38,7 +42,7 @@ struct LowerDelinearizeIndexOps }; class ExpandAffineIndexOpsPass - : public AffineExpandIndexOpsBase<ExpandAffineIndexOpsPass> { + : public impl::AffineExpandIndexOpsBase<ExpandAffineIndexOpsPass> { public: ExpandAffineIndexOpsPass() = default; diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineLoopInvariantCodeMotion.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineLoopInvariantCodeMotion.cpp index 280e928d56de..9794ad63f90f 100644 --- a/mlir/lib/Dialect/Affine/Transforms/AffineLoopInvariantCodeMotion.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineLoopInvariantCodeMotion.cpp @@ -10,7 +10,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Analysis/SliceAnalysis.h" #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/Analysis/AffineStructures.h" @@ -18,9 +19,9 @@ #include "mlir/Dialect/Affine/Analysis/Utils.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" -#include "mlir/Dialect/Affine/Passes.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/Builders.h" @@ -32,6 +33,11 @@ #include "llvm/Support/Debug.h" #include "llvm/Support/raw_ostream.h" +namespace mlir { +#define GEN_PASS_DEF_AFFINELOOPINVARIANTCODEMOTION +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "licm" using namespace mlir; @@ -44,7 +50,7 @@ namespace { /// TODO: This code should be removed once the new LICM pass can handle its /// uses. struct LoopInvariantCodeMotion - : public AffineLoopInvariantCodeMotionBase<LoopInvariantCodeMotion> { + : public impl::AffineLoopInvariantCodeMotionBase<LoopInvariantCodeMotion> { void runOnOperation() override; void runOnAffineForOp(AffineForOp forOp); }; diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp index 28b6b9976bd6..93d7d17d4d9e 100644 --- a/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp @@ -10,10 +10,16 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Passes.h" + +#include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" + +namespace mlir { +#define GEN_PASS_DEF_AFFINELOOPNORMALIZE +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir using namespace mlir; @@ -23,7 +29,7 @@ namespace { /// As currently implemented, this pass cannot fail, but it might skip over ops /// that are already in a normalized form. struct AffineLoopNormalizePass - : public AffineLoopNormalizeBase<AffineLoopNormalizePass> { + : public impl::AffineLoopNormalizeBase<AffineLoopNormalizePass> { void runOnOperation() override { getOperation().walk([](Operation *op) { diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineParallelize.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineParallelize.cpp index dbd55411c0bf..0e0a1a21d7f0 100644 --- a/mlir/lib/Dialect/Affine/Transforms/AffineParallelize.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineParallelize.cpp @@ -11,7 +11,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/Analysis/AffineStructures.h" #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" @@ -19,19 +20,25 @@ #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/IR/AffineValueMap.h" #include "mlir/Dialect/Affine/LoopUtils.h" -#include "mlir/Dialect/Affine/Passes.h" #include "mlir/Dialect/Affine/Passes.h.inc" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "llvm/Support/Debug.h" #include <deque> +namespace mlir { +#define GEN_PASS_DEF_AFFINEPARALLELIZE +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "affine-parallel" using namespace mlir; namespace { /// Convert all parallel affine.for op into 1-D affine.parallel op. -struct AffineParallelize : public AffineParallelizeBase<AffineParallelize> { +struct AffineParallelize + : public impl::AffineParallelizeBase<AffineParallelize> { void runOnOperation() override; }; diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp index efc1676378df..cef2948828e4 100644 --- a/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp @@ -16,19 +16,24 @@ #include "mlir/Dialect/Affine/Passes.h" -#include "PassDetail.h" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/Dominance.h" #include "mlir/Support/LogicalResult.h" #include <algorithm> +namespace mlir { +#define GEN_PASS_DEF_AFFINESCALARREPLACEMENT +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "affine-scalrep" using namespace mlir; namespace { struct AffineScalarReplacement - : public AffineScalarReplacementBase<AffineScalarReplacement> { + : public impl::AffineScalarReplacementBase<AffineScalarReplacement> { void runOnOperation() override; }; diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp index 765bc05a13a9..a451840c1ba1 100644 --- a/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp @@ -6,23 +6,31 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/SCF/Utils/Utils.h" #include "mlir/Transforms/Passes.h" #include "mlir/Transforms/RegionUtils.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_LOOPCOALESCING +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define PASS_NAME "loop-coalescing" #define DEBUG_TYPE PASS_NAME using namespace mlir; namespace { -struct LoopCoalescingPass : public LoopCoalescingBase<LoopCoalescingPass> { +struct LoopCoalescingPass + : public impl::LoopCoalescingBase<LoopCoalescingPass> { /// Walk either an scf.for or an affine.for to find a band to coalesce. template <typename LoopOpTy> diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopFusion.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopFusion.cpp index 81b4e4d65fc4..efa1c96a0693 100644 --- a/mlir/lib/Dialect/Affine/Transforms/LoopFusion.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopFusion.cpp @@ -10,7 +10,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/Analysis/AffineStructures.h" #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" @@ -19,6 +20,7 @@ #include "mlir/Dialect/Affine/LoopFusionUtils.h" #include "mlir/Dialect/Affine/LoopUtils.h" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" @@ -33,6 +35,12 @@ #include "llvm/Support/raw_ostream.h" #include <iomanip> #include <sstream> + +namespace mlir { +#define GEN_PASS_DEF_AFFINELOOPFUSION +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "affine-loop-fusion" using namespace mlir; @@ -47,7 +55,7 @@ namespace { // TODO: Extend this pass to check for fusion preventing dependences, // and add support for more general loop fusion algorithms. -struct LoopFusion : public AffineLoopFusionBase<LoopFusion> { +struct LoopFusion : public impl::AffineLoopFusionBase<LoopFusion> { LoopFusion() = default; LoopFusion(unsigned fastMemorySpace, uint64_t localBufSizeThresholdBytes, bool maximalFusion, enum FusionMode affineFusionMode) { diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopTiling.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopTiling.cpp index 3cdae8ab0265..3b1b48e31aee 100644 --- a/mlir/lib/Dialect/Affine/Transforms/LoopTiling.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopTiling.cpp @@ -10,7 +10,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/Analysis/AffineStructures.h" #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" @@ -18,12 +19,18 @@ #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/IR/AffineValueMap.h" #include "mlir/Dialect/Affine/LoopUtils.h" -#include "mlir/Dialect/Affine/Passes.h" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/Builders.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" + +namespace mlir { +#define GEN_PASS_DEF_AFFINELOOPTILING +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + using namespace mlir; #define DEBUG_TYPE "affine-loop-tile" @@ -31,7 +38,7 @@ using namespace mlir; namespace { /// A pass to perform loop tiling on all suitable loop nests of a Function. -struct LoopTiling : public AffineLoopTilingBase<LoopTiling> { +struct LoopTiling : public impl::AffineLoopTilingBase<LoopTiling> { LoopTiling() = default; explicit LoopTiling(uint64_t cacheSizeBytes, bool avoidMaxMinBounds = true) : avoidMaxMinBounds(avoidMaxMinBounds) { diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp index 8f7cae507a97..55a04ca2f0a5 100644 --- a/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp @@ -9,11 +9,13 @@ // This file implements loop unrolling. // //===----------------------------------------------------------------------===// -#include "PassDetail.h" + +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" -#include "mlir/Dialect/Affine/Passes.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/Builders.h" @@ -21,10 +23,15 @@ #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" -using namespace mlir; +namespace mlir { +#define GEN_PASS_DEF_AFFINELOOPUNROLL +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir #define DEBUG_TYPE "affine-loop-unroll" +using namespace mlir; + namespace { // TODO: this is really a test pass and should be moved out of dialect @@ -34,7 +41,7 @@ namespace { /// full unroll threshold was specified, in which case, fully unrolls all loops /// with trip count less than the specified threshold. The latter is for testing /// purposes, especially for testing outer loop unrolling. -struct LoopUnroll : public AffineLoopUnrollBase<LoopUnroll> { +struct LoopUnroll : public impl::AffineLoopUnrollBase<LoopUnroll> { // Callback to obtain unroll factors; if this has a callable target, takes // precedence over command-line argument or passed argument. const std::function<unsigned(AffineForOp)> getUnrollFactor; diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp index 42f6487e8eeb..be43651c38e7 100644 --- a/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp @@ -33,12 +33,13 @@ // op's, bodies of those loops will not be jammed. //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" -#include "mlir/Dialect/Affine/Passes.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/BlockAndValueMapping.h" @@ -46,14 +47,20 @@ #include "llvm/ADT/DenseMap.h" #include "llvm/Support/CommandLine.h" -using namespace mlir; +namespace mlir { +#define GEN_PASS_DEF_AFFINELOOPUNROLLANDJAM +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir #define DEBUG_TYPE "affine-loop-unroll-jam" +using namespace mlir; + namespace { /// Loop unroll jam pass. Currently, this just unroll jams the first /// outer loop in a Function. -struct LoopUnrollAndJam : public AffineLoopUnrollAndJamBase<LoopUnrollAndJam> { +struct LoopUnrollAndJam + : public impl::AffineLoopUnrollAndJamBase<LoopUnrollAndJam> { explicit LoopUnrollAndJam(Optional<unsigned> unrollJamFactor = None) { if (unrollJamFactor) this->unrollJamFactor = *unrollJamFactor; diff --git a/mlir/lib/Dialect/Affine/Transforms/PassDetail.h b/mlir/lib/Dialect/Affine/Transforms/PassDetail.h deleted file mode 100644 index 37dccb132f44..000000000000 --- a/mlir/lib/Dialect/Affine/Transforms/PassDetail.h +++ /dev/null @@ -1,42 +0,0 @@ -//===- PassDetail.h - Affine 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 DIALECT_AFFINE_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_AFFINE_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Affine/Passes.h" -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { -// Forward declaration from Dialect.h -template <typename ConcreteDialect> -void registerDialect(DialectRegistry ®istry); - -namespace arith { -class ArithmeticDialect; -} // namespace arith - -namespace linalg { -class LinalgDialect; -} // namespace linalg - -namespace memref { -class MemRefDialect; -} // namespace memref - -namespace vector { -class VectorDialect; -} // namespace vector - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Affine/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_AFFINE_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/Affine/Transforms/PipelineDataTransfer.cpp b/mlir/lib/Dialect/Affine/Transforms/PipelineDataTransfer.cpp index ec30c296cb8b..5c2046a03933 100644 --- a/mlir/lib/Dialect/Affine/Transforms/PipelineDataTransfer.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/PipelineDataTransfer.cpp @@ -10,7 +10,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" #include "mlir/Dialect/Affine/Analysis/Utils.h" @@ -18,19 +19,25 @@ #include "mlir/Dialect/Affine/LoopUtils.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Dialect/Arithmetic/Utils/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/Builders.h" #include "mlir/Transforms/Passes.h" #include "llvm/ADT/DenseMap.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_AFFINEPIPELINEDATATRANSFER +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "affine-pipeline-data-transfer" using namespace mlir; namespace { struct PipelineDataTransfer - : public AffinePipelineDataTransferBase<PipelineDataTransfer> { + : public impl::AffinePipelineDataTransferBase<PipelineDataTransfer> { void runOnOperation() override; void runOnAffineForOp(AffineForOp forOp); diff --git a/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp b/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp index cd05d66ea823..bb5b390cb2ea 100644 --- a/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp @@ -10,14 +10,20 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Dialect/Affine/Analysis/Utils.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" -#include "mlir/Dialect/Affine/Passes.h" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/IntegerSet.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_SIMPLIFYAFFINESTRUCTURES +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "simplify-affine-structure" using namespace mlir; @@ -29,7 +35,7 @@ namespace { /// all memrefs with non-trivial layout maps are converted to ones with trivial /// identity layout ones. struct SimplifyAffineStructures - : public SimplifyAffineStructuresBase<SimplifyAffineStructures> { + : public impl::SimplifyAffineStructuresBase<SimplifyAffineStructures> { void runOnOperation() override; /// Utility to simplify an affine attribute and update its entry in the parent diff --git a/mlir/lib/Dialect/Affine/Transforms/SuperVectorize.cpp b/mlir/lib/Dialect/Affine/Transforms/SuperVectorize.cpp index 749ac8dccc45..e59d5543d949 100644 --- a/mlir/lib/Dialect/Affine/Transforms/SuperVectorize.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/SuperVectorize.cpp @@ -11,7 +11,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Affine/Passes.h" + #include "mlir/Analysis/SliceAnalysis.h" #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h" @@ -19,13 +20,20 @@ #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "mlir/Dialect/Vector/Utils/VectorUtils.h" #include "mlir/IR/BlockAndValueMapping.h" +#include "mlir/Pass/Pass.h" #include "mlir/Support/LLVM.h" #include "llvm/ADT/STLExtras.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_AFFINEVECTORIZE +#include "mlir/Dialect/Affine/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace vector; @@ -607,7 +615,7 @@ namespace { /// Base state for the vectorize pass. /// Command line arguments are preempted by non-empty pass arguments. -struct Vectorize : public AffineVectorizeBase<Vectorize> { +struct Vectorize : public impl::AffineVectorizeBase<Vectorize> { Vectorize() = default; Vectorize(ArrayRef<int64_t> virtualVectorSize); void runOnOperation() override; diff --git a/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp index 60c644e0e8c3..72bf79fd86f2 100644 --- a/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp @@ -6,23 +6,29 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Arithmetic/Transforms/Passes.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Arithmetic/Transforms/BufferizableOpInterfaceImpl.h" -#include "mlir/Dialect/Arithmetic/Transforms/Passes.h" #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" +namespace mlir { +namespace arith { +#define GEN_PASS_DEF_ARITHMETICBUFFERIZE +#include "mlir/Dialect/Arithmetic/Transforms/Passes.h.inc" +} // namespace arith +} // namespace mlir + using namespace mlir; using namespace bufferization; namespace { /// Pass to bufferize Arithmetic ops. struct ArithmeticBufferizePass - : public ArithmeticBufferizeBase<ArithmeticBufferizePass> { + : public arith::impl::ArithmeticBufferizeBase<ArithmeticBufferizePass> { ArithmeticBufferizePass(uint64_t alignment = 0, bool constantOpOnly = false) : constantOpOnly(constantOpOnly) { this->alignment = alignment; diff --git a/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp b/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp index afe7aab99af3..e7e6be9b6906 100644 --- a/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp +++ b/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp @@ -6,12 +6,19 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Arithmetic/Transforms/Passes.h" + +#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/IR/TypeUtilities.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +namespace arith { +#define GEN_PASS_DEF_ARITHMETICEXPANDOPS +#include "mlir/Dialect/Arithmetic/Transforms/Passes.h.inc" +} // namespace arith +} // namespace mlir + using namespace mlir; /// Create an integer or index constant. @@ -189,7 +196,7 @@ public: }; struct ArithmeticExpandOpsPass - : public ArithmeticExpandOpsBase<ArithmeticExpandOpsPass> { + : public arith::impl::ArithmeticExpandOpsBase<ArithmeticExpandOpsPass> { void runOnOperation() override { RewritePatternSet patterns(&getContext()); ConversionTarget target(getContext()); diff --git a/mlir/lib/Dialect/Arithmetic/Transforms/PassDetail.h b/mlir/lib/Dialect/Arithmetic/Transforms/PassDetail.h deleted file mode 100644 index 5877ccccd634..000000000000 --- a/mlir/lib/Dialect/Arithmetic/Transforms/PassDetail.h +++ /dev/null @@ -1,34 +0,0 @@ -//===- PassDetail.h - Arithmetic Pass 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 DIALECT_ARITHMETIC_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_ARITHMETIC_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/IR/BuiltinOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace func { -class FuncDialect; -} // namespace func - -namespace memref { -class MemRefDialect; -} // namespace memref - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Arithmetic/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_ARITHMETIC_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/Arithmetic/Transforms/UnsignedWhenEquivalent.cpp b/mlir/lib/Dialect/Arithmetic/Transforms/UnsignedWhenEquivalent.cpp index 82e442851b31..14a028e9ba66 100644 --- a/mlir/lib/Dialect/Arithmetic/Transforms/UnsignedWhenEquivalent.cpp +++ b/mlir/lib/Dialect/Arithmetic/Transforms/UnsignedWhenEquivalent.cpp @@ -8,13 +8,20 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Arithmetic/Transforms/Passes.h" + #include "mlir/Analysis/DataFlow/DeadCodeAnalysis.h" #include "mlir/Analysis/DataFlow/IntegerRangeAnalysis.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" -#include "mlir/Dialect/Arithmetic/Transforms/Passes.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +namespace arith { +#define GEN_PASS_DEF_ARITHMETICUNSIGNEDWHENEQUIVALENT +#include "mlir/Dialect/Arithmetic/Transforms/Passes.h.inc" +} // namespace arith +} // namespace mlir + using namespace mlir; using namespace mlir::arith; using namespace mlir::dataflow; @@ -102,7 +109,7 @@ struct ConvertCmpIToUnsigned : OpConversionPattern<CmpIOp> { }; struct ArithmeticUnsignedWhenEquivalentPass - : public ArithmeticUnsignedWhenEquivalentBase< + : public arith::impl::ArithmeticUnsignedWhenEquivalentBase< ArithmeticUnsignedWhenEquivalentPass> { /// Implementation structure: first find all equivalent ops and collect them, /// then perform all the rewrites in a second pass over the target op. This diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp index c50baf7e6d74..99314c23015d 100644 --- a/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp @@ -10,12 +10,11 @@ // //===----------------------------------------------------------------------===// -#include <utility> +#include "mlir/Dialect/Async/Passes.h" #include "PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Async/IR/Async.h" -#include "mlir/Dialect/Async/Passes.h" #include "mlir/Dialect/Async/Transforms.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/SCF/IR/SCF.h" @@ -26,6 +25,12 @@ #include "mlir/Support/LLVM.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/RegionUtils.h" +#include <utility> + +namespace mlir { +#define GEN_PASS_DEF_ASYNCPARALLELFOR +#include "mlir/Dialect/Async/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::async; @@ -94,7 +99,7 @@ namespace { // } // struct AsyncParallelForPass - : public AsyncParallelForBase<AsyncParallelForPass> { + : public impl::AsyncParallelForBase<AsyncParallelForPass> { AsyncParallelForPass() = default; AsyncParallelForPass(bool asyncDispatch, int32_t numWorkerThreads, diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCounting.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCounting.cpp index 959d6339495e..7db078ad3f0a 100644 --- a/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCounting.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCounting.cpp @@ -11,10 +11,10 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Async/Passes.h" + #include "mlir/Analysis/Liveness.h" #include "mlir/Dialect/Async/IR/Async.h" -#include "mlir/Dialect/Async/Passes.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/ImplicitLocOpBuilder.h" @@ -22,11 +22,17 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/SmallSet.h" -using namespace mlir; -using namespace mlir::async; +namespace mlir { +#define GEN_PASS_DEF_ASYNCRUNTIMEREFCOUNTING +#define GEN_PASS_DEF_ASYNCRUNTIMEPOLICYBASEDREFCOUNTING +#include "mlir/Dialect/Async/Passes.h.inc" +} // namespace mlir #define DEBUG_TYPE "async-runtime-ref-counting" +using namespace mlir; +using namespace mlir::async; + //===----------------------------------------------------------------------===// // Utility functions shared by reference counting passes. //===----------------------------------------------------------------------===// @@ -103,7 +109,7 @@ static LogicalResult walkReferenceCountedValues( namespace { class AsyncRuntimeRefCountingPass - : public AsyncRuntimeRefCountingBase<AsyncRuntimeRefCountingPass> { + : public impl::AsyncRuntimeRefCountingBase<AsyncRuntimeRefCountingPass> { public: AsyncRuntimeRefCountingPass() = default; void runOnOperation() override; @@ -462,7 +468,7 @@ void AsyncRuntimeRefCountingPass::runOnOperation() { namespace { class AsyncRuntimePolicyBasedRefCountingPass - : public AsyncRuntimePolicyBasedRefCountingBase< + : public impl::AsyncRuntimePolicyBasedRefCountingBase< AsyncRuntimePolicyBasedRefCountingPass> { public: AsyncRuntimePolicyBasedRefCountingPass() { initializeDefaultPolicy(); } diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp index 14fe91a90201..5a8a9af6da9c 100644 --- a/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp @@ -10,22 +10,28 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/Async/IR/Async.h" #include "mlir/Dialect/Async/Passes.h" + +#include "mlir/Dialect/Async/IR/Async.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "llvm/ADT/SmallSet.h" #include "llvm/Support/Debug.h" -using namespace mlir; -using namespace mlir::async; +namespace mlir { +#define GEN_PASS_DEF_ASYNCRUNTIMEREFCOUNTINGOPT +#include "mlir/Dialect/Async/Passes.h.inc" +} // namespace mlir #define DEBUG_TYPE "async-ref-counting" +using namespace mlir; +using namespace mlir::async; + namespace { class AsyncRuntimeRefCountingOptPass - : public AsyncRuntimeRefCountingOptBase<AsyncRuntimeRefCountingOptPass> { + : public impl::AsyncRuntimeRefCountingOptBase< + AsyncRuntimeRefCountingOptPass> { public: AsyncRuntimeRefCountingOptPass() = default; void runOnOperation() override; diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp index efdae66969b7..8d45959800a8 100644 --- a/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp @@ -11,11 +11,12 @@ // //===----------------------------------------------------------------------===// +#include "mlir/Dialect/Async/Passes.h" + #include "PassDetail.h" #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Async/IR/Async.h" -#include "mlir/Dialect/Async/Passes.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/SCF/IR/SCF.h" @@ -27,6 +28,11 @@ #include "llvm/ADT/SetVector.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_ASYNCTOASYNCRUNTIME +#include "mlir/Dialect/Async/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::async; @@ -37,7 +43,7 @@ static constexpr const char kAsyncFnPrefix[] = "async_execute_fn"; namespace { class AsyncToAsyncRuntimePass - : public AsyncToAsyncRuntimeBase<AsyncToAsyncRuntimePass> { + : public impl::AsyncToAsyncRuntimeBase<AsyncToAsyncRuntimePass> { public: AsyncToAsyncRuntimePass() = default; void runOnOperation() override; diff --git a/mlir/lib/Dialect/Async/Transforms/PassDetail.h b/mlir/lib/Dialect/Async/Transforms/PassDetail.h index d2d71147fd77..40e3c58722ec 100644 --- a/mlir/lib/Dialect/Async/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Async/Transforms/PassDetail.h @@ -27,9 +27,6 @@ namespace scf { class SCFDialect; } // namespace scf -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Async/Passes.h.inc" - // -------------------------------------------------------------------------- // // Utility functions shared by Async Transformations. // -------------------------------------------------------------------------- // diff --git a/mlir/lib/Dialect/Bufferization/Transforms/AllocTensorElimination.cpp b/mlir/lib/Dialect/Bufferization/Transforms/AllocTensorElimination.cpp index ff308e62e274..4a1d4122a66d 100644 --- a/mlir/lib/Dialect/Bufferization/Transforms/AllocTensorElimination.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/AllocTensorElimination.cpp @@ -6,17 +6,23 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Bufferization/Transforms/Passes.h" #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/AllocTensorElimination.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h" -#include "mlir/Dialect/Bufferization/Transforms/Passes.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/Dominance.h" #include "mlir/Pass/Pass.h" +namespace mlir { +namespace bufferization { +#define GEN_PASS_DEF_ALLOCTENSORELIMINATION +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" +} // namespace bufferization +} // namespace mlir + using namespace mlir; using namespace mlir::bufferization; @@ -234,7 +240,8 @@ mlir::bufferization::insertSliceAnchoredAllocTensorEliminationStep( namespace { struct AllocTensorElimination - : public AllocTensorEliminationBase<AllocTensorElimination> { + : public bufferization::impl::AllocTensorEliminationBase< + AllocTensorElimination> { AllocTensorElimination() = default; void runOnOperation() override; diff --git a/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp b/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp index 2f7ede7adc2e..609236b10d46 100644 --- a/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp @@ -50,15 +50,22 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Bufferization/Transforms/Passes.h" #include "mlir/Dialect/Bufferization/IR/AllocationOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/BufferUtils.h" -#include "mlir/Dialect/Bufferization/Transforms/Passes.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "llvm/ADT/SetOperations.h" +namespace mlir { +namespace bufferization { +#define GEN_PASS_DEF_BUFFERDEALLOCATION +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" +} // namespace bufferization +} // namespace mlir + using namespace mlir; using namespace mlir::bufferization; @@ -633,7 +640,9 @@ struct DefaultAllocationInterface /// The actual buffer deallocation pass that inserts and moves dealloc nodes /// into the right positions. Furthermore, it inserts additional clones if /// necessary. It uses the algorithm described at the top of the file. -struct BufferDeallocationPass : BufferDeallocationBase<BufferDeallocationPass> { +struct BufferDeallocationPass + : public bufferization::impl::BufferDeallocationBase< + BufferDeallocationPass> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<bufferization::BufferizationDialect>(); registry.insert<memref::MemRefDialect>(); diff --git a/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp b/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp index e79e2829b63a..5710d3af8da9 100644 --- a/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp @@ -11,14 +11,24 @@ // allocations and copies during buffer deallocation. The third pass tries to // convert heap-based allocations to stack-based allocations, if possible. -#include "PassDetail.h" -#include "mlir/Dialect/Bufferization/Transforms/BufferUtils.h" #include "mlir/Dialect/Bufferization/Transforms/Passes.h" + +#include "mlir/Dialect/Bufferization/Transforms/BufferUtils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/Operation.h" #include "mlir/Interfaces/LoopLikeInterface.h" #include "mlir/Pass/Pass.h" +namespace mlir { +namespace bufferization { +#define GEN_PASS_DEF_BUFFERHOISTING +#define GEN_PASS_DEF_BUFFERLOOPHOISTING +#define GEN_PASS_DEF_PROMOTEBUFFERSTOSTACK +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" +} // namespace bufferization +} // namespace mlir + using namespace mlir; using namespace mlir::bufferization; @@ -361,7 +371,8 @@ public: /// The buffer hoisting pass that hoists allocation nodes into dominating /// blocks. -struct BufferHoistingPass : BufferHoistingBase<BufferHoistingPass> { +struct BufferHoistingPass + : public bufferization::impl::BufferHoistingBase<BufferHoistingPass> { void runOnOperation() override { // Hoist all allocations into dominator blocks. @@ -372,7 +383,9 @@ struct BufferHoistingPass : BufferHoistingBase<BufferHoistingPass> { }; /// The buffer loop hoisting pass that hoists allocation nodes out of loops. -struct BufferLoopHoistingPass : BufferLoopHoistingBase<BufferLoopHoistingPass> { +struct BufferLoopHoistingPass + : public bufferization::impl::BufferLoopHoistingBase< + BufferLoopHoistingPass> { void runOnOperation() override { // Hoist all allocations out of loops. @@ -385,7 +398,8 @@ struct BufferLoopHoistingPass : BufferLoopHoistingBase<BufferLoopHoistingPass> { /// The promote buffer to stack pass that tries to convert alloc nodes into /// alloca nodes. class PromoteBuffersToStackPass - : public PromoteBuffersToStackBase<PromoteBuffersToStackPass> { + : public bufferization::impl::PromoteBuffersToStackBase< + PromoteBuffersToStackPass> { public: PromoteBuffersToStackPass(unsigned maxAllocSizeInBytes, unsigned maxRankOfAllocatedMemRef) { diff --git a/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp b/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp index d00da763498a..996e7b729c37 100644 --- a/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp @@ -6,13 +6,20 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Bufferization/Transforms/Passes.h" + #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/Operation.h" #include "mlir/Pass/Pass.h" +namespace mlir { +namespace bufferization { +#define GEN_PASS_DEF_BUFFERRESULTSTOOUTPARAMS +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" +} // namespace bufferization +} // namespace mlir + using namespace mlir; /// Return `true` if the given MemRef type has a fully dynamic layout. @@ -179,7 +186,8 @@ mlir::bufferization::promoteBufferResultsToOutParams(ModuleOp module) { namespace { struct BufferResultsToOutParamsPass - : BufferResultsToOutParamsBase<BufferResultsToOutParamsPass> { + : bufferization::impl::BufferResultsToOutParamsBase< + BufferResultsToOutParamsPass> { void runOnOperation() override { if (failed(bufferization::promoteBufferResultsToOutParams(getOperation()))) return signalPassFailure(); diff --git a/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp index a233367b38bf..1a328d871235 100644 --- a/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp @@ -6,14 +6,13 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Bufferization/Transforms/Passes.h" #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotModuleBufferize.h" -#include "mlir/Dialect/Bufferization/Transforms/Passes.h" #include "mlir/Dialect/Bufferization/Transforms/TensorCopyInsertion.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -22,6 +21,15 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/Passes.h" +namespace mlir { +namespace bufferization { +#define GEN_PASS_DEF_FINALIZINGBUFFERIZE +#define GEN_PASS_DEF_BUFFERIZATIONBUFFERIZE +#define GEN_PASS_DEF_ONESHOTBUFFERIZE +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" +} // namespace bufferization +} // namespace mlir + using namespace mlir; using namespace mlir::bufferization; @@ -122,7 +130,8 @@ void mlir::bufferization::populateEliminateBufferizeMaterializationsPatterns( namespace { struct FinalizingBufferizePass - : public FinalizingBufferizeBase<FinalizingBufferizePass> { + : public bufferization::impl::FinalizingBufferizeBase< + FinalizingBufferizePass> { using FinalizingBufferizeBase< FinalizingBufferizePass>::FinalizingBufferizeBase; @@ -164,7 +173,7 @@ parseLayoutMapOption(const std::string &s) { } struct OneShotBufferizePass - : public OneShotBufferizeBase<OneShotBufferizePass> { + : public bufferization::impl::OneShotBufferizeBase<OneShotBufferizePass> { OneShotBufferizePass() {} explicit OneShotBufferizePass(const OneShotBufferizationOptions &options) @@ -255,7 +264,8 @@ private: namespace { struct BufferizationBufferizePass - : public BufferizationBufferizeBase<BufferizationBufferizePass> { + : public bufferization::impl::BufferizationBufferizeBase< + BufferizationBufferizePass> { void runOnOperation() override { BufferizationOptions options = getPartialBufferizationOptions(); options.opFilter.allowDialect<BufferizationDialect>(); diff --git a/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp b/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp index fdbc9a042679..5cbd125b9e80 100644 --- a/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp @@ -27,13 +27,20 @@ // function argument, it is also considered equivalent. A cast is inserted at // the call site in that case. -#include "PassDetail.h" #include "mlir/Dialect/Bufferization/Transforms/Passes.h" + #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/Operation.h" #include "mlir/Pass/Pass.h" +namespace mlir { +namespace bufferization { +#define GEN_PASS_DEF_DROPEQUIVALENTBUFFERRESULTS +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" +} // namespace bufferization +} // namespace mlir + using namespace mlir; /// Return the unique ReturnOp that terminates `funcOp`. @@ -137,7 +144,8 @@ mlir::bufferization::dropEquivalentBufferResults(ModuleOp module) { namespace { struct DropEquivalentBufferResultsPass - : DropEquivalentBufferResultsBase<DropEquivalentBufferResultsPass> { + : bufferization::impl::DropEquivalentBufferResultsBase< + DropEquivalentBufferResultsPass> { void runOnOperation() override { if (failed(bufferization::dropEquivalentBufferResults(getOperation()))) return signalPassFailure(); diff --git a/mlir/lib/Dialect/Bufferization/Transforms/PassDetail.h b/mlir/lib/Dialect/Bufferization/Transforms/PassDetail.h deleted file mode 100644 index c63e85c8f566..000000000000 --- a/mlir/lib/Dialect/Bufferization/Transforms/PassDetail.h +++ /dev/null @@ -1,35 +0,0 @@ -//===- PassDetail.h - Bufferization Pass 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 DIALECT_BUFFERIZATION_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_BUFFERIZATION_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/IR/BuiltinOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace func { -class FuncDialect; -} // namespace func - -namespace memref { -class MemRefDialect; -} // namespace memref - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_BUFFERIZATION_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/Bufferization/Transforms/TensorCopyInsertion.cpp b/mlir/lib/Dialect/Bufferization/Transforms/TensorCopyInsertion.cpp index 8e5aa1cbc807..2b7025c81939 100644 --- a/mlir/lib/Dialect/Bufferization/Transforms/TensorCopyInsertion.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/TensorCopyInsertion.cpp @@ -6,16 +6,22 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Dialect/Bufferization/Transforms/TensorCopyInsertion.h" - -#include "PassDetail.h" +#include "mlir/Dialect/Bufferization/Transforms/Passes.h" #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotModuleBufferize.h" -#include "mlir/Dialect/Bufferization/Transforms/Passes.h" +#include "mlir/Dialect/Bufferization/Transforms/TensorCopyInsertion.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" + +namespace mlir { +namespace bufferization { +#define GEN_PASS_DEF_TENSORCOPYINSERTION +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" +} // namespace bufferization +} // namespace mlir using namespace mlir; using namespace mlir::bufferization; @@ -159,7 +165,8 @@ mlir::bufferization::insertTensorCopies(Operation *op, namespace { struct TensorCopyInsertionPass - : TensorCopyInsertionBase<TensorCopyInsertionPass> { + : public bufferization::impl::TensorCopyInsertionBase< + TensorCopyInsertionPass> { TensorCopyInsertionPass() : options(llvm::None) {} TensorCopyInsertionPass(const OneShotBufferizationOptions &options) : options(options) {} diff --git a/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp b/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp index 327f2dfa187e..5f4fed8e4d49 100644 --- a/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp +++ b/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp @@ -10,20 +10,25 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Func/Transforms/Passes.h" + #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Func/Transforms/FuncConversions.h" -#include "mlir/Dialect/Func/Transforms/Passes.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_FUNCBUFFERIZE +#include "mlir/Dialect/Func/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::func; namespace { -struct FuncBufferizePass : public FuncBufferizeBase<FuncBufferizePass> { +struct FuncBufferizePass : public impl::FuncBufferizeBase<FuncBufferizePass> { using FuncBufferizeBase<FuncBufferizePass>::FuncBufferizeBase; void runOnOperation() override { auto module = getOperation(); diff --git a/mlir/lib/Dialect/Func/Transforms/PassDetail.h b/mlir/lib/Dialect/Func/Transforms/PassDetail.h deleted file mode 100644 index 4b8a2cd8994e..000000000000 --- a/mlir/lib/Dialect/Func/Transforms/PassDetail.h +++ /dev/null @@ -1,33 +0,0 @@ -//===- PassDetail.h - Func 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 DIALECT_FUNC_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_FUNC_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/IR/BuiltinOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -class AtomicRMWOp; - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace memref { -class MemRefDialect; -} // namespace memref - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Func/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_FUNC_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp b/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp index 59758b80357d..92d186b3bd41 100644 --- a/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp +++ b/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp @@ -11,10 +11,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/GPU/Transforms/Passes.h" + #include "mlir/Dialect/Async/IR/Async.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" -#include "mlir/Dialect/GPU/Transforms/Passes.h" #include "mlir/Dialect/GPU/Transforms/Utils.h" #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/Builders.h" @@ -24,9 +25,16 @@ #include "mlir/Transforms/RegionUtils.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +#define GEN_PASS_DEF_GPUASYNCREGIONPASS +#include "mlir/Dialect/GPU/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; + namespace { -class GpuAsyncRegionPass : public GpuAsyncRegionPassBase<GpuAsyncRegionPass> { +class GpuAsyncRegionPass + : public impl::GpuAsyncRegionPassBase<GpuAsyncRegionPass> { struct ThreadTokenCallback; struct DeferWaitCallback; struct SingleTokenUseCallback; diff --git a/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp index df81284edcf8..dd5e9bacde24 100644 --- a/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp +++ b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp @@ -10,14 +10,14 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/GPU/Transforms/Passes.h" + #include "mlir/AsmParser/AsmParser.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/DLTI/DLTI.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" -#include "mlir/Dialect/GPU/Transforms/Passes.h" #include "mlir/Dialect/GPU/Transforms/Utils.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/BlockAndValueMapping.h" @@ -27,6 +27,12 @@ #include "mlir/Support/LLVM.h" #include "mlir/Transforms/RegionUtils.h" +namespace mlir { +#define GEN_PASS_DEF_GPULAUNCHSINKINDEXCOMPUTATIONS +#define GEN_PASS_DEF_GPUKERNELOUTLINING +#include "mlir/Dialect/GPU/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; template <typename OpTy> @@ -239,7 +245,7 @@ namespace { /// Pass that moves ops which are likely an index computation into gpu.launch /// body. class GpuLaunchSinkIndexComputationsPass - : public GpuLaunchSinkIndexComputationsBase< + : public impl::GpuLaunchSinkIndexComputationsBase< GpuLaunchSinkIndexComputationsPass> { public: void runOnOperation() override { @@ -266,7 +272,7 @@ public: /// a separate pass. The external functions can then be annotated with the /// symbol of the cubin accessor function. class GpuKernelOutliningPass - : public GpuKernelOutliningBase<GpuKernelOutliningPass> { + : public impl::GpuKernelOutliningBase<GpuKernelOutliningPass> { public: GpuKernelOutliningPass(StringRef dlStr) { if (!dlStr.empty() && !dataLayoutStr.hasValue()) diff --git a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp index 84f2a907cafb..72e0ebc132e8 100644 --- a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp +++ b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp @@ -11,15 +11,20 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h" +#include "mlir/Dialect/GPU/Transforms/Passes.h" -#include "PassDetail.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" -#include "mlir/Dialect/GPU/Transforms/Passes.h" +#include "mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/IR/AffineMap.h" namespace mlir { +#define GEN_PASS_DEF_GPUMAPPARALLELLOOPSPASS +#include "mlir/Dialect/GPU/Transforms/Passes.h.inc" +} // namespace mlir + +namespace mlir { using scf::ParallelOp; @@ -129,7 +134,7 @@ static void mapParallelOp(ParallelOp parallelOp, namespace { struct GpuMapParallelLoopsPass - : public GpuMapParallelLoopsPassBase<GpuMapParallelLoopsPass> { + : public impl::GpuMapParallelLoopsPassBase<GpuMapParallelLoopsPass> { void runOnOperation() override { for (Region ®ion : getOperation()->getRegions()) { region.walk([](ParallelOp parallelOp) { mapParallelOp(parallelOp); }); diff --git a/mlir/lib/Dialect/GPU/Transforms/PassDetail.h b/mlir/lib/Dialect/GPU/Transforms/PassDetail.h deleted file mode 100644 index 2f5c7dc803ca..000000000000 --- a/mlir/lib/Dialect/GPU/Transforms/PassDetail.h +++ /dev/null @@ -1,25 +0,0 @@ -//===- PassDetail.h - GPU 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 DIALECT_GPU_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_GPU_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Async/IR/Async.h" -#include "mlir/Dialect/DLTI/DLTI.h" -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/Dialect/GPU/IR/GPUDialect.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/GPU/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_GPU_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp index 08fd412ff92c..61c1378d9612 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp @@ -7,11 +7,19 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h" -#include "PassDetail.h" + #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/Block.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinOps.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +namespace LLVM { +#define GEN_PASS_DEF_LLVMLEGALIZEFOREXPORT +#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" +} // namespace LLVM +} // namespace mlir using namespace mlir; @@ -68,7 +76,7 @@ void mlir::LLVM::ensureDistinctSuccessors(Operation *op) { namespace { struct LegalizeForExportPass - : public LLVMLegalizeForExportBase<LegalizeForExportPass> { + : public LLVM::impl::LLVMLegalizeForExportBase<LegalizeForExportPass> { void runOnOperation() override { LLVM::ensureDistinctSuccessors(getOperation()); } diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp index d269aa82ecec..8c33148d1d2d 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp @@ -7,12 +7,20 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h" -#include "PassDetail.h" + #include "mlir/Dialect/LLVMIR/NVVMDialect.h" #include "mlir/IR/Builders.h" #include "mlir/IR/PatternMatch.h" +#include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +namespace NVVM { +#define GEN_PASS_DEF_NVVMOPTIMIZEFORTARGET +#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" +} // namespace NVVM +} // namespace mlir + using namespace mlir; namespace { @@ -32,7 +40,7 @@ private: }; struct NVVMOptimizeForTarget - : public NVVMOptimizeForTargetBase<NVVMOptimizeForTarget> { + : public NVVM::impl::NVVMOptimizeForTargetBase<NVVMOptimizeForTarget> { void runOnOperation() override; void getDependentDialects(DialectRegistry ®istry) const override { diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h b/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h deleted file mode 100644 index deb54ba082a5..000000000000 --- a/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h +++ /dev/null @@ -1,24 +0,0 @@ -//===- PassDetail.h - LLVM 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 DIALECT_LLVMIR_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_LLVMIR_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Pass/Pass.h" - -namespace mlir { -namespace func { -class FuncOp; -} // namespace func - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_LLVMIR_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp index d32f51053e87..099d9c26e558 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp @@ -7,15 +7,22 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h" -#include "PassDetail.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/Pass/Pass.h" + +namespace mlir { +namespace LLVM { +#define GEN_PASS_DEF_LLVMREQUESTCWRAPPERS +#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" +} // namespace LLVM +} // namespace mlir using namespace mlir; namespace { class RequestCWrappersPass - : public LLVMRequestCWrappersBase<RequestCWrappersPass> { + : public LLVM::impl::LLVMRequestCWrappersBase<RequestCWrappersPass> { public: void runOnOperation() override { getOperation()->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), diff --git a/mlir/lib/Dialect/Linalg/Transforms/BubbleUpExtractSlice.cpp b/mlir/lib/Dialect/Linalg/Transforms/BubbleUpExtractSlice.cpp index 78dd1ead3172..c14f91000d89 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/BubbleUpExtractSlice.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/BubbleUpExtractSlice.cpp @@ -12,7 +12,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/Utils/Utils.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" diff --git a/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp index 63433ef50fe2..ca74a17f6cd0 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp @@ -6,12 +6,13 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" +#include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/BufferizableOpInterfaceImpl.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" @@ -19,13 +20,19 @@ #include "mlir/IR/Operation.h" #include "mlir/Pass/Pass.h" +namespace mlir { +#define GEN_PASS_DEF_LINALGBUFFERIZE +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace bufferization; namespace { /// Converts Linalg operations that work on tensor-type operands or results to /// work on buffers. -struct LinalgBufferizePass : public LinalgBufferizeBase<LinalgBufferizePass> { +struct LinalgBufferizePass + : public impl::LinalgBufferizeBase<LinalgBufferizePass> { void runOnOperation() override { BufferizationOptions options = getPartialBufferizationOptions(); options.opFilter.allowDialect<linalg::LinalgDialect>(); diff --git a/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp b/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp index b98b486d38ee..2e975693151f 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp @@ -6,11 +6,12 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" + #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Func/Transforms/FuncConversions.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/OpDefinition.h" #include "mlir/Transforms/DialectConversion.h" @@ -19,6 +20,11 @@ #include <memory> #include <utility> +namespace mlir { +#define GEN_PASS_DEF_LINALGDETENSORIZE +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::linalg; @@ -158,7 +164,8 @@ public: }; /// @see LinalgDetensorize in Linalg/Passes.td for more details. -struct LinalgDetensorize : public LinalgDetensorizeBase<LinalgDetensorize> { +struct LinalgDetensorize + : public impl::LinalgDetensorizeBase<LinalgDetensorize> { LinalgDetensorize() = default; class CostModel { diff --git a/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp b/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp index df3b8e6af043..0d7e41062541 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp @@ -12,10 +12,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" + +#include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/Linalg/Utils/Utils.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" @@ -27,6 +28,11 @@ #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_LINALGFOLDUNITEXTENTDIMS +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "linalg-drop-unit-dims" using namespace mlir; @@ -542,7 +548,7 @@ void mlir::linalg::populateFoldUnitExtentDimsPatterns( namespace { /// Pass that removes unit-extent dims within generic ops. struct LinalgFoldUnitExtentDimsPass - : public LinalgFoldUnitExtentDimsBase<LinalgFoldUnitExtentDimsPass> { + : public impl::LinalgFoldUnitExtentDimsBase<LinalgFoldUnitExtentDimsPass> { void runOnOperation() override { Operation *op = getOperation(); MLIRContext *context = op->getContext(); diff --git a/mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp b/mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp index 2644e50c2684..d63b8f6767f8 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp @@ -9,12 +9,11 @@ // This file implements the linalg dialect Fusion on tensors operations pass. // //===----------------------------------------------------------------------===// -#include <utility> -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/Linalg/Utils/Utils.h" #include "mlir/Dialect/SparseTensor/IR/SparseTensor.h" @@ -24,6 +23,13 @@ #include "mlir/IR/PatternMatch.h" #include "mlir/Support/LLVM.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include <utility> + +namespace mlir { +#define GEN_PASS_DEF_LINALGFOLDUNITEXTENTDIMS +#define GEN_PASS_DEF_LINALGELEMENTWISEOPFUSION +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::linalg; @@ -1753,7 +1759,8 @@ namespace { // favor of test passes that check the functionality of each of the patterns // added here individually. struct LinalgElementwiseOpFusionPass - : public LinalgElementwiseOpFusionBase<LinalgElementwiseOpFusionPass> { + : public impl::LinalgElementwiseOpFusionBase< + LinalgElementwiseOpFusionPass> { void runOnOperation() override { Operation *op = getOperation(); MLIRContext *context = op->getContext(); diff --git a/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp b/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp index 9b5dad05c689..0e115a8ba20f 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp @@ -8,13 +8,17 @@ #include "mlir/Dialect/Linalg/Passes.h" -#include "PassDetail.h" #include "mlir/Dialect/Arithmetic/Utils/Utils.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/Linalg/Utils/Utils.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_CONVERTELEMENTWISETOLINALG +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + using namespace mlir; static bool isElementwiseMappableOpOnRankedTensors(Operation *op) { @@ -121,7 +125,8 @@ void mlir::linalg::populateElementwiseToLinalgConversionPatterns( namespace { class ConvertElementwiseToLinalgPass - : public ConvertElementwiseToLinalgBase<ConvertElementwiseToLinalgPass> { + : public impl::ConvertElementwiseToLinalgBase< + ConvertElementwiseToLinalgPass> { void runOnOperation() final { auto *func = getOperation(); diff --git a/mlir/lib/Dialect/Linalg/Transforms/Fusion.cpp b/mlir/lib/Dialect/Linalg/Transforms/Fusion.cpp index 9d38dd6084fe..ebef65626fd7 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/Fusion.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Fusion.cpp @@ -10,7 +10,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Linalg/Analysis/DependenceAnalysis.h" diff --git a/mlir/lib/Dialect/Linalg/Transforms/FusionOnTensors.cpp b/mlir/lib/Dialect/Linalg/Transforms/FusionOnTensors.cpp index 24caab3b5895..ff28663d479b 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/FusionOnTensors.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/FusionOnTensors.cpp @@ -10,7 +10,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Analysis/SliceAnalysis.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" diff --git a/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp b/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp index 9c5b3fab2240..4f384c33e994 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp @@ -11,9 +11,10 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/Passes.h" + +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/Attributes.h" @@ -24,6 +25,11 @@ #include "llvm/ADT/SmallVector.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_LINALGGENERALIZATION +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "linalg-generalization" using namespace mlir; @@ -67,7 +73,7 @@ FailureOr<GenericOp> mlir::linalg::generalizeNamedOp(RewriterBase &rewriter, namespace { struct LinalgGeneralizationPass - : public LinalgGeneralizationBase<LinalgGeneralizationPass> { + : public impl::LinalgGeneralizationBase<LinalgGeneralizationPass> { void runOnOperation() override; }; diff --git a/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp b/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp index 978220efbf20..518d7376ed2f 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp @@ -6,13 +6,17 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_LINALGINITTENSORTOALLOCTENSOR +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::bufferization; using namespace mlir::linalg; @@ -30,7 +34,8 @@ struct InitTensorLoweringPattern : public OpRewritePattern<InitTensorOp> { }; struct LinalgInitTensorToAllocTensor - : public LinalgInitTensorToAllocTensorBase<LinalgInitTensorToAllocTensor> { + : public impl::LinalgInitTensorToAllocTensorBase< + LinalgInitTensorToAllocTensor> { LinalgInitTensorToAllocTensor() = default; void runOnOperation() override; diff --git a/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp b/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp index b2ad1b24c6b6..e8f5372e4fdc 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp @@ -12,15 +12,21 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_LINALGINLINESCALAROPERANDS +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::linalg; @@ -96,7 +102,8 @@ void mlir::linalg::populateInlineConstantOperandsPatterns( namespace { /// Pass that removes unit-extent dims within generic ops. struct LinalgInlineScalarOperandsPass - : public LinalgInlineScalarOperandsBase<LinalgInlineScalarOperandsPass> { + : public impl::LinalgInlineScalarOperandsBase< + LinalgInlineScalarOperandsPass> { void runOnOperation() override { func::FuncOp funcOp = getOperation(); MLIRContext *context = funcOp.getContext(); diff --git a/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp b/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp index b8a8bbd40a0a..96061dcc621a 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp @@ -11,15 +11,14 @@ // //===----------------------------------------------------------------------===// -#include <utility> +#include "mlir/Dialect/Linalg/Passes.h" -#include "PassDetail.h" #include "mlir/Analysis/SliceAnalysis.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" #include "mlir/Dialect/Affine/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/Hoisting.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/Linalg/Utils/Utils.h" @@ -33,6 +32,20 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/LoopInvariantCodeMotionUtils.h" #include "mlir/Transforms/Passes.h" +#include <utility> + +namespace mlir { +#define GEN_PASS_DEF_LINALGSTRATEGYTILEANDFUSEPASS +#define GEN_PASS_DEF_LINALGSTRATEGYTILEPASS +#define GEN_PASS_DEF_LINALGSTRATEGYPADPASS +#define GEN_PASS_DEF_LINALGSTRATEGYDECOMPOSEPASS +#define GEN_PASS_DEF_LINALGSTRATEGYPEELPASS +#define GEN_PASS_DEF_LINALGSTRATEGYVECTORIZEPASS +#define GEN_PASS_DEF_LINALGSTRATEGYENABLEPASS +#define GEN_PASS_DEF_LINALGSTRATEGYLOWERVECTORSPASS +#define GEN_PASS_DEF_LINALGSTRATEGYREMOVEMARKERSPASS +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::vector; @@ -42,7 +55,8 @@ namespace { /// Configurable pass to apply pattern-based tiling and fusion. struct LinalgStrategyTileAndFusePass - : public LinalgStrategyTileAndFusePassBase<LinalgStrategyTileAndFusePass> { + : public impl::LinalgStrategyTileAndFusePassBase< + LinalgStrategyTileAndFusePass> { LinalgStrategyTileAndFusePass() = default; @@ -79,11 +93,12 @@ struct LinalgStrategyTileAndFusePass /// Configurable pass to apply pattern-based linalg tiling. struct LinalgStrategyTilePass - : public LinalgStrategyTilePassBase<LinalgStrategyTilePass> { + : public impl::LinalgStrategyTilePassBase<LinalgStrategyTilePass> { LinalgStrategyTilePass() = default; - LinalgStrategyTilePass(StringRef opName, LinalgTilingOptions opt, + LinalgStrategyTilePass(StringRef opName, + mlir::linalg::LinalgTilingOptions opt, LinalgTransformationFilter filt) : options(std::move(opt)), filter(std::move(filt)) { this->anchorOpName.setValue(opName.str()); @@ -106,13 +121,13 @@ struct LinalgStrategyTilePass (void)applyPatternsAndFoldGreedily(funcOp, std::move(tilingPattern)); } - LinalgTilingOptions options; + mlir::linalg::LinalgTilingOptions options; LinalgTransformationFilter filter; }; /// Configurable pass to apply hoisting and padding. struct LinalgStrategyPadPass - : public LinalgStrategyPadPassBase<LinalgStrategyPadPass> { + : public impl::LinalgStrategyPadPassBase<LinalgStrategyPadPass> { LinalgStrategyPadPass() = default; @@ -142,11 +157,11 @@ struct LinalgStrategyPadPass LinalgTransformationFilter filter; }; - /// Configurable pass to apply lowering of coarser-grained named linalg ops into /// finer-grained named versions. struct LinalgStrategyDecomposePass - : public LinalgStrategyDecomposePassBase<LinalgStrategyDecomposePass> { + : public impl::LinalgStrategyDecomposePassBase< + LinalgStrategyDecomposePass> { LinalgStrategyDecomposePass() = default; @@ -169,7 +184,7 @@ struct LinalgStrategyDecomposePass /// Configurable pass to apply pattern-based linalg peeling. struct LinalgStrategyPeelPass - : public LinalgStrategyPeelPassBase<LinalgStrategyPeelPass> { + : public impl::LinalgStrategyPeelPassBase<LinalgStrategyPeelPass> { LinalgStrategyPeelPass() = default; @@ -203,7 +218,8 @@ struct LinalgStrategyPeelPass /// Configurable pass to apply pattern-based linalg vectorization. struct LinalgStrategyVectorizePass - : public LinalgStrategyVectorizePassBase<LinalgStrategyVectorizePass> { + : public impl::LinalgStrategyVectorizePassBase< + LinalgStrategyVectorizePass> { LinalgStrategyVectorizePass() = default; @@ -259,7 +275,7 @@ struct LinalgStrategyVectorizePass /// Configurable pass to enable the application of other pattern-based linalg /// passes. struct LinalgStrategyEnablePass - : public LinalgStrategyEnablePassBase<LinalgStrategyEnablePass> { + : public impl::LinalgStrategyEnablePassBase<LinalgStrategyEnablePass> { LinalgStrategyEnablePass(LinalgEnablingOptions opt, LinalgTransformationFilter filt) @@ -309,7 +325,7 @@ struct LinalgStrategyEnablePass /// Configurable pass to lower vector operations. struct LinalgStrategyLowerVectorsPass - : public LinalgStrategyLowerVectorsPassBase< + : public impl::LinalgStrategyLowerVectorsPassBase< LinalgStrategyLowerVectorsPass> { LinalgStrategyLowerVectorsPass(LinalgVectorLoweringOptions opt, @@ -374,7 +390,7 @@ struct LinalgStrategyLowerVectorsPass /// Configurable pass to lower vector operations. struct LinalgStrategyRemoveMarkersPass - : public LinalgStrategyRemoveMarkersPassBase< + : public impl::LinalgStrategyRemoveMarkersPassBase< LinalgStrategyRemoveMarkersPass> { void runOnOperation() override { diff --git a/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp b/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp index 3656bc2a1a06..e73d5aba9ff5 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp @@ -6,11 +6,13 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" + +#include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Arithmetic/Utils/Utils.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/Linalg/Utils/Utils.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" @@ -24,6 +26,13 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +#define GEN_PASS_DEF_LINALGLOWERTOAFFINELOOPS +#define GEN_PASS_DEF_LINALGLOWERTOLOOPS +#define GEN_PASS_DEF_LINALGLOWERTOPARALLELLOOPS +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::linalg; @@ -311,7 +320,7 @@ static void lowerLinalgToLoopsImpl(func::FuncOp funcOp) { } struct LowerToAffineLoops - : public LinalgLowerToAffineLoopsBase<LowerToAffineLoops> { + : public impl::LinalgLowerToAffineLoopsBase<LowerToAffineLoops> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<memref::MemRefDialect>(); } @@ -320,7 +329,7 @@ struct LowerToAffineLoops } }; -struct LowerToLoops : public LinalgLowerToLoopsBase<LowerToLoops> { +struct LowerToLoops : public impl::LinalgLowerToLoopsBase<LowerToLoops> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<memref::MemRefDialect, scf::SCFDialect>(); } @@ -330,7 +339,7 @@ struct LowerToLoops : public LinalgLowerToLoopsBase<LowerToLoops> { }; struct LowerToParallelLoops - : public LinalgLowerToParallelLoopsBase<LowerToParallelLoops> { + : public impl::LinalgLowerToParallelLoopsBase<LowerToParallelLoops> { void runOnOperation() override { lowerLinalgToLoopsImpl<scf::ParallelOp>(getOperation()); } diff --git a/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp b/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp index 1d987ee62c9c..273518f8c382 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp @@ -10,15 +10,20 @@ // canonicalizations of named ops. // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/Linalg/IR/Linalg.h" + #include "mlir/Dialect/Linalg/Passes.h" + +#include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" - #include "llvm/ADT/SmallVector.h" +namespace mlir { +#define GEN_PASS_DEF_LINALGNAMEDOPCONVERSION +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::linalg; @@ -135,7 +140,7 @@ struct SimplifyDepthwiseConvQOp }; struct LinalgNamedOpConversionPass - : public LinalgNamedOpConversionBase<LinalgNamedOpConversionPass> { + : public impl::LinalgNamedOpConversionBase<LinalgNamedOpConversionPass> { LinalgNamedOpConversionPass() = default; LinalgNamedOpConversionPass(const LinalgNamedOpConversionPass &) = default; diff --git a/mlir/lib/Dialect/Linalg/Transforms/PassDetail.h b/mlir/lib/Dialect/Linalg/Transforms/PassDetail.h deleted file mode 100644 index 63f947524cc7..000000000000 --- a/mlir/lib/Dialect/Linalg/Transforms/PassDetail.h +++ /dev/null @@ -1,52 +0,0 @@ -//===- PassDetail.h - Linalg 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 DIALECT_LINALG_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_LINALG_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Affine/IR/AffineOps.h" -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/IR/BuiltinOps.h" -#include "mlir/IR/Dialect.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { -namespace arith { -class ArithmeticDialect; -} // namespace arith - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace linalg { -class LinalgDialect; -} // namespace linalg - -namespace scf { -class SCFDialect; -} // namespace scf - -namespace memref { -class MemRefDialect; -} // namespace memref - -namespace tensor { -class TensorDialect; -} // namespace tensor - -namespace vector { -class VectorDialect; -} // namespace vector - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Linalg/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_LINALG_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/Linalg/Transforms/Promotion.cpp b/mlir/lib/Dialect/Linalg/Transforms/Promotion.cpp index 4c5fbf98cdf0..79ef14e700ac 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/Promotion.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Promotion.cpp @@ -10,7 +10,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Complex/IR/Complex.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" diff --git a/mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp b/mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp index 55ce5ea66ae3..c24f37949a03 100644 --- a/mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp @@ -10,14 +10,13 @@ // //===----------------------------------------------------------------------===// -#include <utility> +#include "mlir/Dialect/Linalg/Passes.h" -#include "PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/Utils/Utils.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" -#include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/Linalg/Utils/Utils.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -28,8 +27,13 @@ #include "mlir/IR/AffineMap.h" #include "mlir/Transforms/FoldUtils.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" - #include "llvm/Support/CommandLine.h" +#include <utility> + +namespace mlir { +#define GEN_PASS_DEF_LINALGTILINGPASS +#include "mlir/Dialect/Linalg/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::linalg; @@ -744,7 +748,7 @@ static void applyExtractSliceOfPadTensorSwapPattern(func::FuncOp funcOp) { } namespace { -struct LinalgTilingPass : public LinalgTilingBase<LinalgTilingPass> { +struct LinalgTilingPass : public impl::LinalgTilingPassBase<LinalgTilingPass> { LinalgTilingPass() = default; LinalgTilingPass(ArrayRef<int64_t> tileSizes, LinalgTilingLoopType loopType) { this->tileSizes = tileSizes; diff --git a/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp b/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp index c4011c55f5a4..0eca196f1089 100644 --- a/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp @@ -12,15 +12,21 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Arithmetic/Transforms/Passes.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" -#include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/IR/TypeUtilities.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +namespace memref { +#define GEN_PASS_DEF_EXPANDOPS +#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc" +} // namespace memref +} // namespace mlir + using namespace mlir; namespace { @@ -119,7 +125,7 @@ public: } }; -struct ExpandOpsPass : public ExpandOpsBase<ExpandOpsPass> { +struct ExpandOpsPass : public memref::impl::ExpandOpsBase<ExpandOpsPass> { void runOnOperation() override { MLIRContext &ctx = getContext(); diff --git a/mlir/lib/Dialect/MemRef/Transforms/FoldMemRefAliasOps.cpp b/mlir/lib/Dialect/MemRef/Transforms/FoldMemRefAliasOps.cpp index 3ce18ec442c0..e7a2ae7464c5 100644 --- a/mlir/lib/Dialect/MemRef/Transforms/FoldMemRefAliasOps.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/FoldMemRefAliasOps.cpp @@ -11,11 +11,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/MemRef/Transforms/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" -#include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Dialect/Utils/IndexingUtils.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "mlir/IR/BuiltinTypes.h" @@ -23,6 +23,13 @@ #include "llvm/ADT/SmallBitVector.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +namespace memref { +#define GEN_PASS_DEF_FOLDMEMREFALIASOPS +#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc" +} // namespace memref +} // namespace mlir + using namespace mlir; //===----------------------------------------------------------------------===// @@ -540,11 +547,8 @@ void memref::populateFoldMemRefAliasOpPatterns(RewritePatternSet &patterns) { namespace { -#define GEN_PASS_CLASSES -#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc" - struct FoldMemRefAliasOpsPass final - : public FoldMemRefAliasOpsBase<FoldMemRefAliasOpsPass> { + : public memref::impl::FoldMemRefAliasOpsBase<FoldMemRefAliasOpsPass> { void runOnOperation() override; }; diff --git a/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp b/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp index 345c71de7af7..53accebe9ca9 100644 --- a/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp @@ -11,7 +11,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Dialect/Func/IR/FuncOps.h" @@ -20,6 +19,13 @@ #include "llvm/ADT/SmallSet.h" #include "llvm/Support/Debug.h" +namespace mlir { +namespace memref { +#define GEN_PASS_DEF_NORMALIZEMEMREFS +#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc" +} // namespace memref +} // namespace mlir + #define DEBUG_TYPE "normalize-memrefs" using namespace mlir; @@ -32,7 +38,8 @@ namespace { /// such functions as normalizable. Also, if a normalizable function is known /// to call a non-normalizable function, we treat that function as /// non-normalizable as well. We assume external functions to be normalizable. -struct NormalizeMemRefs : public NormalizeMemRefsBase<NormalizeMemRefs> { +struct NormalizeMemRefs + : public memref::impl::NormalizeMemRefsBase<NormalizeMemRefs> { void runOnOperation() override; void normalizeFuncOpMemRefs(func::FuncOp funcOp, ModuleOp moduleOp); bool areMemRefsNormalizable(func::FuncOp funcOp); diff --git a/mlir/lib/Dialect/MemRef/Transforms/PassDetail.h b/mlir/lib/Dialect/MemRef/Transforms/PassDetail.h deleted file mode 100644 index ba7a87502dcb..000000000000 --- a/mlir/lib/Dialect/MemRef/Transforms/PassDetail.h +++ /dev/null @@ -1,49 +0,0 @@ -//===- PassDetail.h - MemRef 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 DIALECT_MEMREF_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_MEMREF_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/IR/BuiltinOps.h" -#include "mlir/IR/DialectRegistry.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -class AffineDialect; - -// Forward declaration from Dialect.h -template <typename ConcreteDialect> -void registerDialect(DialectRegistry ®istry); - -namespace arith { -class ArithmeticDialect; -} // namespace arith - -namespace func { -class FuncDialect; -} // namespace func - -namespace memref { -class MemRefDialect; -} // namespace memref - -namespace tensor { -class TensorDialect; -} // namespace tensor - -namespace vector { -class VectorDialect; -} // namespace vector - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_MEMREF_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/MemRef/Transforms/ResolveShapedTypeResultDims.cpp b/mlir/lib/Dialect/MemRef/Transforms/ResolveShapedTypeResultDims.cpp index d90ec22fc0b5..3f3687a458a2 100644 --- a/mlir/lib/Dialect/MemRef/Transforms/ResolveShapedTypeResultDims.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/ResolveShapedTypeResultDims.cpp @@ -11,15 +11,23 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/MemRef/Transforms/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" -#include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Interfaces/InferTypeOpInterface.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +namespace memref { +#define GEN_PASS_DEF_RESOLVERANKEDSHAPETYPERESULTDIMS +#define GEN_PASS_DEF_RESOLVESHAPEDTYPERESULTDIMS +#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc" +} // namespace memref +} // namespace mlir + using namespace mlir; namespace { @@ -108,13 +116,14 @@ struct DimOfReifyRankedShapedTypeOpInterface : public OpRewritePattern<OpTy> { namespace { struct ResolveRankedShapeTypeResultDimsPass final - : public ResolveRankedShapeTypeResultDimsBase< + : public memref::impl::ResolveRankedShapeTypeResultDimsBase< ResolveRankedShapeTypeResultDimsPass> { void runOnOperation() override; }; struct ResolveShapedTypeResultDimsPass final - : public ResolveShapedTypeResultDimsBase<ResolveShapedTypeResultDimsPass> { + : public memref::impl::ResolveShapedTypeResultDimsBase< + ResolveShapedTypeResultDimsPass> { void runOnOperation() override; }; diff --git a/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp b/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp index d24001c4d28a..4d38319f48f9 100644 --- a/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp +++ b/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp @@ -11,7 +11,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" diff --git a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp index da8292b33a80..968afcb4bef8 100644 --- a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp +++ b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp @@ -9,12 +9,13 @@ // This file implements transforms to optimize accesses to shared memory. // //===----------------------------------------------------------------------===// -#include "PassDetail.h" + +#include "mlir/Dialect/NVGPU/Passes.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h" -#include "mlir/Dialect/NVGPU/Passes.h" #include "mlir/Dialect/NVGPU/Transforms/Transforms.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "mlir/Interfaces/SideEffectInterfaces.h" @@ -22,6 +23,13 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/Support/MathExtras.h" +namespace mlir { +namespace nvgpu { +#define GEN_PASS_DEF_OPTIMIZESHAREDMEMORY +#include "mlir/Dialect/NVGPU/Passes.h.inc" +} // namespace nvgpu +} // namespace mlir + using namespace mlir; using namespace mlir::nvgpu; @@ -242,7 +250,7 @@ mlir::nvgpu::optimizeSharedMemoryReadsAndWrites(Operation *parentOp, namespace { class OptimizeSharedMemoryPass - : public OptimizeSharedMemoryBase<OptimizeSharedMemoryPass> { + : public nvgpu::impl::OptimizeSharedMemoryBase<OptimizeSharedMemoryPass> { public: OptimizeSharedMemoryPass() = default; diff --git a/mlir/lib/Dialect/NVGPU/Transforms/PassDetail.h b/mlir/lib/Dialect/NVGPU/Transforms/PassDetail.h deleted file mode 100644 index 884a103543d7..000000000000 --- a/mlir/lib/Dialect/NVGPU/Transforms/PassDetail.h +++ /dev/null @@ -1,33 +0,0 @@ -//===- PassDetail.h - NVGPU 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 DIALECT_NVGPU_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_NVGPU_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/IR/BuiltinOps.h" -#include "mlir/IR/Dialect.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { -namespace arith { -class ArithmeticDialect; -} // namespace arith - -namespace memref { -class MemRefDialect; -} // namespace memref - -namespace vector { -class VectorDialect; -} // namespace vector - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/NVGPU/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_NVGPU_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp b/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp index 961e11a62052..e1ee6a2e0a50 100644 --- a/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp @@ -6,20 +6,25 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Passes.h" + #include "mlir/Dialect/Bufferization/IR/Bufferization.h" +#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_SCFBUFFERIZE +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::scf; namespace { -struct SCFBufferizePass : public SCFBufferizeBase<SCFBufferizePass> { +struct SCFBufferizePass : public impl::SCFBufferizeBase<SCFBufferizePass> { void runOnOperation() override { auto func = getOperation(); auto *context = &getContext(); diff --git a/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp b/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp index 14eb075d8c89..e7a600343e6a 100644 --- a/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp @@ -10,14 +10,19 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Passes.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_SCFFORTOWHILELOOP +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + using namespace llvm; using namespace mlir; using scf::ForOp; @@ -98,7 +103,7 @@ struct ForLoopLoweringPattern : public OpRewritePattern<ForOp> { } }; -struct ForToWhileLoop : public SCFForToWhileLoopBase<ForToWhileLoop> { +struct ForToWhileLoop : public impl::SCFForToWhileLoopBase<ForToWhileLoop> { void runOnOperation() override { auto *parentOp = getOperation(); MLIRContext *ctx = parentOp->getContext(); diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp index 0797dc20fc0f..c153eb6a47d0 100644 --- a/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp @@ -11,11 +11,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/Dialect/SCF/Utils/AffineCanonicalizationUtils.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" @@ -23,6 +23,11 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/TypeSwitch.h" +namespace mlir { +#define GEN_PASS_DEF_SCFFORLOOPCANONICALIZATION +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::scf; @@ -193,7 +198,7 @@ struct AffineOpSCFCanonicalizationPattern : public OpRewritePattern<OpTy> { }; struct SCFForLoopCanonicalization - : public SCFForLoopCanonicalizationBase<SCFForLoopCanonicalization> { + : public impl::SCFForLoopCanonicalizationBase<SCFForLoopCanonicalization> { void runOnOperation() override { auto *parentOp = getOperation(); MLIRContext *ctx = parentOp->getContext(); diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp index c1d04afa920d..7a162dce8193 100644 --- a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp @@ -10,7 +10,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/SCF/Transforms/Patterns.h" @@ -19,6 +18,7 @@ #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Support/MathExtras.h" +#include "llvm/ADT/MapVector.h" using namespace mlir; using namespace mlir::scf; diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp index c5aabe9bb298..24e4c4fd0a86 100644 --- a/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp @@ -10,20 +10,25 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Passes.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/Dialect/SCF/Utils/Utils.h" #include "mlir/IR/BlockAndValueMapping.h" +namespace mlir { +#define GEN_PASS_DEF_SCFFORLOOPRANGEFOLDING +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::scf; namespace { struct ForLoopRangeFolding - : public SCFForLoopRangeFoldingBase<ForLoopRangeFolding> { + : public impl::SCFForLoopRangeFoldingBase<ForLoopRangeFolding> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopSpecialization.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopSpecialization.cpp index 3778bb9a04ed..7e9fd4f525cb 100644 --- a/mlir/lib/Dialect/SCF/Transforms/LoopSpecialization.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/LoopSpecialization.cpp @@ -11,12 +11,12 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Passes.h" + #include "mlir/Dialect/Affine/Analysis/AffineStructures.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/Dialect/SCF/Utils/AffineCanonicalizationUtils.h" #include "mlir/Dialect/Utils/StaticValueUtils.h" @@ -26,6 +26,13 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/DenseMap.h" +namespace mlir { +#define GEN_PASS_DEF_SCFFORLOOPPEELING +#define GEN_PASS_DEF_SCFFORLOOPSPECIALIZATION +#define GEN_PASS_DEF_SCFPARALLELLOOPSPECIALIZATION +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using scf::ForOp; using scf::ParallelOp; @@ -235,7 +242,8 @@ struct ForLoopPeelingPattern : public OpRewritePattern<ForOp> { namespace { struct ParallelLoopSpecialization - : public SCFParallelLoopSpecializationBase<ParallelLoopSpecialization> { + : public impl::SCFParallelLoopSpecializationBase< + ParallelLoopSpecialization> { void runOnOperation() override { getOperation()->walk( [](ParallelOp op) { specializeParallelLoopForUnrolling(op); }); @@ -243,13 +251,13 @@ struct ParallelLoopSpecialization }; struct ForLoopSpecialization - : public SCFForLoopSpecializationBase<ForLoopSpecialization> { + : public impl::SCFForLoopSpecializationBase<ForLoopSpecialization> { void runOnOperation() override { getOperation()->walk([](ForOp op) { specializeForLoopForUnrolling(op); }); } }; -struct ForLoopPeeling : public SCFForLoopPeelingBase<ForLoopPeeling> { +struct ForLoopPeeling : public impl::SCFForLoopPeelingBase<ForLoopPeeling> { void runOnOperation() override { auto *parentOp = getOperation(); MLIRContext *ctx = parentOp->getContext(); diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp index 3403cd55ae4f..34c5e777f948 100644 --- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp @@ -6,21 +6,26 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/SCF/Transforms/Passes.h" + +#include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/SCF/Utils/Utils.h" #include "mlir/Transforms/RegionUtils.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_SCFPARALLELLOOPCOLLAPSING +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "parallel-loop-collapsing" using namespace mlir; namespace { struct ParallelLoopCollapsing - : public SCFParallelLoopCollapsingBase<ParallelLoopCollapsing> { + : public impl::SCFParallelLoopCollapsingBase<ParallelLoopCollapsing> { void runOnOperation() override { Operation *module = getOperation(); diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp index cab8b765661c..e2530572d638 100644 --- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp @@ -10,15 +10,20 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Passes.h" + #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/Builders.h" #include "mlir/IR/OpDefinition.h" +namespace mlir { +#define GEN_PASS_DEF_SCFPARALLELLOOPFUSION +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::scf; @@ -162,7 +167,7 @@ void mlir::scf::naivelyFuseParallelOps(Region ®ion) { namespace { struct ParallelLoopFusion - : public SCFParallelLoopFusionBase<ParallelLoopFusion> { + : public impl::SCFParallelLoopFusionBase<ParallelLoopFusion> { void runOnOperation() override { getOperation()->walk([&](Operation *child) { for (Region ®ion : child->getRegions()) diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp index 942430588b77..afcb2c24a62e 100644 --- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp @@ -10,14 +10,19 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/Dialect/SCF/Utils/Utils.h" +namespace mlir { +#define GEN_PASS_DEF_SCFPARALLELLOOPTILING +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::scf; @@ -181,7 +186,7 @@ mlir::scf::tileParallelLoop(ParallelOp op, ArrayRef<int64_t> tileSizes, namespace { struct ParallelLoopTiling - : public SCFParallelLoopTilingBase<ParallelLoopTiling> { + : public impl::SCFParallelLoopTilingBase<ParallelLoopTiling> { ParallelLoopTiling() = default; explicit ParallelLoopTiling(ArrayRef<int64_t> tileSizes, bool noMinMaxBounds = false) { diff --git a/mlir/lib/Dialect/SCF/Transforms/PassDetail.h b/mlir/lib/Dialect/SCF/Transforms/PassDetail.h deleted file mode 100644 index 98109f7ff4cb..000000000000 --- a/mlir/lib/Dialect/SCF/Transforms/PassDetail.h +++ /dev/null @@ -1,43 +0,0 @@ -//===- PassDetail.h - Loop 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 DIALECT_SCF_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_SCF_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { -// Forward declaration from Dialect.h -template <typename ConcreteDialect> -void registerDialect(DialectRegistry ®istry); - -class AffineDialect; - -namespace arith { -class ArithmeticDialect; -} // namespace arith - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace memref { -class MemRefDialect; -} // namespace memref - -namespace tensor { -class TensorDialect; -} // namespace tensor - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_SCF_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp b/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp index e7408d9a8596..fb15fcb3579e 100644 --- a/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp @@ -6,7 +6,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/SCF/Transforms/Passes.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" diff --git a/mlir/lib/Dialect/SPIRV/Transforms/CanonicalizeGLPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/CanonicalizeGLPass.cpp index 7277e8491fb5..94039c6a8448 100644 --- a/mlir/lib/Dialect/SPIRV/Transforms/CanonicalizeGLPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/CanonicalizeGLPass.cpp @@ -6,18 +6,25 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SPIRV/Transforms/Passes.h" + #include "mlir/Dialect/SPIRV/IR/SPIRVGLCanonicalization.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" -#include "mlir/Dialect/SPIRV/Transforms/Passes.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +namespace spirv { +#define GEN_PASS_DEF_SPIRVCANONICALIZEGL +#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" +} // namespace spirv +} // namespace mlir + using namespace mlir; namespace { class CanonicalizeGLPass final - : public SPIRVCanonicalizeGLBase<CanonicalizeGLPass> { + : public spirv::impl::SPIRVCanonicalizeGLBase<CanonicalizeGLPass> { public: void runOnOperation() override { RewritePatternSet patterns(&getContext()); diff --git a/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp index dd958e522a26..2e8f70c7e434 100644 --- a/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp @@ -13,16 +13,23 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SPIRV/Transforms/Passes.h" + #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" -#include "mlir/Dialect/SPIRV/Transforms/Passes.h" #include "mlir/Dialect/SPIRV/Utils/LayoutUtils.h" #include "mlir/Transforms/DialectConversion.h" using namespace mlir; +namespace mlir { +namespace spirv { +#define GEN_PASS_DEF_SPIRVCOMPOSITETYPELAYOUT +#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" +} // namespace spirv +} // namespace mlir + namespace { class SPIRVGlobalVariableOpLayoutInfoDecoration : public OpRewritePattern<spirv::GlobalVariableOp> { @@ -99,7 +106,7 @@ static void populateSPIRVLayoutInfoPatterns(RewritePatternSet &patterns) { namespace { class DecorateSPIRVCompositeTypeLayoutPass - : public SPIRVCompositeTypeLayoutBase< + : public spirv::impl::SPIRVCompositeTypeLayoutBase< DecorateSPIRVCompositeTypeLayoutPass> { void runOnOperation() override; }; diff --git a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp index 263f00ef1d86..ad296e32fd29 100644 --- a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp @@ -11,15 +11,22 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SPIRV/Transforms/Passes.h" + #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" -#include "mlir/Dialect/SPIRV/Transforms/Passes.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" #include "mlir/Dialect/SPIRV/Utils/LayoutUtils.h" #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/SetVector.h" +namespace mlir { +namespace spirv { +#define GEN_PASS_DEF_SPIRVLOWERABIATTRIBUTES +#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" +} // namespace spirv +} // namespace mlir + using namespace mlir; /// Creates a global variable for an argument based on the ABI info. @@ -165,7 +172,7 @@ public: /// Pass to implement the ABI information specified as attributes. class LowerABIAttributesPass final - : public SPIRVLowerABIAttributesBase<LowerABIAttributesPass> { + : public spirv::impl::SPIRVLowerABIAttributesBase<LowerABIAttributesPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Dialect/SPIRV/Transforms/PassDetail.h b/mlir/lib/Dialect/SPIRV/Transforms/PassDetail.h deleted file mode 100644 index 15f5c54d38f1..000000000000 --- a/mlir/lib/Dialect/SPIRV/Transforms/PassDetail.h +++ /dev/null @@ -1,26 +0,0 @@ -//===- PassDetail.h - SPIRV 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 DIALECT_SPIRV_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_SPIRV_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/IR/BuiltinOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -namespace spirv { -class ModuleOp; -} // namespace spirv - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_SPIRV_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/SPIRV/Transforms/RewriteInsertsPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/RewriteInsertsPass.cpp index c381ee545121..95d21f35e0d0 100644 --- a/mlir/lib/Dialect/SPIRV/Transforms/RewriteInsertsPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/RewriteInsertsPass.cpp @@ -12,12 +12,19 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/Dialect/SPIRV/Transforms/Passes.h" + +#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinOps.h" +namespace mlir { +namespace spirv { +#define GEN_PASS_DEF_SPIRVREWRITEINSERTSPASS +#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" +} // namespace spirv +} // namespace mlir + using namespace mlir; namespace { @@ -25,7 +32,7 @@ namespace { /// Replaces sequential chains of `spirv::CompositeInsertOp` operation into /// `spirv::CompositeConstructOp` operation if possible. class RewriteInsertsPass - : public SPIRVRewriteInsertsPassBase<RewriteInsertsPass> { + : public spirv::impl::SPIRVRewriteInsertsPassBase<RewriteInsertsPass> { public: void runOnOperation() override; diff --git a/mlir/lib/Dialect/SPIRV/Transforms/UnifyAliasedResourcePass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/UnifyAliasedResourcePass.cpp index 8299fd05cda9..140e31619649 100644 --- a/mlir/lib/Dialect/SPIRV/Transforms/UnifyAliasedResourcePass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/UnifyAliasedResourcePass.cpp @@ -11,11 +11,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SPIRV/Transforms/Passes.h" + #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h" -#include "mlir/Dialect/SPIRV/Transforms/Passes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/BuiltinTypes.h" @@ -27,6 +27,13 @@ #include <algorithm> #include <iterator> +namespace mlir { +namespace spirv { +#define GEN_PASS_DEF_SPIRVUNIFYALIASEDRESOURCEPASS +#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" +} // namespace spirv +} // namespace mlir + #define DEBUG_TYPE "spirv-unify-aliased-resource" using namespace mlir; @@ -521,7 +528,8 @@ struct ConvertStore : public ConvertAliasResource<spirv::StoreOp> { namespace { class UnifyAliasedResourcePass final - : public SPIRVUnifyAliasedResourcePassBase<UnifyAliasedResourcePass> { + : public spirv::impl::SPIRVUnifyAliasedResourcePassBase< + UnifyAliasedResourcePass> { public: void runOnOperation() override; }; diff --git a/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp index 3232c984b290..9cb79f80bab0 100644 --- a/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp @@ -11,24 +11,32 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SPIRV/Transforms/Passes.h" + #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h" #include "mlir/Dialect/SPIRV/IR/TargetAndABI.h" -#include "mlir/Dialect/SPIRV/Transforms/Passes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/Visitors.h" #include "llvm/ADT/SetVector.h" #include "llvm/ADT/SmallSet.h" #include "llvm/ADT/StringExtras.h" +namespace mlir { +namespace spirv { +#define GEN_PASS_DEF_SPIRVUPDATEVCE +#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" +} // namespace spirv +} // namespace mlir + using namespace mlir; namespace { /// Pass to deduce minimal version/extension/capability requirements for a /// spirv::ModuleOp. -class UpdateVCEPass final : public SPIRVUpdateVCEBase<UpdateVCEPass> { +class UpdateVCEPass final + : public spirv::impl::SPIRVUpdateVCEBase<UpdateVCEPass> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp index 373552801673..9dadbdbc91ec 100644 --- a/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp @@ -6,21 +6,28 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" -#include "PassDetail.h" +#include "mlir/Dialect/Shape/Transforms/Passes.h" + #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" +#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/Shape/IR/Shape.h" #include "mlir/Dialect/Shape/Transforms/BufferizableOpInterfaceImpl.h" -#include "mlir/Dialect/Shape/Transforms/Passes.h" #include "mlir/Pass/Pass.h" +namespace mlir { +#define GEN_PASS_DEF_SHAPEBUFFERIZE +#include "mlir/Dialect/Shape/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace bufferization; namespace { -struct ShapeBufferizePass : public ShapeBufferizeBase<ShapeBufferizePass> { +struct ShapeBufferizePass + : public impl::ShapeBufferizeBase<ShapeBufferizePass> { void runOnOperation() override { BufferizationOptions options = getPartialBufferizationOptions(); options.opFilter.allowDialect<shape::ShapeDialect>(); diff --git a/mlir/lib/Dialect/Shape/Transforms/PassDetail.h b/mlir/lib/Dialect/Shape/Transforms/PassDetail.h deleted file mode 100644 index 2856871f8b5e..000000000000 --- a/mlir/lib/Dialect/Shape/Transforms/PassDetail.h +++ /dev/null @@ -1,30 +0,0 @@ -//===- PassDetail.h - Shape 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 DIALECT_SHAPE_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_SHAPE_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace memref { -class MemRefDialect; -} // namespace memref - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Shape/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_SHAPE_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/Shape/Transforms/RemoveShapeConstraints.cpp b/mlir/lib/Dialect/Shape/Transforms/RemoveShapeConstraints.cpp index 9375be93f6d8..e1cccd8fd5d6 100644 --- a/mlir/lib/Dialect/Shape/Transforms/RemoveShapeConstraints.cpp +++ b/mlir/lib/Dialect/Shape/Transforms/RemoveShapeConstraints.cpp @@ -6,12 +6,18 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/Shape/IR/Shape.h" #include "mlir/Dialect/Shape/Transforms/Passes.h" + +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/Shape/IR/Shape.h" #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_REMOVESHAPECONSTRAINTS +#include "mlir/Dialect/Shape/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -41,7 +47,7 @@ public: /// Removal pass. class RemoveShapeConstraintsPass - : public RemoveShapeConstraintsBase<RemoveShapeConstraintsPass> { + : public impl::RemoveShapeConstraintsBase<RemoveShapeConstraintsPass> { void runOnOperation() override { MLIRContext &ctx = getContext(); diff --git a/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp b/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp index 4179e604b220..c73512cd1579 100644 --- a/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp +++ b/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp @@ -6,15 +6,21 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Shape/Transforms/Passes.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Shape/IR/Shape.h" -#include "mlir/Dialect/Shape/Transforms/Passes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +#define GEN_PASS_DEF_SHAPETOSHAPELOWERING +#include "mlir/Dialect/Shape/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::shape; @@ -53,7 +59,7 @@ NumElementsOpConverter::matchAndRewrite(NumElementsOp op, namespace { struct ShapeToShapeLowering - : public ShapeToShapeLoweringBase<ShapeToShapeLowering> { + : public impl::ShapeToShapeLoweringBase<ShapeToShapeLowering> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp index 643cff984449..b380c50f98bf 100644 --- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp @@ -20,23 +20,24 @@ #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +#define GEN_PASS_DEF_SPARSIFICATIONPASS +#define GEN_PASS_DEF_SPARSETENSORCONVERSIONPASS +#define GEN_PASS_DEF_SPARSETENSORCODEGEN +#include "mlir/Dialect/SparseTensor/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; using namespace mlir::sparse_tensor; namespace { //===----------------------------------------------------------------------===// -// Passes declaration. -//===----------------------------------------------------------------------===// - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/SparseTensor/Transforms/Passes.h.inc" - -//===----------------------------------------------------------------------===// // Passes implementation. //===----------------------------------------------------------------------===// -struct SparsificationPass : public SparsificationBase<SparsificationPass> { +struct SparsificationPass + : public impl::SparsificationPassBase<SparsificationPass> { SparsificationPass() = default; SparsificationPass(const SparsificationPass &pass) = default; @@ -68,7 +69,7 @@ struct SparsificationPass : public SparsificationBase<SparsificationPass> { }; struct SparseTensorConversionPass - : public SparseTensorConversionBase<SparseTensorConversionPass> { + : public impl::SparseTensorConversionPassBase<SparseTensorConversionPass> { SparseTensorConversionPass() = default; SparseTensorConversionPass(const SparseTensorConversionPass &pass) = default; @@ -145,7 +146,7 @@ struct SparseTensorConversionPass }; struct SparseTensorCodegenPass - : public SparseTensorCodegenBase<SparseTensorCodegenPass> { + : public impl::SparseTensorCodegenBase<SparseTensorCodegenPass> { SparseTensorCodegenPass() = default; SparseTensorCodegenPass(const SparseTensorCodegenPass &pass) = default; diff --git a/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp index caf3f0f42e4a..11dad0f07754 100644 --- a/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp @@ -11,10 +11,10 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" -#include "PassDetail.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" @@ -23,11 +23,19 @@ #include "mlir/IR/ImplicitLocOpBuilder.h" #include "mlir/Transforms/DialectConversion.h" +namespace mlir { +namespace tensor { +#define GEN_PASS_DEF_TENSORBUFFERIZE +#include "mlir/Dialect/Tensor/Transforms/Passes.h.inc" +} // namespace tensor +} // namespace mlir + using namespace mlir; using namespace bufferization; namespace { -struct TensorBufferizePass : public TensorBufferizeBase<TensorBufferizePass> { +struct TensorBufferizePass + : public tensor::impl::TensorBufferizeBase<TensorBufferizePass> { void runOnOperation() override { BufferizationOptions options = getPartialBufferizationOptions(); options.opFilter.allowDialect<tensor::TensorDialect>(); diff --git a/mlir/lib/Dialect/Tensor/Transforms/PassDetail.h b/mlir/lib/Dialect/Tensor/Transforms/PassDetail.h deleted file mode 100644 index 033372b0533b..000000000000 --- a/mlir/lib/Dialect/Tensor/Transforms/PassDetail.h +++ /dev/null @@ -1,34 +0,0 @@ -//===- PassDetail.h - GPU 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 DIALECT_TENSOR_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_TENSOR_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace memref { -class MemRefDialect; -} // namespace memref - -namespace scf { -class SCFDialect; -} // namespace scf - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Tensor/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_TENSOR_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp index b3238904e793..d588836ff108 100644 --- a/mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp @@ -11,11 +11,11 @@ // //===----------------------------------------------------------------------===// +#include "mlir/Dialect/Tosa/Transforms/Passes.h" + #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" -#include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/Dialect/Tosa/Utils/ShapeUtils.h" #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/Builders.h" @@ -26,6 +26,13 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/Support/FormatVariadic.h" +namespace mlir { +namespace tosa { +#define GEN_PASS_DEF_TOSAINFERSHAPES +#include "mlir/Dialect/Tosa/Transforms/Passes.h.inc" +} // namespace tosa +} // namespace mlir + using namespace mlir; using namespace mlir::tosa; @@ -275,7 +282,8 @@ void propagateShapesInRegion(Region ®ion) { /// Pass that performs shape propagation across TOSA operations. This includes /// migrating to within the regions of if/while operations. -struct TosaInferShapes : public TosaInferShapesBase<TosaInferShapes> { +struct TosaInferShapes + : public tosa::impl::TosaInferShapesBase<TosaInferShapes> { public: void runOnOperation() override { func::FuncOp func = getOperation(); diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp index aed8ff5302e5..a217f66cd84c 100644 --- a/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp @@ -10,12 +10,20 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" + +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/Tosa/IR/TosaOps.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +namespace tosa { +#define GEN_PASS_DEF_TOSALAYERWISECONSTANTFOLDPASS +#include "mlir/Dialect/Tosa/Transforms/Passes.h.inc" +} // namespace tosa +} // namespace mlir + using namespace mlir; using namespace mlir::tosa; @@ -35,7 +43,8 @@ void populateTosaOpsCanonicalizationPatterns(MLIRContext *ctx, } struct TosaLayerwiseConstantFoldPass - : public TosaLayerwiseConstantFoldPassBase<TosaLayerwiseConstantFoldPass> { + : public tosa::impl::TosaLayerwiseConstantFoldPassBase< + TosaLayerwiseConstantFoldPass> { void runOnOperation() override { auto *ctx = &getContext(); RewritePatternSet patterns(ctx); diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp index 61f601456f0f..120aba77aa55 100644 --- a/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp @@ -10,14 +10,21 @@ // //===----------------------------------------------------------------------===// +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" -#include "mlir/Dialect/Tosa/IR//TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" +#include "mlir/Dialect/Tosa/IR/TosaOps.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/Dialect/Tosa/Utils/QuantUtils.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +namespace tosa { +#define GEN_PASS_DEF_TOSAMAKEBROADCASTABLE +#include "mlir/Dialect/Tosa/Transforms/Passes.h.inc" +} // namespace tosa +} // namespace mlir + using namespace mlir; using namespace mlir::tosa; @@ -232,7 +239,7 @@ namespace { /// Pass that enables broadcast by making all input arrays have the same /// number of dimensions. Insert RESHAPE operations to lower rank operand struct TosaMakeBroadcastable - : public TosaMakeBroadcastableBase<TosaMakeBroadcastable> { + : public tosa::impl::TosaMakeBroadcastableBase<TosaMakeBroadcastable> { public: void runOnOperation() override { auto func = getOperation(); diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp index 78b8cb3084af..cef903a39e45 100644 --- a/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp @@ -12,18 +12,27 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Dialect/Tosa/IR/TosaOps.h" -#include "mlir/Dialect/Tosa/Transforms/PassDetail.h" #include "mlir/Dialect/Tosa/Transforms/Passes.h" + +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/Tosa/IR/TosaOps.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +namespace mlir { +namespace tosa { +#define GEN_PASS_DEF_TOSAOPTIONALDECOMPOSITIONS +#include "mlir/Dialect/Tosa/Transforms/Passes.h.inc" +} // namespace tosa +} // namespace mlir + using namespace mlir; namespace { struct TosaOptionalDecompositions - : public TosaOptionalDecompositionsBase<TosaOptionalDecompositions> { + : public tosa::impl::TosaOptionalDecompositionsBase< + TosaOptionalDecompositions> { void runOnOperation() override { auto *ctx = &getContext(); RewritePatternSet patterns(ctx); diff --git a/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp b/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp index f70e2b6475aa..dbb03e0b3609 100644 --- a/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp +++ b/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp @@ -11,12 +11,20 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Dialect/Transform/IR/TransformInterfaces.h" #include "mlir/Dialect/Transform/Transforms/Passes.h" + +#include "mlir/Dialect/Transform/IR/TransformInterfaces.h" #include "mlir/Interfaces/SideEffectInterfaces.h" #include "mlir/Pass/Pass.h" #include "llvm/ADT/SetOperations.h" +namespace mlir { +namespace transform { +#define GEN_PASS_DEF_CHECKUSES +#include "mlir/Dialect/Transform/Transforms/Passes.h.inc" +} // namespace transform +} // namespace mlir + using namespace mlir; namespace { @@ -361,12 +369,9 @@ private: DenseMap<Block *, llvm::SmallPtrSet<Block *, 4>> reachableFromCache; }; -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Transform/Transforms/Passes.h.inc" - //// A simple pass that warns about any use of a value by a transform operation // that may be using the value after it has been freed. -class CheckUsesPass : public CheckUsesBase<CheckUsesPass> { +class CheckUsesPass : public transform::impl::CheckUsesBase<CheckUsesPass> { public: void runOnOperation() override { auto &analysis = getAnalysis<TransformOpMemFreeAnalysis>(); diff --git a/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp index f98eeda1bde5..ee99a99b5610 100644 --- a/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp @@ -11,20 +11,29 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" -#include "PassDetail.h" + #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "mlir/Dialect/Vector/Transforms/BufferizableOpInterfaceImpl.h" #include "mlir/Dialect/Vector/Transforms/Passes.h" +namespace mlir { +namespace vector { +#define GEN_PASS_DEF_VECTORBUFFERIZE +#include "mlir/Dialect/Vector/Transforms/Passes.h.inc" +} // namespace vector +} // namespace mlir + using namespace mlir; using namespace bufferization; namespace { -struct VectorBufferizePass : public VectorBufferizeBase<VectorBufferizePass> { +struct VectorBufferizePass + : public vector::impl::VectorBufferizeBase<VectorBufferizePass> { void runOnOperation() override { BufferizationOptions options = getPartialBufferizationOptions(); options.opFilter.allowDialect<vector::VectorDialect>(); diff --git a/mlir/lib/Dialect/Vector/Transforms/PassDetail.h b/mlir/lib/Dialect/Vector/Transforms/PassDetail.h deleted file mode 100644 index 305aae47e5bb..000000000000 --- a/mlir/lib/Dialect/Vector/Transforms/PassDetail.h +++ /dev/null @@ -1,30 +0,0 @@ -//===- PassDetail.h - Vector 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 DIALECT_VECTOR_TRANSFORMS_PASSDETAIL_H_ -#define DIALECT_VECTOR_TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/Pass/Pass.h" - -namespace mlir { - -namespace bufferization { -class BufferizationDialect; -} // namespace bufferization - -namespace memref { -class MemRefDialect; -} // namespace memref - -#define GEN_PASS_CLASSES -#include "mlir/Dialect/Vector/Transforms/Passes.h.inc" - -} // namespace mlir - -#endif // DIALECT_VECTOR_TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Reducer/OptReductionPass.cpp b/mlir/lib/Reducer/OptReductionPass.cpp index a7f09b4c923f..8618de5eeee7 100644 --- a/mlir/lib/Reducer/OptReductionPass.cpp +++ b/mlir/lib/Reducer/OptReductionPass.cpp @@ -14,18 +14,22 @@ #include "mlir/Pass/PassManager.h" #include "mlir/Pass/PassRegistry.h" -#include "mlir/Reducer/PassDetail.h" #include "mlir/Reducer/Passes.h" #include "mlir/Reducer/Tester.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_OPTREDUCTION +#include "mlir/Reducer/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "mlir-reduce" using namespace mlir; namespace { -class OptReductionPass : public OptReductionBase<OptReductionPass> { +class OptReductionPass : public impl::OptReductionBase<OptReductionPass> { public: /// Runs the pass instance in the pass pipeline. void runOnOperation() override; diff --git a/mlir/lib/Reducer/ReductionTreePass.cpp b/mlir/lib/Reducer/ReductionTreePass.cpp index 05f0f749166e..152c8fb2d57f 100644 --- a/mlir/lib/Reducer/ReductionTreePass.cpp +++ b/mlir/lib/Reducer/ReductionTreePass.cpp @@ -16,7 +16,6 @@ #include "mlir/IR/DialectInterface.h" #include "mlir/IR/OpDefinition.h" -#include "mlir/Reducer/PassDetail.h" #include "mlir/Reducer/Passes.h" #include "mlir/Reducer/ReductionNode.h" #include "mlir/Reducer/ReductionPatternInterface.h" @@ -29,6 +28,11 @@ #include "llvm/Support/Allocator.h" #include "llvm/Support/ManagedStatic.h" +namespace mlir { +#define GEN_PASS_DEF_REDUCTIONTREE +#include "mlir/Reducer/Passes.h.inc" +} // namespace mlir + using namespace mlir; /// We implicitly number each operation in the region and if an operation's @@ -183,7 +187,7 @@ public: /// This class defines the Reduction Tree Pass. It provides a framework to /// to implement a reduction pass using a tree structure to keep track of the /// generated reduced variants. -class ReductionTreePass : public ReductionTreeBase<ReductionTreePass> { +class ReductionTreePass : public impl::ReductionTreeBase<ReductionTreePass> { public: ReductionTreePass() = default; ReductionTreePass(const ReductionTreePass &pass) = default; diff --git a/mlir/lib/Transforms/CSE.cpp b/mlir/lib/Transforms/CSE.cpp index 783a2260275d..959cec3c0b33 100644 --- a/mlir/lib/Transforms/CSE.cpp +++ b/mlir/lib/Transforms/CSE.cpp @@ -11,11 +11,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/IR/Dominance.h" #include "mlir/Interfaces/SideEffectInterfaces.h" #include "mlir/Pass/Pass.h" -#include "mlir/Transforms/Passes.h" #include "llvm/ADT/DenseMapInfo.h" #include "llvm/ADT/Hashing.h" #include "llvm/ADT/ScopedHashTable.h" @@ -23,6 +23,11 @@ #include "llvm/Support/RecyclingAllocator.h" #include <deque> +namespace mlir { +#define GEN_PASS_DEF_CSE +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { @@ -53,7 +58,7 @@ struct SimpleOperationInfo : public llvm::DenseMapInfo<Operation *> { namespace { /// Simple common sub-expression elimination. -struct CSE : public CSEBase<CSE> { +struct CSE : public impl::CSEBase<CSE> { /// Shared implementation of operation elimination and scoped map definitions. using AllocatorTy = llvm::RecyclingAllocator< llvm::BumpPtrAllocator, diff --git a/mlir/lib/Transforms/Canonicalizer.cpp b/mlir/lib/Transforms/Canonicalizer.cpp index 3f6dbe933b2c..a4215629a964 100644 --- a/mlir/lib/Transforms/Canonicalizer.cpp +++ b/mlir/lib/Transforms/Canonicalizer.cpp @@ -11,16 +11,21 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" -#include "mlir/Transforms/Passes.h" + +namespace mlir { +#define GEN_PASS_DEF_CANONICALIZER +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir using namespace mlir; namespace { /// Canonicalize operations in nested regions. -struct Canonicalizer : public CanonicalizerBase<Canonicalizer> { +struct Canonicalizer : public impl::CanonicalizerBase<Canonicalizer> { Canonicalizer() = default; Canonicalizer(const GreedyRewriteConfig &config, ArrayRef<std::string> disabledPatterns, diff --git a/mlir/lib/Transforms/ControlFlowSink.cpp b/mlir/lib/Transforms/ControlFlowSink.cpp index 5b0cfc269511..0fcdb9c51fa6 100644 --- a/mlir/lib/Transforms/ControlFlowSink.cpp +++ b/mlir/lib/Transforms/ControlFlowSink.cpp @@ -13,19 +13,24 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/IR/Dominance.h" #include "mlir/Interfaces/ControlFlowInterfaces.h" #include "mlir/Interfaces/SideEffectInterfaces.h" #include "mlir/Transforms/ControlFlowSinkUtils.h" -#include "mlir/Transforms/Passes.h" #include "mlir/Transforms/SideEffectUtils.h" +namespace mlir { +#define GEN_PASS_DEF_CONTROLFLOWSINK +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { /// A control-flow sink pass. -struct ControlFlowSink : public ControlFlowSinkBase<ControlFlowSink> { +struct ControlFlowSink : public impl::ControlFlowSinkBase<ControlFlowSink> { void runOnOperation() override; }; } // end anonymous namespace diff --git a/mlir/lib/Transforms/Inliner.cpp b/mlir/lib/Transforms/Inliner.cpp index 5ce32b14185e..1fdeeafc8ebd 100644 --- a/mlir/lib/Transforms/Inliner.cpp +++ b/mlir/lib/Transforms/Inliner.cpp @@ -13,7 +13,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/Analysis/CallGraph.h" #include "mlir/IR/Threading.h" #include "mlir/Interfaces/CallInterfaces.h" @@ -21,10 +22,14 @@ #include "mlir/Pass/PassManager.h" #include "mlir/Support/DebugStringHelper.h" #include "mlir/Transforms/InliningUtils.h" -#include "mlir/Transforms/Passes.h" #include "llvm/ADT/SCCIterator.h" #include "llvm/Support/Debug.h" +namespace mlir { +#define GEN_PASS_DEF_INLINER +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + #define DEBUG_TYPE "inlining" using namespace mlir; @@ -582,7 +587,7 @@ static LogicalResult inlineCallsInSCC(Inliner &inliner, CGUseList &useList, //===----------------------------------------------------------------------===// namespace { -class InlinerPass : public InlinerBase<InlinerPass> { +class InlinerPass : public impl::InlinerBase<InlinerPass> { public: InlinerPass(); InlinerPass(const InlinerPass &) = default; diff --git a/mlir/lib/Transforms/LocationSnapshot.cpp b/mlir/lib/Transforms/LocationSnapshot.cpp index a042d07335bb..f020be2a2206 100644 --- a/mlir/lib/Transforms/LocationSnapshot.cpp +++ b/mlir/lib/Transforms/LocationSnapshot.cpp @@ -7,13 +7,19 @@ //===----------------------------------------------------------------------===// #include "mlir/Transforms/LocationSnapshot.h" -#include "PassDetail.h" + #include "mlir/IR/AsmState.h" #include "mlir/IR/Builders.h" +#include "mlir/Pass/Pass.h" #include "mlir/Support/FileUtilities.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/ToolOutputFile.h" +namespace mlir { +#define GEN_PASS_DEF_LOCATIONSNAPSHOT +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; /// This function generates new locations from the given IR by snapshotting the @@ -123,7 +129,7 @@ LogicalResult mlir::generateLocationsFromIR(StringRef fileName, StringRef tag, namespace { struct LocationSnapshotPass - : public LocationSnapshotBase<LocationSnapshotPass> { + : public impl::LocationSnapshotBase<LocationSnapshotPass> { LocationSnapshotPass() = default; LocationSnapshotPass(OpPrintingFlags flags, StringRef fileName, StringRef tag) : flags(flags) { diff --git a/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp b/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp index 35e0f48b2958..471193701718 100644 --- a/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp +++ b/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp @@ -10,18 +10,23 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/Interfaces/LoopLikeInterface.h" #include "mlir/Transforms/LoopInvariantCodeMotionUtils.h" -#include "mlir/Transforms/Passes.h" #include "mlir/Transforms/SideEffectUtils.h" +namespace mlir { +#define GEN_PASS_DEF_LOOPINVARIANTCODEMOTION +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { /// Loop invariant code motion (LICM) pass. struct LoopInvariantCodeMotion - : public LoopInvariantCodeMotionBase<LoopInvariantCodeMotion> { + : public impl::LoopInvariantCodeMotionBase<LoopInvariantCodeMotion> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Transforms/OpStats.cpp b/mlir/lib/Transforms/OpStats.cpp index d02cb8a3c895..6a71e1f02edc 100644 --- a/mlir/lib/Transforms/OpStats.cpp +++ b/mlir/lib/Transforms/OpStats.cpp @@ -6,19 +6,24 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/Operation.h" #include "mlir/IR/OperationSupport.h" -#include "mlir/Transforms/Passes.h" #include "llvm/ADT/DenseMap.h" #include "llvm/Support/Format.h" #include "llvm/Support/raw_ostream.h" +namespace mlir { +#define GEN_PASS_DEF_PRINTOPSTATS +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { -struct PrintOpStatsPass : public PrintOpStatsBase<PrintOpStatsPass> { +struct PrintOpStatsPass : public impl::PrintOpStatsBase<PrintOpStatsPass> { explicit PrintOpStatsPass(raw_ostream &os) : os(os) {} explicit PrintOpStatsPass(raw_ostream &os, bool printAsJSON) : os(os) { diff --git a/mlir/lib/Transforms/PassDetail.h b/mlir/lib/Transforms/PassDetail.h deleted file mode 100644 index 7c1f53929fe4..000000000000 --- a/mlir/lib/Transforms/PassDetail.h +++ /dev/null @@ -1,21 +0,0 @@ -//===- PassDetail.h - 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 TRANSFORMS_PASSDETAIL_H_ -#define TRANSFORMS_PASSDETAIL_H_ - -#include "mlir/Pass/Pass.h" -#include "mlir/Pass/PassManager.h" -#include "mlir/Transforms/Passes.h" - -namespace mlir { -#define GEN_PASS_CLASSES -#include "mlir/Transforms/Passes.h.inc" -} // namespace mlir - -#endif // TRANSFORMS_PASSDETAIL_H_ diff --git a/mlir/lib/Transforms/SCCP.cpp b/mlir/lib/Transforms/SCCP.cpp index a0a06352a0e8..385ff25aeda1 100644 --- a/mlir/lib/Transforms/SCCP.cpp +++ b/mlir/lib/Transforms/SCCP.cpp @@ -14,7 +14,8 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/Analysis/DataFlow/ConstantPropagationAnalysis.h" #include "mlir/Analysis/DataFlow/DeadCodeAnalysis.h" #include "mlir/IR/Builders.h" @@ -22,7 +23,11 @@ #include "mlir/Interfaces/SideEffectInterfaces.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/FoldUtils.h" -#include "mlir/Transforms/Passes.h" + +namespace mlir { +#define GEN_PASS_DEF_SCCP +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir using namespace mlir; using namespace mlir::dataflow; @@ -109,7 +114,7 @@ static void rewrite(DataFlowSolver &solver, MLIRContext *context, //===----------------------------------------------------------------------===// namespace { -struct SCCP : public SCCPBase<SCCP> { +struct SCCP : public impl::SCCPBase<SCCP> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Transforms/StripDebugInfo.cpp b/mlir/lib/Transforms/StripDebugInfo.cpp index 99e11517b7e0..87f5f76d3d43 100644 --- a/mlir/lib/Transforms/StripDebugInfo.cpp +++ b/mlir/lib/Transforms/StripDebugInfo.cpp @@ -6,16 +6,21 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/Operation.h" #include "mlir/Pass/Pass.h" -#include "mlir/Transforms/Passes.h" + +namespace mlir { +#define GEN_PASS_DEF_STRIPDEBUGINFO +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir using namespace mlir; namespace { -struct StripDebugInfo : public StripDebugInfoBase<StripDebugInfo> { +struct StripDebugInfo : public impl::StripDebugInfoBase<StripDebugInfo> { void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Transforms/SymbolDCE.cpp b/mlir/lib/Transforms/SymbolDCE.cpp index c197f5025c46..bf85cecaa46e 100644 --- a/mlir/lib/Transforms/SymbolDCE.cpp +++ b/mlir/lib/Transforms/SymbolDCE.cpp @@ -11,14 +11,19 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/IR/SymbolTable.h" #include "mlir/Transforms/Passes.h" +#include "mlir/IR/SymbolTable.h" + +namespace mlir { +#define GEN_PASS_DEF_SYMBOLDCE +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { -struct SymbolDCE : public SymbolDCEBase<SymbolDCE> { +struct SymbolDCE : public impl::SymbolDCEBase<SymbolDCE> { void runOnOperation() override; /// Compute the liveness of the symbols within the given symbol table. diff --git a/mlir/lib/Transforms/SymbolPrivatize.cpp b/mlir/lib/Transforms/SymbolPrivatize.cpp index 4aa7897da0ad..58fafa7e2b58 100644 --- a/mlir/lib/Transforms/SymbolPrivatize.cpp +++ b/mlir/lib/Transforms/SymbolPrivatize.cpp @@ -11,14 +11,19 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/IR/SymbolTable.h" #include "mlir/Transforms/Passes.h" +#include "mlir/IR/SymbolTable.h" + +namespace mlir { +#define GEN_PASS_DEF_SYMBOLPRIVATIZE +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { -struct SymbolPrivatize : public SymbolPrivatizeBase<SymbolPrivatize> { +struct SymbolPrivatize : public impl::SymbolPrivatizeBase<SymbolPrivatize> { explicit SymbolPrivatize(ArrayRef<std::string> excludeSymbols); LogicalResult initialize(MLIRContext *context) override; void runOnOperation() override; diff --git a/mlir/lib/Transforms/TopologicalSort.cpp b/mlir/lib/Transforms/TopologicalSort.cpp index afa0b78fbf25..04b52259abad 100644 --- a/mlir/lib/Transforms/TopologicalSort.cpp +++ b/mlir/lib/Transforms/TopologicalSort.cpp @@ -6,14 +6,21 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/IR/RegionKindInterface.h" #include "mlir/Transforms/TopologicalSortUtils.h" +namespace mlir { +#define GEN_PASS_DEF_TOPOLOGICALSORT +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; namespace { -struct TopologicalSortPass : public TopologicalSortBase<TopologicalSortPass> { +struct TopologicalSortPass + : public impl::TopologicalSortBase<TopologicalSortPass> { void runOnOperation() override { // Topologically sort the regions of the operation without SSA dominance. getOperation()->walk([](RegionKindInterface op) { diff --git a/mlir/lib/Transforms/ViewOpGraph.cpp b/mlir/lib/Transforms/ViewOpGraph.cpp index 1cd0164f6dd3..9afe2f5dc00c 100644 --- a/mlir/lib/Transforms/ViewOpGraph.cpp +++ b/mlir/lib/Transforms/ViewOpGraph.cpp @@ -7,15 +7,22 @@ //===----------------------------------------------------------------------===// #include "mlir/Transforms/ViewOpGraph.h" -#include "PassDetail.h" + #include "mlir/IR/Block.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Operation.h" +#include "mlir/Pass/Pass.h" #include "mlir/Support/IndentedOstream.h" +#include "llvm/ADT/StringMap.h" #include "llvm/Support/Format.h" #include "llvm/Support/GraphWriter.h" #include <utility> +namespace mlir { +#define GEN_PASS_DEF_VIEWOPGRAPH +#include "mlir/Transforms/Passes.h.inc" +} // namespace mlir + using namespace mlir; static const StringRef kLineStyleControlFlow = "dashed"; @@ -72,7 +79,7 @@ public: /// This pass generates a Graphviz dataflow visualization of an MLIR operation. /// Note: See https://www.graphviz.org/doc/info/lang.html for more information /// about the Graphviz DOT language. -class PrintOpPass : public ViewOpGraphBase<PrintOpPass> { +class PrintOpPass : public impl::ViewOpGraphBase<PrintOpPass> { public: PrintOpPass(raw_ostream &os) : os(os) {} PrintOpPass(const PrintOpPass &o) : PrintOpPass(o.os.getOStream()) {} |