diff --git a/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h b/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h --- a/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h +++ b/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h @@ -21,7 +21,12 @@ RewritePatternSet &patterns, amdgpu::Chipset chipset); -std::unique_ptr createConvertAMDGPUToROCDLPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertAMDGPUToROCDL +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/AffineToStandard/AffineToStandard.h b/mlir/include/mlir/Conversion/AffineToStandard/AffineToStandard.h --- a/mlir/include/mlir/Conversion/AffineToStandard/AffineToStandard.h +++ b/mlir/include/mlir/Conversion/AffineToStandard/AffineToStandard.h @@ -45,6 +45,13 @@ /// primitives). std::unique_ptr createLowerAffinePass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertAffineToStandard +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_AFFINETOSTANDARD_AFFINETOSTANDARD_H diff --git a/mlir/include/mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h b/mlir/include/mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h --- a/mlir/include/mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h +++ b/mlir/include/mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h @@ -17,12 +17,22 @@ class RewritePatternSet; class Pass; +struct ConvertArithmeticToLLVMPassOptions; + namespace arith { + void populateArithmeticToLLVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns); -std::unique_ptr createConvertArithmeticToLLVMPass(); } // namespace arith + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertArithmeticToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_ARITHMETICTOLLVM_ARITHMETICTOLLVM_H diff --git a/mlir/include/mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h b/mlir/include/mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h --- a/mlir/include/mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h +++ b/mlir/include/mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h @@ -18,12 +18,27 @@ class RewritePatternSet; class Pass; +struct ConvertArithmeticToSPIRVPassOptions; + namespace arith { + void populateArithmeticToSPIRVPatterns(SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns); -std::unique_ptr> createConvertArithmeticToSPIRVPass(); } // namespace arith + +std::unique_ptr> createConvertArithmeticToSPIRVPass(); + +std::unique_ptr> createConvertArithmeticToSPIRVPass( + const ConvertArithmeticToSPIRVPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertArithmeticToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_ARITHMETICTOSPIRV_ARITHMETICTOSPIRV_H diff --git a/mlir/include/mlir/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.h b/mlir/include/mlir/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.h --- a/mlir/include/mlir/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.h +++ b/mlir/include/mlir/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.h @@ -16,10 +16,12 @@ /// See createConvertArmNeon2dToIntrPass. void populateConvertArmNeon2dToIntrPatterns(RewritePatternSet &patterns); -/// Creates a pass to lower Arm NEON 2D ops to intrinsics, i.e. -/// equivalent ops operating on flattened 1D vectors and mapping more -/// directly to the corresponding Arm NEON instruction. -std::unique_ptr createConvertArmNeon2dToIntrPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertArmNeon2dToIntr +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h b/mlir/include/mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h --- a/mlir/include/mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h +++ b/mlir/include/mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h @@ -36,6 +36,13 @@ TypeConverter &typeConverter, RewritePatternSet &patterns, ConversionTarget &target); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertAsyncToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_ASYNCTOLLVM_ASYNCTOLLVM_H diff --git a/mlir/include/mlir/Conversion/BufferizationToMemRef/BufferizationToMemRef.h b/mlir/include/mlir/Conversion/BufferizationToMemRef/BufferizationToMemRef.h --- a/mlir/include/mlir/Conversion/BufferizationToMemRef/BufferizationToMemRef.h +++ b/mlir/include/mlir/Conversion/BufferizationToMemRef/BufferizationToMemRef.h @@ -20,7 +20,13 @@ void populateBufferizationToMemRefConversionPatterns( RewritePatternSet &patterns); -std::unique_ptr createBufferizationToMemRefPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertBufferizationToMemRef +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_BUFFERIZATIONTOMEMREF_BUFFERIZATIONTOMEMREF_H diff --git a/mlir/include/mlir/Conversion/CMakeLists.txt b/mlir/include/mlir/Conversion/CMakeLists.txt --- a/mlir/include/mlir/Conversion/CMakeLists.txt +++ b/mlir/include/mlir/Conversion/CMakeLists.txt @@ -1,6 +1,7 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Conversion) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Conversion) mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix Conversion) mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix Conversion) add_public_tablegen_target(MLIRConversionPassIncGen) diff --git a/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h b/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h --- a/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h +++ b/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h @@ -38,8 +38,12 @@ void populateComplexToLLVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns); -/// Create a pass to convert Complex operations to the LLVMIR dialect. -std::unique_ptr createConvertComplexToLLVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertComplexToLLVM +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/ComplexToLibm/ComplexToLibm.h b/mlir/include/mlir/Conversion/ComplexToLibm/ComplexToLibm.h --- a/mlir/include/mlir/Conversion/ComplexToLibm/ComplexToLibm.h +++ b/mlir/include/mlir/Conversion/ComplexToLibm/ComplexToLibm.h @@ -22,6 +22,13 @@ /// Create a pass to convert Complex operations to libm calls. std::unique_ptr> createConvertComplexToLibmPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertComplexToLibm +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_COMPLEXTOLIBM_COMPLEXTOLIBM_H_ diff --git a/mlir/include/mlir/Conversion/ComplexToStandard/ComplexToStandard.h b/mlir/include/mlir/Conversion/ComplexToStandard/ComplexToStandard.h --- a/mlir/include/mlir/Conversion/ComplexToStandard/ComplexToStandard.h +++ b/mlir/include/mlir/Conversion/ComplexToStandard/ComplexToStandard.h @@ -17,8 +17,12 @@ /// Populate the given list with patterns that convert from Complex to Standard. void populateComplexToStandardConversionPatterns(RewritePatternSet &patterns); -/// Create a pass to convert Complex operations to the Standard dialect. -std::unique_ptr createConvertComplexToStandardPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertComplexToStandard +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h b/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h --- a/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h +++ b/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h @@ -20,16 +20,25 @@ class RewritePatternSet; class Pass; +struct ConvertControlFlowToLLVMPassOptions; + namespace cf { + /// Collect the patterns to convert from the ControlFlow dialect to LLVM. The /// conversion patterns capture the LLVMTypeConverter by reference meaning the /// references have to remain alive during the entire pattern lifetime. void populateControlFlowToLLVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns); -/// Creates a pass to convert the ControlFlow dialect into the LLVMIR dialect. -std::unique_ptr createConvertControlFlowToLLVMPass(); } // namespace cf + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertControlFlowToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_CONTROLFLOWTOLLVM_CONTROLFLOWTOLLVM_H diff --git a/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.h b/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.h --- a/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.h @@ -18,9 +18,22 @@ namespace mlir { class ModuleOp; +struct ConvertControlFlowToSPIRVPassOptions; + /// Creates a pass to convert ControlFlow ops to SPIR-V ops. std::unique_ptr> createConvertControlFlowToSPIRVPass(); +/// Creates a pass to convert ControlFlow ops to SPIR-V ops. +std::unique_ptr> createConvertControlFlowToSPIRVPass( + const ConvertControlFlowToSPIRVPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertControlFlowToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_CONTROLFLOWTOSPIRV_CONTROLFLOWTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h b/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h --- a/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h +++ b/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h @@ -10,6 +10,7 @@ #define MLIR_CONVERSION_FUNCTOLLVM_CONVERTFUNCTOLLVMPASS_H_ #include +#include namespace mlir { class LowerToLLVMOptions; @@ -18,10 +19,20 @@ class OperationPass; class Pass; +struct ConvertFuncToLLVMPassOptions; + /// Creates a pass to convert the Func dialect into the LLVMIR dialect. std::unique_ptr> createConvertFuncToLLVMPass(); + std::unique_ptr> -createConvertFuncToLLVMPass(const LowerToLLVMOptions &options); +createConvertFuncToLLVMPass(const ConvertFuncToLLVMPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertFuncToLLVM +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRVPass.h b/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRVPass.h --- a/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRVPass.h @@ -18,9 +18,22 @@ namespace mlir { class ModuleOp; +struct ConvertFuncToSPIRVPassOptions; + /// Creates a pass to convert Func ops to SPIR-V ops. std::unique_ptr> createConvertFuncToSPIRVPass(); +/// Creates a pass to convert Func ops to SPIR-V ops. +std::unique_ptr> +createConvertFuncToSPIRVPass(const ConvertFuncToSPIRVPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertFuncToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_FUNCTOSPIRV_FUNCTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h b/mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h --- a/mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h +++ b/mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h @@ -30,6 +30,8 @@ template class OperationPass; +struct GpuToLLVMConversionPassPassOptions; + namespace gpu { class GPUModuleOp; } // namespace gpu @@ -50,8 +52,10 @@ /// This pass does not generate code to call GPU runtime APIs directly but /// instead uses a small wrapper library that exports a stable and conveniently /// typed ABI on top of GPU runtimes such as CUDA or ROCm (HIP). -std::unique_ptr> -createGpuToLLVMConversionPass(bool kernelBarePtrCallConv = false); +std::unique_ptr> createGpuToLLVMConversionPass(); + +std::unique_ptr> createGpuToLLVMConversionPass( + const GpuToLLVMConversionPassPassOptions &options); /// Collect a set of patterns to convert from the GPU dialect to LLVM and /// populate converter for gpu types. @@ -60,6 +64,13 @@ StringRef gpuBinaryAnnotation = {}, bool kernelBarePtrCallConv = false); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_GpuToLLVMConversionPass +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_GPUCOMMON_GPUCOMMONPASS_H_ diff --git a/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h b/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h --- a/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h +++ b/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h @@ -20,6 +20,8 @@ template class OperationPass; +struct ConvertGpuOpsToNVVMOpsPassOptions; + namespace gpu { class GPUModuleOp; class MMAMatrixType; @@ -38,11 +40,20 @@ void populateGpuWMMAToNVVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns); -/// Creates a pass that lowers GPU dialect operations to NVVM counterparts. The -/// index bitwidth used for the lowering of the device side index computations -/// is configurable. +/// Creates a pass that lowers GPU dialect operations to NVVM counterparts. +std::unique_ptr> +createLowerGpuOpsToNVVMOpsPass(); + +/// Creates a pass that lowers GPU dialect operations to NVVM counterparts. std::unique_ptr> createLowerGpuOpsToNVVMOpsPass( - unsigned indexBitwidth = kDeriveIndexBitwidthFromDataLayout); + const ConvertGpuOpsToNVVMOpsPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertGpuOpsToNVVMOps +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h b/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h --- a/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h +++ b/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h @@ -20,6 +20,8 @@ template class OperationPass; +struct ConvertGpuOpsToROCDLOpsPassOptions; + namespace gpu { class GPUModuleOp; } // namespace gpu @@ -34,15 +36,21 @@ /// Configure target to convert from the GPU dialect to ROCDL. void configureGpuToROCDLConversionLegality(ConversionTarget &target); -/// Creates a pass that lowers GPU dialect operations to ROCDL counterparts. The -/// index bitwidth used for the lowering of the device side index computations -/// is configurable. +/// Creates a pass that lowers GPU dialect operations to ROCDL counterparts. +std::unique_ptr> +createLowerGpuOpsToROCDLOpsPass(); + +/// Creates a pass that lowers GPU dialect operations to ROCDL counterparts. std::unique_ptr> createLowerGpuOpsToROCDLOpsPass( - const std::string &chipset = "gfx900", - unsigned indexBitwidth = kDeriveIndexBitwidthFromDataLayout, - bool useBarePtrCallConv = false, - gpu::amd::Runtime runtime = gpu::amd::Runtime::Unknown); + const ConvertGpuOpsToROCDLOpsPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertGpuOpsToROCDLOps +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h b/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h --- a/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h @@ -21,12 +21,23 @@ template class OperationPass; +struct ConvertGPUToSPIRVPassOptions; + /// Creates a pass to convert GPU kernel ops to corresponding SPIR-V ops. For a /// gpu.func to be converted, it should have a spv.entry_point_abi attribute. -/// If `mapMemorySpace` is true, performs MemRef memory space to SPIR-V mapping -/// according to default Vulkan rules first. +/// If the `mapMemorySpace` option is true, performs MemRef memory space to +/// SPIR-V mapping according to default Vulkan rules first. +std::unique_ptr> createConvertGPUToSPIRVPass(); + std::unique_ptr> -createConvertGPUToSPIRVPass(bool mapMemorySpace = false); +createConvertGPUToSPIRVPass(const ConvertGPUToSPIRVPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertGPUToSPIRV +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir #endif // MLIR_CONVERSION_GPUTOSPIRV_GPUTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h b/mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h --- a/mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h +++ b/mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h @@ -30,5 +30,13 @@ std::unique_ptr> createConvertGpuLaunchFuncToVulkanLaunchFuncPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertVulkanLaunchFuncToVulkanCalls +#define GEN_PASS_DECL_ConvertGpuLaunchFuncToVulkanLaunchFunc +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_GPUTOVULKAN_CONVERTGPUTOVULKANPASS_H diff --git a/mlir/include/mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h b/mlir/include/mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h --- a/mlir/include/mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h +++ b/mlir/include/mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h @@ -25,6 +25,13 @@ /// Create a pass to convert Linalg operations to the LLVMIR dialect. std::unique_ptr> createConvertLinalgToLLVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertLinalgToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_LINALGTOLLVM_LINALGTOLLVM_H_ diff --git a/mlir/include/mlir/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.h b/mlir/include/mlir/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.h --- a/mlir/include/mlir/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.h @@ -21,6 +21,13 @@ /// Creates and returns a pass to convert Linalg ops to SPIR-V ops. std::unique_ptr> createLinalgToSPIRVPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertLinalgToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_LINALGTOSPIRV_LINALGTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/LinalgToStandard/LinalgToStandard.h b/mlir/include/mlir/Conversion/LinalgToStandard/LinalgToStandard.h --- a/mlir/include/mlir/Conversion/LinalgToStandard/LinalgToStandard.h +++ b/mlir/include/mlir/Conversion/LinalgToStandard/LinalgToStandard.h @@ -47,6 +47,13 @@ /// Create a pass to convert Linalg operations to the Standard dialect. std::unique_ptr> createConvertLinalgToStandardPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertLinalgToStandard +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_LINALGTOSTANDARD_LINALGTOSTANDARD_H_ diff --git a/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h b/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h --- a/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h +++ b/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h @@ -20,7 +20,13 @@ void populateMathToLLVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns); -std::unique_ptr createConvertMathToLLVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertMathToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_MATHTOLLVM_MATHTOLLVM_H diff --git a/mlir/include/mlir/Conversion/MathToLibm/MathToLibm.h b/mlir/include/mlir/Conversion/MathToLibm/MathToLibm.h --- a/mlir/include/mlir/Conversion/MathToLibm/MathToLibm.h +++ b/mlir/include/mlir/Conversion/MathToLibm/MathToLibm.h @@ -21,6 +21,13 @@ /// Create a pass to convert Math operations to libm calls. std::unique_ptr> createConvertMathToLibmPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertMathToLibm +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_MATHTOLIBM_MATHTOLIBM_H_ diff --git a/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRVPass.h b/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRVPass.h --- a/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRVPass.h @@ -21,6 +21,13 @@ /// Creates a pass to convert Math ops to SPIR-V ops. std::unique_ptr> createConvertMathToSPIRVPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertMathToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_MATHTOSPIRV_MATHTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h b/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h --- a/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h +++ b/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h @@ -21,7 +21,13 @@ void populateMemRefToLLVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns); -std::unique_ptr createMemRefToLLVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertMemRefToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_MEMREFTOLLVM_MEMREFTOLLVM_H diff --git a/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.h b/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.h --- a/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.h @@ -19,6 +19,8 @@ namespace mlir { class ModuleOp; +struct ConvertMemRefToSPIRVPassOptions; + /// Creates a pass to map numeric MemRef memory spaces to symbolic SPIR-V /// storage classes. The mapping is read from the command-line option. std::unique_ptr> createMapMemRefStorageClassPass(); @@ -26,6 +28,18 @@ /// Creates a pass to convert MemRef ops to SPIR-V ops. std::unique_ptr> createConvertMemRefToSPIRVPass(); +/// Creates a pass to convert MemRef ops to SPIR-V ops. +std::unique_ptr> +createConvertMemRefToSPIRVPass(const ConvertMemRefToSPIRVPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertMemRefToSPIRV +#define GEN_PASS_DECL_MapMemRefStorageClass +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_MEMREFTOSPIRV_MEMREFTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h b/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h --- a/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h +++ b/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h @@ -19,7 +19,12 @@ void populateNVGPUToNVVMConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns); -std::unique_ptr createConvertNVGPUToNVVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertNVGPUToNVVM +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/OpenACCToLLVM/ConvertOpenACCToLLVM.h b/mlir/include/mlir/Conversion/OpenACCToLLVM/ConvertOpenACCToLLVM.h --- a/mlir/include/mlir/Conversion/OpenACCToLLVM/ConvertOpenACCToLLVM.h +++ b/mlir/include/mlir/Conversion/OpenACCToLLVM/ConvertOpenACCToLLVM.h @@ -70,6 +70,13 @@ /// Create a pass to convert the OpenACC dialect into the LLVMIR dialect. std::unique_ptr> createConvertOpenACCToLLVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertOpenACCToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_OPENACCTOLLVM_CONVERTOPENACCTOLLVM_H diff --git a/mlir/include/mlir/Conversion/OpenACCToSCF/ConvertOpenACCToSCF.h b/mlir/include/mlir/Conversion/OpenACCToSCF/ConvertOpenACCToSCF.h --- a/mlir/include/mlir/Conversion/OpenACCToSCF/ConvertOpenACCToSCF.h +++ b/mlir/include/mlir/Conversion/OpenACCToSCF/ConvertOpenACCToSCF.h @@ -23,6 +23,13 @@ /// Create a pass to convert the OpenACC dialect into the LLVMIR dialect. std::unique_ptr> createConvertOpenACCToSCFPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertOpenACCToSCF +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_OPENACCTOSCF_CONVERTOPENACCTOSCF_H diff --git a/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h b/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h --- a/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h +++ b/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h @@ -31,6 +31,13 @@ /// Create a pass to convert OpenMP operations to the LLVMIR dialect. std::unique_ptr> createConvertOpenMPToLLVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertOpenMPToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_OPENMPTOLLVM_CONVERTOPENMPTOLLVM_H diff --git a/mlir/include/mlir/Conversion/PDLToPDLInterp/PDLToPDLInterp.h b/mlir/include/mlir/Conversion/PDLToPDLInterp/PDLToPDLInterp.h --- a/mlir/include/mlir/Conversion/PDLToPDLInterp/PDLToPDLInterp.h +++ b/mlir/include/mlir/Conversion/PDLToPDLInterp/PDLToPDLInterp.h @@ -23,6 +23,13 @@ /// Creates and returns a pass to convert PDL ops to PDL interpreter ops. std::unique_ptr> createPDLToPDLInterpPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertPDLToPDLInterp +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_PDLTOPDLINTERP_PDLTOPDLINTERP_H diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td --- a/mlir/include/mlir/Conversion/Passes.td +++ b/mlir/include/mlir/Conversion/Passes.td @@ -82,7 +82,6 @@ let description = [{ This pass converts supported AMDGPU ops to ROCDL dialect intrinsics. }]; - let constructor = "mlir::createConvertAMDGPUToROCDLPass()"; let dependentDialects = [ "LLVM::LLVMDialect", "ROCDL::ROCDLDialect", @@ -101,7 +100,6 @@ let description = [{ This pass converts supported Arithmetic ops to LLVM dialect instructions. }]; - let constructor = "mlir::arith::createConvertArithmeticToLLVMPass()"; let dependentDialects = ["LLVM::LLVMDialect"]; let options = [ Option<"indexBitwidth", "index-bitwidth", "unsigned", @@ -116,7 +114,7 @@ def ConvertArithmeticToSPIRV : Pass<"convert-arith-to-spirv"> { let summary = "Convert Arithmetic dialect to SPIR-V dialect"; - let constructor = "mlir::arith::createConvertArithmeticToSPIRVPass()"; + let constructor = "mlir::createConvertArithmeticToSPIRVPass()"; let dependentDialects = ["spirv::SPIRVDialect"]; let options = [ Option<"emulateNon32BitScalarTypes", "emulate-non-32-bit-scalar-types", @@ -132,7 +130,11 @@ def ConvertArmNeon2dToIntr : Pass<"arm-neon-2d-to-intr"> { let summary = "Convert Arm NEON structured ops to intrinsics"; - let constructor = "mlir::createConvertArmNeon2dToIntrPass()"; + let description = [{ + Lower Arm NEON 2D ops to intrinsics, i.e. equivalent ops operating on + flattened 1D vectors and mapping more directly to the corresponding Arm + NEON instruction. + }]; let dependentDialects = ["arm_neon::ArmNeonDialect", "vector::VectorDialect"]; } @@ -185,7 +187,6 @@ and hence does not resolve any memory leaks. }]; - let constructor = "mlir::createBufferizationToMemRefPass()"; let dependentDialects = ["arith::ArithmeticDialect", "memref::MemRefDialect"]; } @@ -195,7 +196,6 @@ def ConvertComplexToLLVM : Pass<"convert-complex-to-llvm"> { let summary = "Convert Complex dialect to LLVM dialect"; - let constructor = "mlir::createConvertComplexToLLVMPass()"; let dependentDialects = ["LLVM::LLVMDialect"]; } @@ -220,7 +220,6 @@ def ConvertComplexToStandard : Pass<"convert-complex-to-standard"> { let summary = "Convert Complex dialect to standard dialect"; - let constructor = "mlir::createConvertComplexToStandardPass()"; let dependentDialects = ["math::MathDialect"]; } @@ -237,7 +236,6 @@ IR dialect operations, the pass will fail. Any LLVM IR operations or types already present in the IR will be kept as is. }]; - let constructor = "mlir::cf::createConvertControlFlowToLLVMPass()"; let dependentDialects = ["LLVM::LLVMDialect"]; let options = [ Option<"indexBitwidth", "index-bitwidth", "unsigned", @@ -331,6 +329,10 @@ let summary = "Convert GPU dialect to LLVM dialect with GPU runtime calls"; let constructor = "mlir::createGpuToLLVMConversionPass()"; let dependentDialects = ["LLVM::LLVMDialect"]; + let options = [ + Option<"kernelBarePtrCallConv", "kernel-bare-ptr-call-conv", "bool", + /*default=*/"false", "Use bare pointer calling convention">, + ]; } def LowerHostCodeToLLVM : Pass<"lower-host-to-llvm", "ModuleOp"> { @@ -408,6 +410,12 @@ }]; let constructor = "mlir::createConvertGPUToSPIRVPass()"; let dependentDialects = ["spirv::SPIRVDialect"]; + let options = [ + Option<"mapMemorySpace", "map-memory-space", + "bool", /*default=*/"false", + "Performs MemRef memory space to SPIR-V mapping according to default" + " Vulkan rules first"> + ]; } //===----------------------------------------------------------------------===// @@ -497,7 +505,6 @@ let description = [{ This pass converts supported Math ops to LLVM dialect intrinsics. }]; - let constructor = "mlir::createConvertMathToLLVMPass()"; let dependentDialects = ["LLVM::LLVMDialect"]; } @@ -518,7 +525,6 @@ def ConvertMemRefToLLVM : Pass<"convert-memref-to-llvm", "ModuleOp"> { let summary = "Convert operations from the MemRef dialect to the LLVM " "dialect"; - let constructor = "mlir::createMemRefToLLVMPass()"; let dependentDialects = ["LLVM::LLVMDialect"]; let options = [ Option<"useAlignedAlloc", "use-aligned-alloc", "bool", /*default=*/"false", @@ -568,7 +574,6 @@ let description = [{ This pass converts supported NVGPU ops to NVVM dialect intrinsics. }]; - let constructor = "mlir::createConvertNVGPUToNVVMPass()"; let dependentDialects = [ "NVVM::NVVMDialect", ]; @@ -638,17 +643,15 @@ and the producer operation is converted by another pass, each of which produces an unrealized cast. This pass can be used to clean up the IR. }]; - let constructor = "mlir::createReconcileUnrealizedCastsPass()"; } //===----------------------------------------------------------------------===// // SCFToControlFlow //===----------------------------------------------------------------------===// -def SCFToControlFlow : Pass<"convert-scf-to-cf"> { +def ConvertSCFToControlFlow : Pass<"convert-scf-to-cf"> { let summary = "Convert SCF dialect to ControlFlow dialect, replacing structured" " control flow with a CFG"; - let constructor = "mlir::createConvertSCFToCFPass()"; let dependentDialects = ["cf::ControlFlowDialect"]; } @@ -728,7 +731,6 @@ can happen at a different part of the program than general shape computation lowering. }]; - let constructor = "mlir::createConvertShapeConstraintsPass()"; let dependentDialects = ["cf::ControlFlowDialect", "scf::SCFDialect"]; } @@ -872,7 +874,6 @@ def ConvertVectorToGPU : Pass<"convert-vector-to-gpu"> { let summary = "Lower the operations from the vector dialect into the GPU " "dialect"; - let constructor = "mlir::createConvertVectorToGPUPass()"; let dependentDialects = [ "memref::MemRefDialect", "gpu::GPUDialect", "AffineDialect", "vector::VectorDialect", "nvgpu::NVGPUDialect" diff --git a/mlir/include/mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h b/mlir/include/mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h --- a/mlir/include/mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h +++ b/mlir/include/mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h @@ -15,13 +15,17 @@ class Pass; class RewritePatternSet; -/// Creates a pass that eliminates noop `unrealized_conversion_cast` operation -/// sequences. -std::unique_ptr createReconcileUnrealizedCastsPass(); - /// Populates `patterns` with rewrite patterns that eliminate noop /// `unrealized_conversion_cast` operation sequences. void populateReconcileUnrealizedCastsPatterns(RewritePatternSet &patterns); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ReconcileUnrealizedCasts +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_RECONCILEUNREALIZEDCASTS_RECONCILEUNREALIZEDCASTS_H_ diff --git a/mlir/include/mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h b/mlir/include/mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h --- a/mlir/include/mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h +++ b/mlir/include/mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h @@ -19,9 +19,12 @@ /// operations within the ControlFlow dialect. void populateSCFToControlFlowConversionPatterns(RewritePatternSet &patterns); -/// Creates a pass to convert SCF operations to CFG branch-based operation in -/// the ControlFlow dialect. -std::unique_ptr createConvertSCFToCFPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertSCFToControlFlow +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/SCFToGPU/SCFToGPUPass.h b/mlir/include/mlir/Conversion/SCFToGPU/SCFToGPUPass.h --- a/mlir/include/mlir/Conversion/SCFToGPU/SCFToGPUPass.h +++ b/mlir/include/mlir/Conversion/SCFToGPU/SCFToGPUPass.h @@ -18,6 +18,8 @@ class InterfacePass; class Pass; +struct ConvertAffineForToGPUPassOptions; + /// Create a pass that converts loop nests into GPU kernels. It considers /// top-level affine.for operations as roots of loop nests and converts them to /// the gpu.launch operations if possible. @@ -26,16 +28,25 @@ /// parallelization is performed, it is under the responsibility of the caller /// to strip-mine the loops and to perform the dependence analysis before /// calling the conversion. -std::unique_ptr> -createAffineForToGPUPass(unsigned numBlockDims, unsigned numThreadDims); std::unique_ptr> createAffineForToGPUPass(); +std::unique_ptr> +createAffineForToGPUPass(const ConvertAffineForToGPUPassOptions &options); + /// Creates a pass that converts scf.parallel operations into a gpu.launch /// operation. The mapping of loop dimensions to launch dimensions is derived /// from mapping attributes. See ParallelToGpuLaunchLowering::matchAndRewrite /// for a description of the used attributes. std::unique_ptr createParallelLoopToGpuPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertAffineForToGPU +#define GEN_PASS_DECL_ConvertParallelLoopToGpu +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_SCFTOGPU_SCFTOGPUPASS_H_ diff --git a/mlir/include/mlir/Conversion/SCFToOpenMP/SCFToOpenMP.h b/mlir/include/mlir/Conversion/SCFToOpenMP/SCFToOpenMP.h --- a/mlir/include/mlir/Conversion/SCFToOpenMP/SCFToOpenMP.h +++ b/mlir/include/mlir/Conversion/SCFToOpenMP/SCFToOpenMP.h @@ -18,6 +18,13 @@ std::unique_ptr> createConvertSCFToOpenMPPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertSCFToOpenMP +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_SCFTOOPENMP_SCFTOOPENMP_H diff --git a/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRVPass.h b/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRVPass.h --- a/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRVPass.h @@ -21,6 +21,13 @@ /// Creates a pass to convert SCF ops into SPIR-V ops. std::unique_ptr> createConvertSCFToSPIRVPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_SCFToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_SCFTOSPIRV_SCFTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h b/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h --- a/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h +++ b/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h @@ -33,6 +33,14 @@ /// Creates a pass to convert SPIR-V operations to the LLVMIR dialect. std::unique_ptr> createConvertSPIRVToLLVMPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_LowerHostCodeToLLVM +#define GEN_PASS_DECL_ConvertSPIRVToLLVM +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_SPIRVTOLLVM_SPIRVTOLLVMPASS_H diff --git a/mlir/include/mlir/Conversion/ShapeToStandard/ShapeToStandard.h b/mlir/include/mlir/Conversion/ShapeToStandard/ShapeToStandard.h --- a/mlir/include/mlir/Conversion/ShapeToStandard/ShapeToStandard.h +++ b/mlir/include/mlir/Conversion/ShapeToStandard/ShapeToStandard.h @@ -26,7 +26,13 @@ void populateConvertShapeConstraintsConversionPatterns( RewritePatternSet &patterns); -std::unique_ptr createConvertShapeConstraintsPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertShapeToStandard +#define GEN_PASS_DECL_ConvertShapeConstraints +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/TensorToLinalg/TensorToLinalgPass.h b/mlir/include/mlir/Conversion/TensorToLinalg/TensorToLinalgPass.h --- a/mlir/include/mlir/Conversion/TensorToLinalg/TensorToLinalgPass.h +++ b/mlir/include/mlir/Conversion/TensorToLinalg/TensorToLinalgPass.h @@ -21,6 +21,13 @@ /// Creates a pass to convert Tensor ops to Linalg ops. std::unique_ptr> createConvertTensorToLinalgPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertTensorToLinalg +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_TENSORTOLINALG_TENSORTOLINALGPASS_H diff --git a/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRVPass.h b/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRVPass.h --- a/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRVPass.h @@ -18,9 +18,22 @@ namespace mlir { class ModuleOp; +struct ConvertTensorToSPIRVPassOptions; + /// Creates a pass to convert Tensor ops to SPIR-V ops. std::unique_ptr> createConvertTensorToSPIRVPass(); +/// Creates a pass to convert Tensor ops to SPIR-V ops. +std::unique_ptr> +createConvertTensorToSPIRVPass(const ConvertTensorToSPIRVPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertTensorToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_TENSORTOSPIRV_TENSORTOSPIRVPASS_H diff --git a/mlir/include/mlir/Conversion/TosaToArith/TosaToArith.h b/mlir/include/mlir/Conversion/TosaToArith/TosaToArith.h --- a/mlir/include/mlir/Conversion/TosaToArith/TosaToArith.h +++ b/mlir/include/mlir/Conversion/TosaToArith/TosaToArith.h @@ -16,16 +16,29 @@ #include "mlir/Pass/Pass.h" namespace mlir { + +struct TosaToArithPassOptions; + namespace tosa { std::unique_ptr createTosaToArith(); +std::unique_ptr createTosaToArith(const TosaToArithPassOptions &options); + void populateTosaToArithConversionPatterns(RewritePatternSet *patterns); void populateTosaRescaleToArithConversionPatterns(RewritePatternSet *patterns, bool include32Bit = false); } // namespace tosa + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_TosaToArith +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_TOSATOARITH_TOSATOARITH_H diff --git a/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h b/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h --- a/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h +++ b/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h @@ -36,6 +36,15 @@ void populateTosaToLinalgNamedConversionPatterns(RewritePatternSet *patterns); } // namespace tosa + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_TosaToLinalg +#define GEN_PASS_DECL_TosaToLinalgNamed +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_TOSATOLINALG_TOSATOLINALG_H diff --git a/mlir/include/mlir/Conversion/TosaToSCF/TosaToSCF.h b/mlir/include/mlir/Conversion/TosaToSCF/TosaToSCF.h --- a/mlir/include/mlir/Conversion/TosaToSCF/TosaToSCF.h +++ b/mlir/include/mlir/Conversion/TosaToSCF/TosaToSCF.h @@ -26,6 +26,14 @@ void addTosaToSCFPasses(OpPassManager &pm); } // namespace tosa + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_TosaToSCF +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_TOSATOSCF_TOSATOSCF_H diff --git a/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h b/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h --- a/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h +++ b/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h @@ -23,6 +23,14 @@ void populateTosaToTensorConversionPatterns(RewritePatternSet *patterns); } // namespace tosa + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_TosaToTensor +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_TOSATOTENSOR_TOSATOTENSOR_H diff --git a/mlir/include/mlir/Conversion/VectorToGPU/VectorToGPU.h b/mlir/include/mlir/Conversion/VectorToGPU/VectorToGPU.h --- a/mlir/include/mlir/Conversion/VectorToGPU/VectorToGPU.h +++ b/mlir/include/mlir/Conversion/VectorToGPU/VectorToGPU.h @@ -34,8 +34,12 @@ /// the vector operations are left untouched. LogicalResult convertVectorToNVVMCompatibleMMASync(Operation *rootOp); -/// Convert from vector to GPU ops. -std::unique_ptr createConvertVectorToGPUPass(bool useNvGpu = false); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertVectorToGPU +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h b/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h --- a/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h +++ b/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h @@ -16,45 +16,7 @@ template class OperationPass; -/// Options to control Vector to LLVM lowering. -/// -/// This should kept in sync with VectorToLLVM options defined for the -/// ConvertVectorToLLVM pass in include/mlir/Conversion/Passes.td -struct LowerVectorToLLVMOptions { - LowerVectorToLLVMOptions() {} - - LowerVectorToLLVMOptions &enableReassociateFPReductions(bool b = true) { - reassociateFPReductions = b; - return *this; - } - LowerVectorToLLVMOptions &enableIndexOptimizations(bool b = true) { - force32BitVectorIndices = b; - return *this; - } - LowerVectorToLLVMOptions &enableArmNeon(bool b = true) { - armNeon = b; - return *this; - } - LowerVectorToLLVMOptions &enableArmSVE(bool b = true) { - armSVE = b; - return *this; - } - LowerVectorToLLVMOptions &enableAMX(bool b = true) { - amx = b; - return *this; - } - LowerVectorToLLVMOptions &enableX86Vector(bool b = true) { - x86Vector = b; - return *this; - } - - bool reassociateFPReductions{false}; - bool force32BitVectorIndices{true}; - bool armNeon{false}; - bool armSVE{false}; - bool amx{false}; - bool x86Vector{false}; -}; +struct ConvertVectorToLLVMPassOptions; /// Collect a set of patterns to convert from Vector contractions to LLVM Matrix /// Intrinsics. To lower to assembly, the LLVM flag -lower-matrix-intrinsics @@ -68,8 +30,17 @@ bool reassociateFPReductions = false, bool force32BitVectorIndices = false); /// Create a pass to convert vector operations to the LLVMIR dialect. -std::unique_ptr> createConvertVectorToLLVMPass( - const LowerVectorToLLVMOptions &options = LowerVectorToLLVMOptions()); +std::unique_ptr> createConvertVectorToLLVMPass(); + +std::unique_ptr> +createConvertVectorToLLVMPass(const ConvertVectorToLLVMPassOptions &options); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertVectorToLLVM +#include "mlir/Conversion/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Conversion/VectorToSCF/VectorToSCF.h b/mlir/include/mlir/Conversion/VectorToSCF/VectorToSCF.h --- a/mlir/include/mlir/Conversion/VectorToSCF/VectorToSCF.h +++ b/mlir/include/mlir/Conversion/VectorToSCF/VectorToSCF.h @@ -83,6 +83,13 @@ std::unique_ptr createConvertVectorToSCFPass( const VectorTransferToSCFOptions &options = VectorTransferToSCFOptions()); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertVectorToSCF +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_VECTORTOSCF_VECTORTOSCF_H_ diff --git a/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRVPass.h b/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRVPass.h --- a/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRVPass.h +++ b/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRVPass.h @@ -21,6 +21,13 @@ /// Creates a pass to convert Vector Ops to SPIR-V ops. std::unique_ptr> createConvertVectorToSPIRVPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertVectorToSPIRV +#include "mlir/Conversion/Passes.h.inc" + } // namespace mlir #endif // MLIR_CONVERSION_VECTORTOSPIRV_VECTORTOSPIRVPASS_H diff --git a/mlir/include/mlir/Dialect/Affine/CMakeLists.txt b/mlir/include/mlir/Dialect/Affine/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Affine/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Affine/CMakeLists.txt @@ -2,6 +2,7 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Affine) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Affine) add_public_tablegen_target(MLIRAffinePassIncGen) add_mlir_doc(Passes AffinePasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/Affine/Passes.h b/mlir/include/mlir/Dialect/Affine/Passes.h --- a/mlir/include/mlir/Dialect/Affine/Passes.h +++ b/mlir/include/mlir/Dialect/Affine/Passes.h @@ -110,6 +110,25 @@ /// Overload relying on pass options for initialization. std::unique_ptr> createSuperVectorizePass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_AffineDataCopyGeneration +#define GEN_PASS_DECL_AffineLoopFusion +#define GEN_PASS_DECL_AffineLoopInvariantCodeMotion +#define GEN_PASS_DECL_AffineLoopNormalize +#define GEN_PASS_DECL_AffineLoopTiling +#define GEN_PASS_DECL_AffineLoopUnroll +#define GEN_PASS_DECL_AffineLoopUnrollAndJam +#define GEN_PASS_DECL_AffineParallelize +#define GEN_PASS_DECL_AffineScalarReplacement +#define GEN_PASS_DECL_AffineVectorize +#define GEN_PASS_DECL_LoopCoalescing +#define GEN_PASS_DECL_AffinePipelineDataTransfer +#define GEN_PASS_DECL_SimplifyAffineStructures +#include "mlir/Dialect/Affine/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/Arithmetic/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Arithmetic/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Arithmetic/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Arithmetic/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Arithmetic) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Arithmetic) add_public_tablegen_target(MLIRArithmeticTransformsIncGen) add_mlir_doc(Passes ArithmeticPasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/Arithmetic/Transforms/Passes.h b/mlir/include/mlir/Dialect/Arithmetic/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Arithmetic/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Arithmetic/Transforms/Passes.h @@ -30,6 +30,15 @@ /// equivalent. std::unique_ptr createArithmeticUnsignedWhenEquivalentPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ArithmeticBufferize +#define GEN_PASS_DECL_ArithmeticExpandOps +#define GEN_PASS_DECL_ArithmeticUnsignedWhenEquivalent +#include "mlir/Dialect/Arithmetic/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/Async/CMakeLists.txt b/mlir/include/mlir/Dialect/Async/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Async/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Async/CMakeLists.txt @@ -2,6 +2,7 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Async) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Async) mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix Async) mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix Async) add_public_tablegen_target(MLIRAsyncPassIncGen) diff --git a/mlir/include/mlir/Dialect/Async/Passes.h b/mlir/include/mlir/Dialect/Async/Passes.h --- a/mlir/include/mlir/Dialect/Async/Passes.h +++ b/mlir/include/mlir/Dialect/Async/Passes.h @@ -32,6 +32,17 @@ std::unique_ptr createAsyncRuntimePolicyBasedRefCountingPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_AsyncParallelFor +#define GEN_PASS_DECL_AsyncRuntimePolicyBasedRefCounting +#define GEN_PASS_DECL_AsyncRuntimeRefCounting +#define GEN_PASS_DECL_AsyncRuntimeRefCountingOpt +#define GEN_PASS_DECL_AsyncToAsyncRuntime +#include "mlir/Dialect/Async/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/Bufferization/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Bufferization/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Bufferization/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Bufferization/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Bufferization) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Bufferization) add_public_tablegen_target(MLIRBufferizationPassIncGen) add_dependencies(mlir-headers MLIRBufferizationPassIncGen) diff --git a/mlir/include/mlir/Dialect/Bufferization/Transforms/Passes.h b/mlir/include/mlir/Dialect/Bufferization/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Bufferization/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Bufferization/Transforms/Passes.h @@ -83,6 +83,23 @@ std::unique_ptr createTensorCopyInsertionPass(const OneShotBufferizationOptions &options); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_AllocTensorElimination +#define GEN_PASS_DECL_BufferizationBufferize +#define GEN_PASS_DECL_BufferDeallocation +#define GEN_PASS_DECL_BufferHoisting +#define GEN_PASS_DECL_BufferLoopHoisting +#define GEN_PASS_DECL_BufferResultsToOutParams +#define GEN_PASS_DECL_DropEquivalentBufferResults +#define GEN_PASS_DECL_FinalizingBufferize +#define GEN_PASS_DECL_OneShotBufferize +#define GEN_PASS_DECL_PromoteBuffersToStack +#define GEN_PASS_DECL_TensorCopyInsertion +#include "mlir/Dialect/Bufferization/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/Func/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Func/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Func/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Func/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Func) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Func) add_public_tablegen_target(MLIRFuncTransformsIncGen) add_mlir_doc(Passes FuncPasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/Func/Transforms/Passes.h b/mlir/include/mlir/Dialect/Func/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Func/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Func/Transforms/Passes.h @@ -28,6 +28,13 @@ /// Creates an instance of func bufferization pass. std::unique_ptr createFuncBufferizePass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_FuncBufferize +#include "mlir/Dialect/Func/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name GPU) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name GPU) mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix GPU) mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix GPU) add_public_tablegen_target(MLIRGPUPassIncGen) diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h @@ -124,6 +124,16 @@ StringRef features, int optLevel); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_GpuAsyncRegionPass +#define GEN_PASS_DECL_GpuKernelOutlining +#define GEN_PASS_DECL_GpuLaunchSinkIndexComputations +#define GEN_PASS_DECL_GpuMapParallelLoopsPass +#include "mlir/Dialect/GPU/Transforms/Passes.h.inc" + /// Generate the code for registering passes. #define GEN_PASS_REGISTRATION #include "mlir/Dialect/GPU/Transforms/Passes.h.inc" diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/LLVMIR/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name LLVM) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name LLVM) add_public_tablegen_target(MLIRLLVMPassIncGen) add_mlir_doc(Passes LLVMPasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h @@ -28,6 +28,13 @@ /// be translated to LLVM IR. std::unique_ptr createLegalizeForExportPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_LLVMLegalizeForExport +#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" + } // namespace LLVM } // namespace mlir diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h @@ -19,6 +19,13 @@ /// Creates a pass that optimizes LLVM IR for the NVVM target. std::unique_ptr createOptimizeForTargetPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_NVVMOptimizeForTarget +#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" + } // namespace NVVM } // namespace mlir diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h @@ -15,7 +15,16 @@ class Pass; namespace LLVM { + std::unique_ptr createRequestCWrappersPass(); + +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_LLVMRequestCWrappers +#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" + } // namespace LLVM } // namespace mlir diff --git a/mlir/include/mlir/Dialect/Linalg/CMakeLists.txt b/mlir/include/mlir/Dialect/Linalg/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Linalg/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Linalg/CMakeLists.txt @@ -3,6 +3,7 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Linalg) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Linalg) mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix Linalg) mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix Linalg) add_public_tablegen_target(MLIRLinalgPassIncGen) diff --git a/mlir/include/mlir/Dialect/Linalg/Passes.h b/mlir/include/mlir/Dialect/Linalg/Passes.h --- a/mlir/include/mlir/Dialect/Linalg/Passes.h +++ b/mlir/include/mlir/Dialect/Linalg/Passes.h @@ -149,6 +149,36 @@ std::unique_ptr> createLinalgStrategyRemoveMarkersPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ConvertElementwiseToLinalg +#define GEN_PASS_DECL_LinalgBufferize +#define GEN_PASS_DECL_LinalgDetensorize +#define GEN_PASS_DECL_LinalgFoldUnitExtentDims +#define GEN_PASS_DECL_LinalgGeneralization +#define GEN_PASS_DECL_LinalgElementwiseOpFusion +#define GEN_PASS_DECL_LinalgInitTensorToAllocTensor +#define GEN_PASS_DECL_LinalgInlineScalarOperands +#define GEN_PASS_DECL_LinalgLowerToAffineLoops +#define GEN_PASS_DECL_LinalgLowerToLoops +#define GEN_PASS_DECL_LinalgLowerToParallelLoops +#define GEN_PASS_DECL_LinalgNamedOpConversion +#define GEN_PASS_DECL_LinalgStrategyDecomposePass +#define GEN_PASS_DECL_LinalgStrategyEnablePass +#define GEN_PASS_DECL_LinalgStrategyGeneralizePass +#define GEN_PASS_DECL_LinalgStrategyInterchangePass +#define GEN_PASS_DECL_LinalgStrategyLowerVectorsPass +#define GEN_PASS_DECL_LinalgStrategyPadPass +#define GEN_PASS_DECL_LinalgStrategyPeelPass +#define GEN_PASS_DECL_LinalgStrategyRemoveMarkersPass +#define GEN_PASS_DECL_LinalgStrategyTilePass +#define GEN_PASS_DECL_LinalgStrategyTileAndFusePass +#define GEN_PASS_DECL_LinalgStrategyVectorizePass +#define GEN_PASS_DECL_LinalgTiling +#include "mlir/Dialect/Linalg/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/MemRef/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/MemRef/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/MemRef/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/MemRef/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name MemRef) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name MemRef) add_public_tablegen_target(MLIRMemRefPassIncGen) add_dependencies(mlir-headers MLIRMemRefPassIncGen) diff --git a/mlir/include/mlir/Dialect/MemRef/Transforms/Passes.h b/mlir/include/mlir/Dialect/MemRef/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/MemRef/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/MemRef/Transforms/Passes.h @@ -111,6 +111,17 @@ /// in terms of shapes of its input operands. std::unique_ptr createResolveShapedTypeResultDimsPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_ExpandOps +#define GEN_PASS_DECL_FoldSubViewOps +#define GEN_PASS_DECL_NormalizeMemRefs +#define GEN_PASS_DECL_ResolveRankedShapeTypeResultDims +#define GEN_PASS_DECL_ResolveShapedTypeResultDims +#include "mlir/Dialect/MemRef/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/NVGPU/CMakeLists.txt b/mlir/include/mlir/Dialect/NVGPU/CMakeLists.txt --- a/mlir/include/mlir/Dialect/NVGPU/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/NVGPU/CMakeLists.txt @@ -2,6 +2,7 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name NVGPU) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name NVGPU) mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix NVGPU) mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix NVGPU) add_public_tablegen_target(MLIRNVGPUPassIncGen) diff --git a/mlir/include/mlir/Dialect/NVGPU/Passes.h b/mlir/include/mlir/Dialect/NVGPU/Passes.h --- a/mlir/include/mlir/Dialect/NVGPU/Passes.h +++ b/mlir/include/mlir/Dialect/NVGPU/Passes.h @@ -22,6 +22,13 @@ } // namespace nvgpu +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_OptimizeSharedMemory +#include "mlir/Dialect/NVGPU/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/SCF/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/SCF/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/SCF/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/SCF/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name SCF) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name SCF) add_public_tablegen_target(MLIRSCFPassIncGen) add_dependencies(mlir-headers MLIRSCFPassIncGen) diff --git a/mlir/include/mlir/Dialect/SCF/Transforms/Passes.h b/mlir/include/mlir/Dialect/SCF/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/SCF/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/SCF/Transforms/Passes.h @@ -59,6 +59,22 @@ // Creates a pass which lowers for loops into while loops. std::unique_ptr createForToWhileLoopPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_SCFBufferize +#define GEN_PASS_DECL_SCFForLoopCanonicalization +#define GEN_PASS_DECL_SCFForLoopRangeFolding +#define GEN_PASS_DECL_SCFForLoopPeeling +#define GEN_PASS_DECL_SCFForLoopSpecialization +#define GEN_PASS_DECL_SCFForToWhileLoop +#define GEN_PASS_DECL_SCFParallelLoopCollapsing +#define GEN_PASS_DECL_SCFParallelLoopFusion +#define GEN_PASS_DECL_SCFParallelLoopSpecialization +#define GEN_PASS_DECL_SCFParallelLoopTiling +#include "mlir/Dialect/SCF/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/SPIRV/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/SPIRV/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/SPIRV/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/SPIRV/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name SPIRV) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name SPIRV) add_public_tablegen_target(MLIRSPIRVPassIncGen) add_dependencies(mlir-headers MLIRSPIRVPassIncGen) diff --git a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h @@ -68,6 +68,18 @@ std::unique_ptr> createUnifyAliasedResourcePass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_SPIRVCanonicalizeGL +#define GEN_PASS_DECL_SPIRVCompositeTypeLayout +#define GEN_PASS_DECL_SPIRVLowerABIAttributes +#define GEN_PASS_DECL_SPIRVRewriteInsertsPass +#define GEN_PASS_DECL_SPIRVUnifyAliasedResourcePass +#define GEN_PASS_DECL_SPIRVUpdateVCE +#include "mlir/Dialect/SPIRV/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td --- a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td @@ -12,7 +12,7 @@ include "mlir/Pass/PassBase.td" def SPIRVCompositeTypeLayout - : Pass<"decorate-spirv-composite-type-layout", "ModuleOp"> { + : Pass<"decorate-spirv-composite-type-layout", "mlir::ModuleOp"> { let summary = "Decorate SPIR-V composite type with layout info"; let constructor = "mlir::spirv::createDecorateSPIRVCompositeTypeLayoutPass()"; } diff --git a/mlir/include/mlir/Dialect/Shape/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Shape/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Shape/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Shape/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Shape) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Shape) add_public_tablegen_target(MLIRShapeTransformsIncGen) add_mlir_doc(Passes ShapePasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/Shape/Transforms/Passes.h b/mlir/include/mlir/Dialect/Shape/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Shape/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Shape/Transforms/Passes.h @@ -50,6 +50,15 @@ // level. std::unique_ptr> createShapeBufferizePass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_RemoveShapeConstraints +#define GEN_PASS_DECL_ShapeBufferize +#define GEN_PASS_DECL_ShapeToShapeLowering +#include "mlir/Dialect/Shape/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/SparseTensor/Pipelines/Passes.h b/mlir/include/mlir/Dialect/SparseTensor/Pipelines/Passes.h --- a/mlir/include/mlir/Dialect/SparseTensor/Pipelines/Passes.h +++ b/mlir/include/mlir/Dialect/SparseTensor/Pipelines/Passes.h @@ -95,18 +95,6 @@ desc("Enables the use of X86Vector dialect while lowering the vector " "dialect."), init(false)}; - - /// Projects out the options for `createConvertVectorToLLVMPass`. - LowerVectorToLLVMOptions lowerVectorToLLVMOptions() const { - LowerVectorToLLVMOptions opts{}; - opts.enableReassociateFPReductions(reassociateFPReductions); - opts.enableIndexOptimizations(indexOptimizations); - opts.enableArmNeon(armNeon); - opts.enableArmSVE(armSVE); - opts.enableAMX(amx); - opts.enableX86Vector(x86Vector); - return opts; - } }; //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/SparseTensor/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/SparseTensor/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/SparseTensor/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/SparseTensor/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name SparseTensor) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name SparseTensor) mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix SparseTensor) mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix SparseTensor) add_public_tablegen_target(MLIRSparseTensorPassIncGen) diff --git a/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h b/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h @@ -147,6 +147,14 @@ std::unique_ptr createDenseBufferizationPass( const bufferization::OneShotBufferizationOptions &options); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_SparseTensorConversion +#define GEN_PASS_DECL_Sparsification +#include "mlir/Dialect/SparseTensor/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration. //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/Tensor/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Tensor/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Tensor/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Tensor/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Tensor) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Tensor) add_public_tablegen_target(MLIRTensorTransformsIncGen) add_mlir_doc(Passes TensorPasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h b/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h @@ -15,6 +15,13 @@ /// Creates an instance of `tensor` dialect bufferization pass. std::unique_ptr createTensorBufferizePass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_TensorBufferize +#include "mlir/Dialect/Tensor/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/Tosa/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Tosa/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Tosa/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Tosa/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name TosaOpt) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name TosaOpt) add_public_tablegen_target(MLIRTosaPassIncGen) add_dependencies(mlir-headers MLIRTosaPassIncGen) diff --git a/mlir/include/mlir/Dialect/Tosa/Transforms/PassDetail.h b/mlir/include/mlir/Dialect/Tosa/Transforms/PassDetail.h --- a/mlir/include/mlir/Dialect/Tosa/Transforms/PassDetail.h +++ b/mlir/include/mlir/Dialect/Tosa/Transforms/PassDetail.h @@ -16,7 +16,7 @@ namespace mlir { -#define GEN_PASS_CLASSES +#define GEN_PASS_DECL_TosaLayerwiseConstantFoldPass #include "mlir/Dialect/Tosa/Transforms/Passes.h.inc" } // namespace mlir diff --git a/mlir/include/mlir/Dialect/Tosa/Transforms/Passes.h b/mlir/include/mlir/Dialect/Tosa/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Tosa/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Tosa/Transforms/Passes.h @@ -35,6 +35,19 @@ std::unique_ptr createTosaTestQuantUtilAPIPass(); std::unique_ptr createTosaOptionalDecompositions(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_TosaInferShapes +#define GEN_PASS_DECL_TosaMakeBroadcastable +#define GEN_PASS_DECL_TosaOptionalDecompositions +#include "mlir/Dialect/Tosa/Transforms/Passes.h.inc" + +//===----------------------------------------------------------------------===// +// Registration. +//===----------------------------------------------------------------------===// + #define GEN_PASS_REGISTRATION #include "mlir/Dialect/Tosa/Transforms/Passes.h.inc" diff --git a/mlir/include/mlir/Dialect/Transform/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Transform/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Transform/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Transform/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Transform) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Transform) add_public_tablegen_target(MLIRTransformDialectTransformsIncGen) add_mlir_doc(Passes TransformPasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/Transform/Transforms/Passes.h b/mlir/include/mlir/Dialect/Transform/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Transform/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Transform/Transforms/Passes.h @@ -18,6 +18,13 @@ namespace transform { std::unique_ptr createCheckUsesPass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_CheckUses +#include "mlir/Dialect/Transform/Transforms/Passes.h.inc" + #define GEN_PASS_REGISTRATION #include "mlir/Dialect/Transform/Transforms/Passes.h.inc" } // namespace transform diff --git a/mlir/include/mlir/Dialect/Vector/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/Vector/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Dialect/Vector/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/Vector/Transforms/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Vector) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Vector) add_public_tablegen_target(MLIRVectorTransformsIncGen) add_mlir_doc(Passes VectorPasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Dialect/Vector/Transforms/Passes.h b/mlir/include/mlir/Dialect/Vector/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/Vector/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Vector/Transforms/Passes.h @@ -16,6 +16,13 @@ /// Creates an instance of the `vector` dialect bufferization pass. std::unique_ptr createVectorBufferizePass(); +//===----------------------------------------------------------------------===// +// Declarations. +//===----------------------------------------------------------------------===// + +#define GEN_PASS_DECL_VectorBufferize +#include "mlir/Dialect/Vector/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Pass/PassBase.td b/mlir/include/mlir/Pass/PassBase.td --- a/mlir/include/mlir/Pass/PassBase.td +++ b/mlir/include/mlir/Pass/PassBase.td @@ -76,7 +76,7 @@ string description = ""; // A C++ constructor call to create an instance of this pass. - code constructor = [{}]; + code constructor = ?; // A list of dialects this pass may produce entities in. list dependentDialects = []; diff --git a/mlir/include/mlir/Reducer/CMakeLists.txt b/mlir/include/mlir/Reducer/CMakeLists.txt --- a/mlir/include/mlir/Reducer/CMakeLists.txt +++ b/mlir/include/mlir/Reducer/CMakeLists.txt @@ -1,5 +1,6 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Reducer) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Reducer) add_public_tablegen_target(MLIRReducerIncGen) add_mlir_doc(Passes ReducerPasses ./ -gen-pass-doc) diff --git a/mlir/include/mlir/Reducer/PassDetail.h b/mlir/include/mlir/Reducer/PassDetail.h --- a/mlir/include/mlir/Reducer/PassDetail.h +++ b/mlir/include/mlir/Reducer/PassDetail.h @@ -14,9 +14,6 @@ namespace mlir { -#define GEN_PASS_CLASSES -#include "mlir/Reducer/Passes.h.inc" - } // namespace mlir #endif // MLIR_REDUCER_PASSDETAIL_H diff --git a/mlir/include/mlir/Reducer/Passes.h b/mlir/include/mlir/Reducer/Passes.h --- a/mlir/include/mlir/Reducer/Passes.h +++ b/mlir/include/mlir/Reducer/Passes.h @@ -16,6 +16,10 @@ std::unique_ptr createOptReductionPass(); +#define GEN_PASS_DECL_OptReduction +#define GEN_PASS_DECL_ReductionTree +#include "mlir/Reducer/Passes.h.inc" + /// Generate the code for registering reducer passes. #define GEN_PASS_REGISTRATION #include "mlir/Reducer/Passes.h.inc" diff --git a/mlir/include/mlir/TableGen/Pass.h b/mlir/include/mlir/TableGen/Pass.h --- a/mlir/include/mlir/TableGen/Pass.h +++ b/mlir/include/mlir/TableGen/Pass.h @@ -92,7 +92,7 @@ StringRef getDescription() const; /// Return the C++ constructor call to create an instance of this pass. - StringRef getConstructor() const; + Optional getConstructor() const; /// Return the dialects this pass needs to be registered. ArrayRef getDependentDialects() const; diff --git a/mlir/include/mlir/Transforms/CMakeLists.txt b/mlir/include/mlir/Transforms/CMakeLists.txt --- a/mlir/include/mlir/Transforms/CMakeLists.txt +++ b/mlir/include/mlir/Transforms/CMakeLists.txt @@ -1,6 +1,7 @@ set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Transforms) +mlir_tablegen(Passes.cpp.inc -gen-pass-defs -name Transforms) mlir_tablegen(Transforms.capi.h.inc -gen-pass-capi-header --prefix Transforms) mlir_tablegen(Transforms.capi.cpp.inc -gen-pass-capi-impl --prefix Transforms) add_public_tablegen_target(MLIRTransformsPassIncGen) diff --git a/mlir/include/mlir/Transforms/LocationSnapshot.h b/mlir/include/mlir/Transforms/LocationSnapshot.h --- a/mlir/include/mlir/Transforms/LocationSnapshot.h +++ b/mlir/include/mlir/Transforms/LocationSnapshot.h @@ -61,6 +61,9 @@ /// Overload utilizing pass options for initialization. std::unique_ptr createLocationSnapshotPass(); +#define GEN_PASS_DECL_LocationSnapshot +#include "mlir/Transforms/Passes.h.inc" + } // namespace mlir #endif // MLIR_TRANSFORMS_LOCATIONSNAPSHOT_H diff --git a/mlir/include/mlir/Transforms/Passes.h b/mlir/include/mlir/Transforms/Passes.h --- a/mlir/include/mlir/Transforms/Passes.h +++ b/mlir/include/mlir/Transforms/Passes.h @@ -15,6 +15,7 @@ #define MLIR_TRANSFORMS_PASSES_H #include "mlir/Pass/Pass.h" +#include "mlir/Pass/PassManager.h" #include "mlir/Transforms/LocationSnapshot.h" #include "mlir/Transforms/ViewOpGraph.h" #include "llvm/Support/Debug.h" @@ -101,6 +102,19 @@ /// their producers. std::unique_ptr createTopologicalSortPass(); +#define GEN_PASS_DECL_Canonicalizer +#define GEN_PASS_DECL_ControlFlowSink +#define GEN_PASS_DECL_CSE +#define GEN_PASS_DECL_SymbolDCE +#define GEN_PASS_DECL_Inliner +#define GEN_PASS_DECL_LoopInvariantCodeMotion +#define GEN_PASS_DECL_PrintOpStats +#define GEN_PASS_DECL_SCCP +#define GEN_PASS_DECL_StripDebugInfo +#define GEN_PASS_DECL_SymbolPrivatize +#define GEN_PASS_DECL_TopologicalSort +#include "mlir/Transforms/Passes.h.inc" + //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Transforms/ViewOpGraph.h b/mlir/include/mlir/Transforms/ViewOpGraph.h --- a/mlir/include/mlir/Transforms/ViewOpGraph.h +++ b/mlir/include/mlir/Transforms/ViewOpGraph.h @@ -22,6 +22,9 @@ /// Creates a pass to print op graphs. std::unique_ptr createPrintOpGraphPass(raw_ostream &os = llvm::errs()); +#define GEN_PASS_DECL_ViewOpGraph +#include "mlir/Transforms/Passes.h.inc" + } // namespace mlir #endif // MLIR_TRANSFORMS_VIEWOPGRAPH_H_ diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp --- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp +++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp @@ -7,15 +7,20 @@ //===----------------------------------------------------------------------===// #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 "../PassDetail.h" + using namespace mlir; using namespace mlir::amdgpu; +#define GEN_PASS_DEF_ConvertAMDGPUToROCDL +#include "mlir/Conversion/Passes.cpp.inc" + static Value createI32Constant(ConversionPatternRewriter &rewriter, Location loc, int32_t value) { Type llvmI32 = rewriter.getI32Type(); @@ -270,7 +275,7 @@ struct ConvertAMDGPUToROCDLPass : public ConvertAMDGPUToROCDLBase { - ConvertAMDGPUToROCDLPass() = default; + using ConvertAMDGPUToROCDLBase::ConvertAMDGPUToROCDLBase; void runOnOperation() override { MLIRContext *ctx = &getContext(); @@ -307,3 +312,7 @@ std::unique_ptr mlir::createConvertAMDGPUToROCDLPass() { return std::make_unique(); } + +std::unique_ptr mlir::createConvertAMDGPUToROCDLPass(const ConvertAMDGPUToROCDLPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp b/mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp --- 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" @@ -26,9 +25,14 @@ #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/Passes.h" +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::vector; +#define GEN_PASS_DEF_ConvertAffineToStandard +#include "mlir/Conversion/Passes.cpp.inc" + /// Given a range of values, emit the code that reduces them with "min" or "max" /// depending on the provided comparison predicate. The predicate defines which /// comparison to perform, "lt" for "min", "gt" for "max" and is used for the @@ -545,6 +549,8 @@ namespace { class LowerAffinePass : public ConvertAffineToStandardBase { + using ConvertAffineToStandardBase::ConvertAffineToStandardBase; + void runOnOperation() override { RewritePatternSet patterns(&getContext()); populateAffineToStdConversionPatterns(patterns); diff --git a/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp b/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp --- a/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp +++ b/mlir/lib/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.cpp @@ -7,15 +7,20 @@ //===----------------------------------------------------------------------===// #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 "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertArithmeticToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + namespace { //===----------------------------------------------------------------------===// @@ -259,7 +264,7 @@ namespace { struct ConvertArithmeticToLLVMPass : public ConvertArithmeticToLLVMBase { - ConvertArithmeticToLLVMPass() = default; + using ConvertArithmeticToLLVMBase::ConvertArithmeticToLLVMBase; void runOnOperation() override { LLVMConversionTarget target(getContext()); @@ -326,6 +331,10 @@ // clang-format on } -std::unique_ptr mlir::arith::createConvertArithmeticToLLVMPass() { +std::unique_ptr mlir::createConvertArithmeticToLLVMPass() { return std::make_unique(); } + +std::unique_ptr mlir::createConvertArithmeticToLLVMPass(const ConvertArithmeticToLLVMPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp b/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp --- a/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp +++ b/mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp @@ -7,8 +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" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" @@ -17,10 +16,16 @@ #include "mlir/IR/BuiltinTypes.h" #include "llvm/Support/Debug.h" +#include "../PassDetail.h" +#include "../SPIRVCommon/Pattern.h" + #define DEBUG_TYPE "arith-to-spirv-pattern" using namespace mlir; +#define GEN_PASS_DEF_ConvertArithmeticToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // Operation Conversion //===----------------------------------------------------------------------===// @@ -911,6 +916,8 @@ namespace { struct ConvertArithmeticToSPIRVPass : public ConvertArithmeticToSPIRVBase { + using ConvertArithmeticToSPIRVBase::ConvertArithmeticToSPIRVBase; + void runOnOperation() override { Operation *op = getOperation(); auto targetAttr = spirv::lookupTargetEnvOrDefault(op); @@ -940,7 +947,10 @@ }; } // namespace -std::unique_ptr> -mlir::arith::createConvertArithmeticToSPIRVPass() { +std::unique_ptr> mlir::createConvertArithmeticToSPIRVPass() { return std::make_unique(); } + +std::unique_ptr> mlir::createConvertArithmeticToSPIRVPass(const ConvertArithmeticToSPIRVPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp b/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp --- 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,9 +15,14 @@ #include "mlir/Pass/PassRegistry.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::arm_neon; +#define GEN_PASS_DEF_ConvertArmNeon2dToIntr +#include "mlir/Conversion/Passes.cpp.inc" + namespace { class Sdot2dLoweringPattern : public OpRewritePattern { @@ -48,6 +53,8 @@ class ConvertArmNeon2dToIntr : public ConvertArmNeon2dToIntrBase { + using ConvertArmNeon2dToIntrBase::ConvertArmNeon2dToIntrBase; + void runOnOperation() override { auto *context = &getContext(); diff --git a/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp b/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp --- 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,11 +23,16 @@ #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/TypeSwitch.h" +#include "../PassDetail.h" + #define DEBUG_TYPE "convert-async-to-llvm" using namespace mlir; using namespace mlir::async; +#define GEN_PASS_DEF_ConvertAsyncToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // Async Runtime C API declaration. //===----------------------------------------------------------------------===// @@ -976,6 +980,8 @@ namespace { struct ConvertAsyncToLLVMPass : public ConvertAsyncToLLVMBase { + using ConvertAsyncToLLVMBase::ConvertAsyncToLLVMBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp b/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp --- a/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp +++ b/mlir/lib/Conversion/BufferizationToMemRef/BufferizationToMemRef.cpp @@ -12,7 +12,7 @@ //===----------------------------------------------------------------------===// #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" @@ -20,8 +20,13 @@ #include "mlir/Support/LogicalResult.h" #include "mlir/Transforms/DialectConversion.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertBufferizationToMemRef +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// The CloneOpConversion transforms all bufferization clone operations into /// memref alloc and memref copy operations. In the dynamic-shape case, it also @@ -83,7 +88,7 @@ namespace { struct BufferizationToMemRefPass : public ConvertBufferizationToMemRefBase { - BufferizationToMemRefPass() = default; + using ConvertBufferizationToMemRefBase::ConvertBufferizationToMemRefBase; void runOnOperation() override { RewritePatternSet patterns(&getContext()); @@ -101,6 +106,6 @@ }; } // namespace -std::unique_ptr mlir::createBufferizationToMemRefPass() { +std::unique_ptr mlir::createConvertBufferizationToMemRefPass() { return std::make_unique(); } diff --git a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp --- a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp +++ b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp @@ -8,16 +8,20 @@ #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 "../PassDetail.h" + using namespace mlir; using namespace mlir::LLVM; +#define GEN_PASS_DEF_ConvertComplexToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // ComplexStructBuilder implementation. //===----------------------------------------------------------------------===// @@ -319,6 +323,8 @@ namespace { struct ConvertComplexToLLVMPass : public ConvertComplexToLLVMBase { + using ConvertComplexToLLVMBase::ConvertComplexToLLVMBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp b/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp --- a/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp +++ b/mlir/lib/Conversion/ComplexToLibm/ComplexToLibm.cpp @@ -8,13 +8,17 @@ #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 "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertComplexToLibm +#include "mlir/Conversion/Passes.cpp.inc" + namespace { // Functor to resolve the function name corresponding to the given complex // result type. @@ -118,6 +122,8 @@ namespace { struct ConvertComplexToLibmPass : public ConvertComplexToLibmBase { + using ConvertComplexToLibmBase::ConvertComplexToLibmBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp b/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp --- a/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp +++ b/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp @@ -8,19 +8,22 @@ #include "mlir/Conversion/ComplexToStandard/ComplexToStandard.h" -#include -#include - -#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/Transforms/DialectConversion.h" +#include +#include + +#include "../PassDetail.h" using namespace mlir; +#define GEN_PASS_DEF_ConvertComplexToStandard +#include "mlir/Conversion/Passes.cpp.inc" + namespace { struct AbsOpConversion : public OpConversionPattern { using OpConversionPattern::OpConversionPattern; @@ -1065,6 +1068,8 @@ namespace { struct ConvertComplexToStandardPass : public ConvertComplexToStandardBase { + using ConvertComplexToStandardBase::ConvertComplexToStandardBase; + void runOnOperation() override; }; diff --git a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp --- 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" @@ -25,8 +25,13 @@ #include "llvm/ADT/StringRef.h" #include +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertControlFlowToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + #define PASS_NAME "convert-cf-to-llvm" namespace { @@ -197,7 +202,7 @@ /// A pass converting MLIR operations into the LLVM IR dialect. struct ConvertControlFlowToLLVM : public ConvertControlFlowToLLVMBase { - ConvertControlFlowToLLVM() = default; + using ConvertControlFlowToLLVMBase::ConvertControlFlowToLLVMBase; /// Run the dialect converter on the module. void runOnOperation() override { @@ -218,6 +223,11 @@ }; } // namespace -std::unique_ptr mlir::cf::createConvertControlFlowToLLVMPass() { +std::unique_ptr mlir::createConvertControlFlowToLLVMPass() { return std::make_unique(); } + +std::unique_ptr mlir::createConvertControlFlowToLLVMPass( + const ConvertControlFlowToLLVMPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp b/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp --- a/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp +++ b/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRVPass.cpp @@ -11,17 +11,24 @@ //===----------------------------------------------------------------------===// #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" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertControlFlowToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR ControlFlow operations into the SPIR-V dialect. -class ConvertControlFlowToSPIRVPass +struct ConvertControlFlowToSPIRVPass : public ConvertControlFlowToSPIRVBase { + using ConvertControlFlowToSPIRVBase::ConvertControlFlowToSPIRVBase; + void runOnOperation() override; }; } // namespace @@ -48,3 +55,8 @@ std::unique_ptr> mlir::createConvertControlFlowToSPIRVPass() { return std::make_unique(); } + +std::unique_ptr> mlir::createConvertControlFlowToSPIRVPass( + const ConvertControlFlowToSPIRVPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp --- 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/Conversion/FuncToLLVM/ConvertFuncToLLVM.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" @@ -43,8 +43,13 @@ #include #include +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertFuncToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + #define PASS_NAME "convert-func-to-llvm" /// Only retain those attributes that are not constructed by @@ -664,14 +669,7 @@ /// A pass converting Func operations into the LLVM IR dialect. struct ConvertFuncToLLVMPass : public ConvertFuncToLLVMBase { - ConvertFuncToLLVMPass() = default; - ConvertFuncToLLVMPass(bool useBarePtrCallConv, unsigned indexBitwidth, - bool useAlignedAlloc, - const llvm::DataLayout &dataLayout) { - this->useBarePtrCallConv = useBarePtrCallConv; - this->indexBitwidth = indexBitwidth; - this->dataLayout = dataLayout.getStringRepresentation(); - } + using ConvertFuncToLLVMBase::ConvertFuncToLLVMBase; /// Run the dialect converter on the module. void runOnOperation() override { @@ -718,15 +716,6 @@ } std::unique_ptr> -mlir::createConvertFuncToLLVMPass(const LowerToLLVMOptions &options) { - auto allocLowering = options.allocLowering; - // There is no way to provide additional patterns for pass, so - // AllocLowering::None will always fail. - assert(allocLowering != LowerToLLVMOptions::AllocLowering::None && - "ConvertFuncToLLVMPass doesn't support AllocLowering::None"); - bool useAlignedAlloc = - (allocLowering == LowerToLLVMOptions::AllocLowering::AlignedAlloc); - return std::make_unique( - options.useBarePtrCallConv, options.getIndexBitwidth(), useAlignedAlloc, - options.dataLayout); +mlir::createConvertFuncToLLVMPass(const ConvertFuncToLLVMPassOptions &options) { + return std::make_unique(options); } diff --git a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp --- a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp +++ b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRVPass.cpp @@ -11,17 +11,24 @@ //===----------------------------------------------------------------------===// #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" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertFuncToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR Func operations into the SPIR-V dialect. -class ConvertFuncToSPIRVPass +struct ConvertFuncToSPIRVPass : public ConvertFuncToSPIRVBase { + using ConvertFuncToSPIRVBase::ConvertFuncToSPIRVBase; + void runOnOperation() override; }; } // namespace @@ -49,3 +56,8 @@ std::unique_ptr> mlir::createConvertFuncToSPIRVPass() { return std::make_unique(); } + +std::unique_ptr> mlir::createConvertFuncToSPIRVPass( + const ConvertFuncToSPIRVPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp b/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp --- 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" @@ -33,13 +32,17 @@ #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/BuiltinTypes.h" - #include "llvm/ADT/STLExtras.h" #include "llvm/Support/Error.h" #include "llvm/Support/FormatVariadic.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_GpuToLLVMConversionPass +#include "mlir/Conversion/Passes.cpp.inc" + static constexpr const char *kGpuBinaryStorageSuffix = "_gpubin_cst"; namespace { @@ -47,13 +50,7 @@ class GpuToLLVMConversionPass : public GpuToLLVMConversionPassBase { public: - GpuToLLVMConversionPass() = default; - - GpuToLLVMConversionPass(bool kernelBarePtrCallConv) - : GpuToLLVMConversionPass() { - if (this->kernelBarePtrCallConv.getNumOccurrences() == 0) - this->kernelBarePtrCallConv = kernelBarePtrCallConv; - } + using GpuToLLVMConversionPassBase::GpuToLLVMConversionPassBase; GpuToLLVMConversionPass(const GpuToLLVMConversionPass &other) : GpuToLLVMConversionPassBase(other) {} @@ -896,9 +893,12 @@ return success(); } -std::unique_ptr> -mlir::createGpuToLLVMConversionPass(bool kernelBarePtrCallConv) { - return std::make_unique(kernelBarePtrCallConv); +std::unique_ptr> mlir::createGpuToLLVMConversionPass() { + return std::make_unique(); +} + +std::unique_ptr> mlir::createGpuToLLVMConversionPass(const GpuToLLVMConversionPassPassOptions &options) { + return std::make_unique(options); } void mlir::populateGpuToLLVMConversionPatterns(LLVMTypeConverter &converter, diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp --- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp +++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp @@ -33,13 +33,16 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/Support/FormatVariadic.h" +#include "../PassDetail.h" #include "../GPUCommon/GPUOpsLowering.h" #include "../GPUCommon/IndexIntrinsicsOpLowering.h" #include "../GPUCommon/OpToFuncCallLowering.h" -#include "../PassDetail.h" using namespace mlir; +#define GEN_PASS_DEF_ConvertGpuOpsToNVVMOps +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// Convert gpu dialect shfl mode enum to the equivalent nvvm one. @@ -153,10 +156,7 @@ /// code. struct LowerGpuOpsToNVVMOpsPass : public ConvertGpuOpsToNVVMOpsBase { - LowerGpuOpsToNVVMOpsPass() = default; - LowerGpuOpsToNVVMOpsPass(unsigned indexBitwidth) { - this->indexBitwidth = indexBitwidth; - } + using ConvertGpuOpsToNVVMOpsBase::ConvertGpuOpsToNVVMOpsBase; void runOnOperation() override { gpu::GPUModuleOp m = getOperation(); @@ -287,7 +287,11 @@ "__nv_tanh"); } +std::unique_ptr> mlir::createLowerGpuOpsToNVVMOpsPass() { + return std::make_unique(); +} + std::unique_ptr> -mlir::createLowerGpuOpsToNVVMOpsPass(unsigned indexBitwidth) { - return std::make_unique(indexBitwidth); +mlir::createLowerGpuOpsToNVVMOpsPass(const ConvertGpuOpsToNVVMOpsPassOptions &options) { + return std::make_unique(options); } diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp --- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp +++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp @@ -11,9 +11,9 @@ // //===----------------------------------------------------------------------===// -#include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h" #include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h" +#include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h" #include "mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h" #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h" @@ -34,13 +34,16 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/Support/FormatVariadic.h" +#include "../PassDetail.h" #include "../GPUCommon/GPUOpsLowering.h" #include "../GPUCommon/IndexIntrinsicsOpLowering.h" #include "../GPUCommon/OpToFuncCallLowering.h" -#include "../PassDetail.h" using namespace mlir; +#define GEN_PASS_DEF_ConvertGpuOpsToROCDLOps +#include "mlir/Conversion/Passes.cpp.inc" + /// Returns true if the given `gpu.func` can be safely called using the bare /// pointer calling convention. static bool canBeCalledWithBarePointers(gpu::GPUFuncOp func) { @@ -63,19 +66,7 @@ // code. struct LowerGpuOpsToROCDLOpsPass : public ConvertGpuOpsToROCDLOpsBase { - LowerGpuOpsToROCDLOpsPass() = default; - LowerGpuOpsToROCDLOpsPass(const std::string &chipset, unsigned indexBitwidth, - bool useBarePtrCallConv, - gpu::amd::Runtime runtime) { - if (this->chipset.getNumOccurrences() == 0) - this->chipset = chipset; - if (this->indexBitwidth.getNumOccurrences() == 0) - this->indexBitwidth = indexBitwidth; - if (this->useBarePtrCallConv.getNumOccurrences() == 0) - this->useBarePtrCallConv = useBarePtrCallConv; - if (this->runtime.getNumOccurrences() == 0) - this->runtime = runtime; - } + using ConvertGpuOpsToROCDLOpsBase::ConvertGpuOpsToROCDLOpsBase; void runOnOperation() override { gpu::GPUModuleOp m = getOperation(); @@ -219,11 +210,10 @@ "__ocml_tanh_f64"); } -std::unique_ptr> -mlir::createLowerGpuOpsToROCDLOpsPass(const std::string &chipset, - unsigned indexBitwidth, - bool useBarePtrCallConv, - gpu::amd::Runtime runtime) { - return std::make_unique( - chipset, indexBitwidth, useBarePtrCallConv, runtime); +std::unique_ptr> mlir::createLowerGpuOpsToROCDLOpsPass() { + return std::make_unique(); +} + +std::unique_ptr> mlir::createLowerGpuOpsToROCDLOpsPass(const ConvertGpuOpsToROCDLOpsPassOptions &options) { + return std::make_unique(options); } diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp --- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp +++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp @@ -12,19 +12,23 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h" +#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h" -#include "../PassDetail.h" #include "mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h" #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h" -#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h" #include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertGPUToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// Pass to lower GPU Dialect to SPIR-V. The pass only converts the gpu.func ops /// inside gpu.module ops. i.e., the function that are referenced in @@ -37,12 +41,9 @@ /// 2) Lower the body of the spirv::ModuleOp. class GPUToSPIRVPass : public ConvertGPUToSPIRVBase { public: - explicit GPUToSPIRVPass(bool mapMemorySpace) - : mapMemorySpace(mapMemorySpace) {} - void runOnOperation() override; + using ConvertGPUToSPIRVBase::ConvertGPUToSPIRVBase; -private: - bool mapMemorySpace; + void runOnOperation() override; }; } // namespace @@ -92,7 +93,11 @@ return signalPassFailure(); } +std::unique_ptr> mlir::createConvertGPUToSPIRVPass() { + return std::make_unique(); +} + std::unique_ptr> -mlir::createConvertGPUToSPIRVPass(bool mapMemorySpace) { - return std::make_unique(mapMemorySpace); +mlir::createConvertGPUToSPIRVPass(const ConvertGPUToSPIRVPassOptions &options) { + return std::make_unique(options); } diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp --- 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" @@ -25,8 +25,13 @@ #include "mlir/IR/BuiltinTypes.h" #include "mlir/Target/SPIRV/Serialization.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertGpuLaunchFuncToVulkanLaunchFunc +#include "mlir/Conversion/Passes.cpp.inc" + static constexpr const char *kSPIRVBlobAttrName = "spirv_blob"; static constexpr const char *kSPIRVEntryPointAttrName = "spirv_entry_point"; static constexpr const char *kVulkanLaunch = "vulkanLaunch"; @@ -41,6 +46,9 @@ : public ConvertGpuLaunchFuncToVulkanLaunchFuncBase< ConvertGpuLaunchFuncToVulkanLaunchFunc> { public: + using ConvertGpuLaunchFuncToVulkanLaunchFuncBase:: + ConvertGpuLaunchFuncToVulkanLaunchFuncBase; + void runOnOperation() override; private: diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp --- a/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp +++ b/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp @@ -14,18 +14,23 @@ // //===----------------------------------------------------------------------===// -#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 "llvm/ADT/SmallString.h" #include "llvm/Support/FormatVariadic.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertVulkanLaunchFuncToVulkanCalls +#define GEN_PASS_DEF_ConvertGpuLaunchFuncToVulkanLaunchFunc +#include "mlir/Conversion/Passes.cpp.inc" + static constexpr const char *kCInterfaceVulkanLaunch = "_mlir_ciface_vulkanLaunch"; static constexpr const char *kDeinitVulkan = "deinitVulkan"; @@ -55,6 +60,9 @@ class VulkanLaunchFuncToVulkanCallsPass : public ConvertVulkanLaunchFuncToVulkanCallsBase< VulkanLaunchFuncToVulkanCallsPass> { +public: + using ConvertVulkanLaunchFuncToVulkanCallsBase::ConvertVulkanLaunchFuncToVulkanCallsBase; + private: void initializeCachedTypes() { llvmFloatType = Float32Type::get(&getContext()); diff --git a/mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp b/mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp --- 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,10 +40,15 @@ #include "llvm/Support/Allocator.h" #include "llvm/Support/ErrorHandling.h" +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::LLVM; using namespace mlir::linalg; +#define GEN_PASS_DEF_ConvertLinalgToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + template static Type getPtrToElementType(T containerType, LLVMTypeConverter &lowering) { return LLVMPointerType::get( @@ -75,6 +79,8 @@ namespace { struct ConvertLinalgToLLVMPass : public ConvertLinalgToLLVMBase { + using ConvertLinalgToLLVMBase::ConvertLinalgToLLVMBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.cpp b/mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRVPass.cpp --- 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 "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertLinalgToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR Linalg ops into SPIR-V ops. -class LinalgToSPIRVPass : public ConvertLinalgToSPIRVBase { +struct LinalgToSPIRVPass : public ConvertLinalgToSPIRVBase { + using ConvertLinalgToSPIRVBase::ConvertLinalgToSPIRVBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp b/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp --- 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" @@ -17,9 +16,14 @@ #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_ConvertLinalgToStandard +#include "mlir/Conversion/Passes.cpp.inc" + /// Helper function to extract the operand types that are passed to the /// generated CallOp. MemRefTypes have their layout canonicalized since the /// information is not used in signature generation. @@ -122,6 +126,8 @@ namespace { struct ConvertLinalgToStandardPass : public ConvertLinalgToStandardBase { + using ConvertLinalgToStandardBase::ConvertLinalgToStandardBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp --- a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp +++ b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #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" @@ -15,8 +15,13 @@ #include "mlir/Dialect/Math/IR/Math.h" #include "mlir/IR/TypeUtilities.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertMathToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + namespace { using AbsFOpLowering = VectorConvertToLLVMPattern; using CeilOpLowering = VectorConvertToLLVMPattern; @@ -249,7 +254,7 @@ struct ConvertMathToLLVMPass : public ConvertMathToLLVMBase { - ConvertMathToLLVMPass() = default; + using ConvertMathToLLVMBase::ConvertMathToLLVMBase; void runOnOperation() override { RewritePatternSet patterns(&getContext()); diff --git a/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp b/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp --- 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/Math/IR/Math.h" @@ -18,8 +17,13 @@ #include "mlir/IR/BuiltinDialect.h" #include "mlir/IR/PatternMatch.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertMathToLibm +#include "mlir/Conversion/Passes.cpp.inc" + namespace { // Pattern to convert vector operations to scalar operations. This is needed as // libm calls require scalars. @@ -173,6 +177,8 @@ namespace { struct ConvertMathToLibmPass : public ConvertMathToLibmBase { + using ConvertMathToLibmBase::ConvertMathToLibmBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp b/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp --- a/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp +++ b/mlir/lib/Conversion/MathToSPIRV/MathToSPIRVPass.cpp @@ -11,17 +11,24 @@ //===----------------------------------------------------------------------===// #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 "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertMathToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR Math operations into the SPIR-V dialect. -class ConvertMathToSPIRVPass +struct ConvertMathToSPIRVPass : public ConvertMathToSPIRVBase { + using ConvertMathToSPIRVBase::ConvertMathToSPIRVBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp b/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp --- 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" @@ -22,8 +22,13 @@ #include "mlir/IR/BlockAndValueMapping.h" #include "llvm/ADT/SmallBitVector.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertMemRefToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + namespace { bool isStaticStrideOrOffset(int64_t strideOrOffset) { @@ -2043,7 +2048,7 @@ namespace { struct MemRefToLLVMPass : public ConvertMemRefToLLVMBase { - MemRefToLLVMPass() = default; + using ConvertMemRefToLLVMBase::ConvertMemRefToLLVMBase; void runOnOperation() override { Operation *op = getOperation(); @@ -2071,6 +2076,10 @@ }; } // namespace -std::unique_ptr mlir::createMemRefToLLVMPass() { +std::unique_ptr mlir::createConvertMemRefToLLVMPass() { return std::make_unique(); } + +std::unique_ptr mlir::createConvertMemRefToLLVMPass(const ConvertMemRefToLLVMPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/MemRefToSPIRV/MapMemRefStorageClassPass.cpp b/mlir/lib/Conversion/MemRefToSPIRV/MapMemRefStorageClassPass.cpp --- 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,10 +23,15 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/Support/Debug.h" +#include "../PassDetail.h" + #define DEBUG_TYPE "mlir-map-memref-storage-class" using namespace mlir; +#define GEN_PASS_DEF_MapMemRefStorageClass +#include "mlir/Conversion/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // Mappings //===----------------------------------------------------------------------===// diff --git a/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp b/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp --- a/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp +++ b/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRVPass.cpp @@ -11,17 +11,24 @@ //===----------------------------------------------------------------------===// #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" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertMemRefToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR MemRef operations into the SPIR-V dialect. -class ConvertMemRefToSPIRVPass +struct ConvertMemRefToSPIRVPass : public ConvertMemRefToSPIRVBase { + using ConvertMemRefToSPIRVBase::ConvertMemRefToSPIRVBase; + void runOnOperation() override; }; } // namespace @@ -59,3 +66,8 @@ std::unique_ptr> mlir::createConvertMemRefToSPIRVPass() { return std::make_unique(); } + +std::unique_ptr> mlir::createConvertMemRefToSPIRVPass( + const ConvertMemRefToSPIRVPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp --- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp +++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp @@ -7,15 +7,20 @@ //===----------------------------------------------------------------------===// #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 "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertNVGPUToNVVM +#include "mlir/Conversion/Passes.cpp.inc" + /// Returns the type for the intrinsic given the vectorResultType of the /// `gpu.mma.sync` operation. static Type inferIntrinsicResultType(Type vectorResultType) { @@ -328,7 +333,7 @@ struct ConvertNVGPUToNVVMPass : public ConvertNVGPUToNVVMBase { - ConvertNVGPUToNVVMPass() = default; + using ConvertNVGPUToNVVMBase::ConvertNVGPUToNVVMBase; void runOnOperation() override { RewritePatternSet patterns(&getContext()); diff --git a/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp b/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp --- a/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp +++ b/mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp @@ -6,15 +6,20 @@ // //===----------------------------------------------------------------------===// -#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 "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertOpenACCToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // DataDescriptor implementation //===----------------------------------------------------------------------===// @@ -149,6 +154,8 @@ namespace { struct ConvertOpenACCToLLVMPass : public ConvertOpenACCToLLVMBase { + using ConvertOpenACCToLLVMBase::ConvertOpenACCToLLVMBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp b/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp --- a/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp +++ b/mlir/lib/Conversion/OpenACCToSCF/OpenACCToSCF.cpp @@ -6,15 +6,20 @@ // //===----------------------------------------------------------------------===// -#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/Transforms/DialectConversion.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertOpenACCToSCF +#include "mlir/Conversion/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // Conversion patterns //===----------------------------------------------------------------------===// @@ -57,6 +62,8 @@ namespace { struct ConvertOpenACCToSCFPass : public ConvertOpenACCToSCFBase { + using ConvertOpenACCToSCFBase::ConvertOpenACCToSCFBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp b/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp --- 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" @@ -19,8 +18,13 @@ #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/OpenMP/OpenMPDialect.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertOpenMPToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pattern that converts the region arguments in a single-region OpenMP /// operation to the LLVM dialect. The body of the region is not modified and is @@ -133,6 +137,8 @@ namespace { struct ConvertOpenMPToLLVMPass : public ConvertOpenMPToLLVMBase { + using ConvertOpenMPToLLVMBase::ConvertOpenMPToLLVMBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp b/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp --- a/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp +++ b/mlir/lib/Conversion/PDLToPDLInterp/PDLToPDLInterp.cpp @@ -7,8 +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" #include "mlir/Dialect/PDLInterp/IR/PDLInterp.h" @@ -20,9 +19,15 @@ #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/TypeSwitch.h" +#include "../PassDetail.h" +#include "PredicateTree.h" + using namespace mlir; using namespace mlir::pdl_to_pdl_interp; +#define GEN_PASS_DEF_ConvertPDLToPDLInterp +#include "mlir/Conversion/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // PatternLowering //===----------------------------------------------------------------------===// @@ -910,6 +915,8 @@ namespace { struct PDLToPDLInterpPass : public ConvertPDLToPDLInterpBase { + using ConvertPDLToPDLInterpBase::ConvertPDLToPDLInterpBase; + void runOnOperation() final; }; } // namespace diff --git a/mlir/lib/Conversion/PassDetail.h b/mlir/lib/Conversion/PassDetail.h --- a/mlir/lib/Conversion/PassDetail.h +++ b/mlir/lib/Conversion/PassDetail.h @@ -107,9 +107,6 @@ 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 --- a/mlir/lib/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.cpp +++ b/mlir/lib/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.cpp @@ -7,14 +7,19 @@ //===----------------------------------------------------------------------===// #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" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ReconcileUnrealizedCasts +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// Folds the DAGs of `unrealized_conversion_cast`s that have as exit types @@ -104,7 +109,7 @@ /// Pass to simplify and eliminate unrealized conversion casts. struct ReconcileUnrealizedCasts : public ReconcileUnrealizedCastsBase { - ReconcileUnrealizedCasts() = default; + using ReconcileUnrealizedCastsBase::ReconcileUnrealizedCastsBase; void runOnOperation() override { RewritePatternSet patterns(&getContext()); diff --git a/mlir/lib/Conversion/SCFToControlFlow/SCFToControlFlow.cpp b/mlir/lib/Conversion/SCFToControlFlow/SCFToControlFlow.cpp --- 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,20 @@ #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/Passes.h" +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::scf; +#define GEN_PASS_DEF_ConvertSCFToControlFlow +#include "mlir/Conversion/Passes.cpp.inc" + namespace { struct SCFToControlFlowPass - : public SCFToControlFlowBase { + : public ConvertSCFToControlFlowBase { + using ConvertSCFToControlFlowBase::ConvertSCFToControlFlowBase; + void runOnOperation() override; }; @@ -631,6 +638,6 @@ signalPassFailure(); } -std::unique_ptr mlir::createConvertSCFToCFPass() { +std::unique_ptr mlir::createConvertSCFToControlFlowPass() { return std::make_unique(); } diff --git a/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp b/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp --- a/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp +++ b/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp @@ -7,31 +7,32 @@ //===----------------------------------------------------------------------===// #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/Transforms/DialectConversion.h" - #include "llvm/ADT/ArrayRef.h" #include "llvm/Support/CommandLine.h" +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::scf; +#define GEN_PASS_DEF_ConvertAffineForToGPU +#define GEN_PASS_DEF_ConvertParallelLoopToGpu +#include "mlir/Conversion/Passes.cpp.inc" + 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() = default; - ForLoopMapper(unsigned numBlockDims, unsigned numThreadDims) { - this->numBlockDims = numBlockDims; - this->numThreadDims = numThreadDims; - } + using ConvertAffineForToGPUBase::ConvertAffineForToGPUBase; void runOnOperation() override { for (Operation &op : @@ -47,6 +48,8 @@ struct ParallelLoopToGpuPass : public ConvertParallelLoopToGpuBase { + using ConvertParallelLoopToGpuBase::ConvertParallelLoopToGpuBase; + void runOnOperation() override { RewritePatternSet patterns(&getContext()); populateParallelLoopToGPUPatterns(patterns); @@ -62,15 +65,17 @@ } // namespace -std::unique_ptr> -mlir::createAffineForToGPUPass(unsigned numBlockDims, unsigned numThreadDims) { - return std::make_unique(numBlockDims, numThreadDims); -} std::unique_ptr> mlir::createAffineForToGPUPass() { return std::make_unique(); } +std::unique_ptr> +mlir::createAffineForToGPUPass( + const ConvertAffineForToGPUPassOptions &options) { + return std::make_unique(options); +} + std::unique_ptr mlir::createParallelLoopToGpuPass() { return std::make_unique(); } diff --git a/mlir/lib/Conversion/SCFToOpenMP/SCFToOpenMP.cpp b/mlir/lib/Conversion/SCFToOpenMP/SCFToOpenMP.cpp --- 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" @@ -24,8 +24,13 @@ #include "mlir/IR/SymbolTable.h" #include "mlir/Transforms/DialectConversion.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertSCFToOpenMP +#include "mlir/Conversion/Passes.cpp.inc" + /// Matches a block containing a "simple" reduction. The expected shape of the /// block is as follows. /// @@ -444,6 +449,8 @@ /// A pass converting SCF operations to OpenMP operations. struct SCFToOpenMPPass : public ConvertSCFToOpenMPBase { + using ConvertSCFToOpenMPBase::ConvertSCFToOpenMPBase; + /// 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 --- a/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRVPass.cpp +++ b/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRVPass.cpp @@ -11,20 +11,26 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/SCFToSPIRV/SCFToSPIRVPass.h" +#include "mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h" -#include "../PassDetail.h" #include "mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h" #include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h" #include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h" -#include "mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_SCFToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { struct SCFToSPIRVPass : public SCFToSPIRVBase { + using SCFToSPIRVBase::SCFToSPIRVBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp --- a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp +++ b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp @@ -11,7 +11,6 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h" #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" @@ -27,13 +26,17 @@ #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/SymbolTable.h" #include "mlir/Transforms/DialectConversion.h" - #include "llvm/ADT/DenseMap.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Support/FormatVariadic.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_LowerHostCodeToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + static constexpr const char kSPIRVModule[] = "__spv__"; //===----------------------------------------------------------------------===// @@ -270,9 +273,10 @@ } }; -class LowerHostCodeToLLVM +struct LowerHostCodeToLLVM : public LowerHostCodeToLLVMBase { -public: + using LowerHostCodeToLLVMBase::LowerHostCodeToLLVMBase; + void runOnOperation() override { ModuleOp module = getOperation(); diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp --- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp +++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp @@ -11,18 +11,26 @@ //===----------------------------------------------------------------------===// #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 "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_LowerHostCodeToLLVM +#define GEN_PASS_DEF_ConvertSPIRVToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR SPIR-V operations into LLVM dialect. -class ConvertSPIRVToLLVMPass +struct ConvertSPIRVToLLVMPass : public ConvertSPIRVToLLVMBase { + using ConvertSPIRVToLLVMBase::ConvertSPIRVToLLVMBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Conversion/ShapeToStandard/ConvertShapeConstraints.cpp b/mlir/lib/Conversion/ShapeToStandard/ConvertShapeConstraints.cpp --- 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" +#include "../PassDetail.h" + using namespace mlir; + +#define GEN_PASS_DEF_ConvertShapeConstraints +#include "mlir/Conversion/Passes.cpp.inc" + namespace { #include "ShapeToStandard.cpp.inc" } // namespace @@ -48,8 +53,10 @@ // eager (side-effecting) error handling code. After eager error handling code // is emitted, witnesses are satisfied, so they are replace with // `shape.const_witness true`. -class ConvertShapeConstraints +struct ConvertShapeConstraints : public ConvertShapeConstraintsBase { + using ConvertShapeConstraintsBase::ConvertShapeConstraintsBase; + 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 --- a/mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp +++ b/mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp @@ -23,6 +23,10 @@ using namespace mlir::shape; using namespace mlir::scf; +#define GEN_PASS_DEF_ConvertShapeToStandard +#define GEN_PASS_DEF_ConvertShapeConstraints +#include "mlir/Conversion/Passes.cpp.inc" + /// Conversion patterns. namespace { class AnyOpConversion : public OpConversionPattern { @@ -679,8 +683,9 @@ namespace { /// Conversion pass. -class ConvertShapeToStandardPass +struct ConvertShapeToStandardPass : public ConvertShapeToStandardBase { + using ConvertShapeToStandardBase::ConvertShapeToStandardBase; void runOnOperation() override; }; diff --git a/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp b/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp --- a/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp +++ b/mlir/lib/Conversion/TensorToLinalg/TensorToLinalgPass.cpp @@ -11,18 +11,25 @@ //===----------------------------------------------------------------------===// #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" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertTensorToLinalg +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR Tensor operations into the Linalg dialect. -class ConvertTensorToLinalgPass +struct ConvertTensorToLinalgPass : public ConvertTensorToLinalgBase { + using ConvertTensorToLinalgBase::ConvertTensorToLinalgBase; + 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 --- a/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRVPass.cpp +++ b/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRVPass.cpp @@ -11,19 +11,26 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/TensorToSPIRV/TensorToSPIRVPass.h" -#include "../PassDetail.h" +#include "mlir/Conversion/TensorToSPIRV/TensorToSPIRV.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" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertTensorToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// A pass converting MLIR Tensor operations into the SPIR-V dialect. -class ConvertTensorToSPIRVPass +struct ConvertTensorToSPIRVPass : public ConvertTensorToSPIRVBase { + using ConvertTensorToSPIRVBase::ConvertTensorToSPIRVBase; + void runOnOperation() override { MLIRContext *context = &getContext(); Operation *op = getOperation(); @@ -52,3 +59,8 @@ std::unique_ptr> mlir::createConvertTensorToSPIRVPass() { return std::make_unique(); } + +std::unique_ptr> mlir::createConvertTensorToSPIRVPass( + const ConvertTensorToSPIRVPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/TosaToArith/TosaToArith.cpp b/mlir/lib/Conversion/TosaToArith/TosaToArith.cpp --- a/mlir/lib/Conversion/TosaToArith/TosaToArith.cpp +++ b/mlir/lib/Conversion/TosaToArith/TosaToArith.cpp @@ -11,15 +11,21 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/TosaToArith/TosaToArith.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Tosa/IR/TosaOps.h" #include "mlir/IR/PatternMatch.h" #include "mlir/IR/TypeUtilities.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "../PassDetail.h" + using namespace mlir; using namespace tosa; +#define GEN_PASS_DEF_ConvertTosaToArith +#include "mlir/Conversion/Passes.cpp.inc" + namespace { class ConstOpConverter : public OpRewritePattern { diff --git a/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp b/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp --- a/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp +++ b/mlir/lib/Conversion/TosaToArith/TosaToArithPass.cpp @@ -10,8 +10,8 @@ // //===----------------------------------------------------------------------===// -#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" @@ -21,12 +21,18 @@ #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "../PassDetail.h" + using namespace mlir; -using namespace tosa; +using namespace mlir::tosa; + +#define GEN_PASS_DEF_TosaToArith +#include "mlir/Conversion/Passes.cpp.inc" namespace { struct TosaToArith : public TosaToArithBase { -public: + using TosaToArithBase::TosaToArithBase; + void runOnOperation() override { RewritePatternSet patterns(&getContext()); ConversionTarget target(getContext()); @@ -51,3 +57,8 @@ std::unique_ptr mlir::tosa::createTosaToArith() { return std::make_unique(); } + +std::unique_ptr +mlir::tosa::createTosaToArith(const TosaToArithPassOptions &options) { + return std::make_unique(options); +} diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp --- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp +++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp @@ -20,6 +20,7 @@ #include "mlir/Dialect/Tosa/IR/TosaOps.h" #include "mlir/Dialect/Tosa/Utils/CoversionUtils.h" #include "mlir/Dialect/Utils/ReshapeOpsUtils.h" +#include "mlir/IR/FunctionInterfaces.h" #include "mlir/IR/Matchers.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Transforms/DialectConversion.h" @@ -30,6 +31,10 @@ using namespace mlir; using namespace mlir::tosa; +#define GEN_PASS_DEF_TosaToLinalg +#define GEN_PASS_DEF_TosaToLinalgNamed +#include "mlir/Conversion/Passes.cpp.inc" + template static arith::ConstantOp createConstFromIntAttribute(Operation *op, const std::string &attrName, diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp --- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp +++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp @@ -11,6 +11,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/TosaToLinalg/TosaToLinalg.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Math/IR/Math.h" @@ -24,9 +25,10 @@ #include "mlir/IR/PatternMatch.h" #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" - #include +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::tosa; diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp --- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp +++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamedPass.cpp @@ -10,8 +10,8 @@ // //===----------------------------------------------------------------------===// -#include "../PassDetail.h" #include "mlir/Conversion/TosaToLinalg/TosaToLinalg.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Math/IR/Math.h" @@ -26,11 +26,17 @@ #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_TosaToLinalgNamed +#include "mlir/Conversion/Passes.cpp.inc" + namespace { struct TosaToLinalgNamed : public TosaToLinalgNamedBase { -public: + using TosaToLinalgNamedBase::TosaToLinalgNamedBase; + void getDependentDialects(DialectRegistry ®istry) const override { registry .insert { -public: + using TosaToLinalgBase::TosaToLinalgBase; + void getDependentDialects(DialectRegistry ®istry) const override { registry .insert { -public: + using TosaToSCFBase::TosaToSCFBase; + void runOnOperation() override { RewritePatternSet patterns(&getContext()); ConversionTarget target(getContext()); diff --git a/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp b/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp --- a/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp +++ b/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp @@ -11,12 +11,15 @@ //===----------------------------------------------------------------------===// #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/IR/PatternMatch.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "../PassDetail.h" + using namespace mlir; using namespace tosa; diff --git a/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp b/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp --- a/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp +++ b/mlir/lib/Conversion/TosaToTensor/TosaToTensorPass.cpp @@ -25,9 +25,13 @@ using namespace mlir; using namespace tosa; +#define GEN_PASS_DEF_TosaToTensor +#include "mlir/Conversion/Passes.cpp.inc" + namespace { struct TosaToTensor : public TosaToTensorBase { -public: + using TosaToTensorBase::TosaToTensorBase; + void runOnOperation() override { RewritePatternSet patterns(&getContext()); ConversionTarget target(getContext()); diff --git a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp --- a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp +++ b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp @@ -10,12 +10,8 @@ // //===----------------------------------------------------------------------===// -#include - -#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" @@ -30,9 +26,16 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/Passes.h" #include "llvm/ADT/TypeSwitch.h" +#include + +#include "../PassDetail.h" +#include "NvGpuSupport.h" using namespace mlir; +#define GEN_PASS_DEF_ConvertVectorToGPU +#include "mlir/Conversion/Passes.cpp.inc" + /// For a vector TransferOpType `xferOp`, an empty `indices` vector, and an /// AffineMap representing offsets to apply to indices, the function fills /// `indices` with the original indices plus the offsets. The offsets are @@ -883,10 +886,7 @@ struct ConvertVectorToGPUPass : public ConvertVectorToGPUBase { - - explicit ConvertVectorToGPUPass(bool useNvGpu_) { - useNvGpu.setValue(useNvGpu_); - } + using ConvertVectorToGPUBase::ConvertVectorToGPUBase; void runOnOperation() override { RewritePatternSet patterns(&getContext()); @@ -906,6 +906,10 @@ } // namespace -std::unique_ptr mlir::createConvertVectorToGPUPass(bool useNvGpu) { - return std::make_unique(useNvGpu); +std::unique_ptr mlir::createConvertVectorToGPUPass() { + return std::make_unique(); +} + +std::unique_ptr mlir::createConvertVectorToGPUPass(const ConvertVectorToGPUPassOptions &options) { + return std::make_unique(options); } diff --git a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.cpp b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.cpp --- 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" @@ -26,20 +24,19 @@ #include "mlir/Dialect/X86Vector/X86VectorDialect.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "../PassDetail.h" + using namespace mlir; using namespace mlir::vector; +#define GEN_PASS_DEF_ConvertVectorToLLVM +#include "mlir/Conversion/Passes.cpp.inc" + namespace { struct LowerVectorToLLVMPass : public ConvertVectorToLLVMBase { - LowerVectorToLLVMPass(const LowerVectorToLLVMOptions &options) { - this->reassociateFPReductions = options.reassociateFPReductions; - this->force32BitVectorIndices = options.force32BitVectorIndices; - this->armNeon = options.armNeon; - this->armSVE = options.armSVE; - this->amx = options.amx; - this->x86Vector = options.x86Vector; - } + using ConvertVectorToLLVMBase::ConvertVectorToLLVMBase; + // Override explicitly to allow conditional dialect dependence. void getDependentDialects(DialectRegistry ®istry) const override { registry.insert(); @@ -113,7 +110,11 @@ signalPassFailure(); } -std::unique_ptr> -mlir::createConvertVectorToLLVMPass(const LowerVectorToLLVMOptions &options) { +std::unique_ptr> mlir::createConvertVectorToLLVMPass() { + return std::make_unique(); +} + +std::unique_ptr> mlir::createConvertVectorToLLVMPass( + const ConvertVectorToLLVMPassOptions &options) { return std::make_unique(options); } diff --git a/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp b/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp --- a/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp +++ b/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp @@ -10,11 +10,8 @@ // //===----------------------------------------------------------------------===// -#include - #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" @@ -25,11 +22,17 @@ #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/Passes.h" +#include + +#include "../PassDetail.h" using namespace mlir; using vector::TransferReadOp; using vector::TransferWriteOp; +#define GEN_PASS_DEF_ConvertVectorToSCF +#include "mlir/Conversion/Passes.cpp.inc" + namespace { /// Attribute name used for labeling transfer ops during progressive lowering. diff --git a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp --- a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp +++ b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRVPass.cpp @@ -11,19 +11,25 @@ //===----------------------------------------------------------------------===// #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" +#include "../PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ConvertVectorToSPIRV +#include "mlir/Conversion/Passes.cpp.inc" + namespace { struct ConvertVectorToSPIRVPass : public ConvertVectorToSPIRVBase { + using ConvertVectorToSPIRVBase::ConvertVectorToSPIRVBase; + void runOnOperation() override; }; } // namespace diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp --- a/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp @@ -19,7 +19,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Affine/Analysis/Utils.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" @@ -32,10 +31,15 @@ #include "llvm/Support/Debug.h" #include +#include "PassDetail.h" + #define DEBUG_TYPE "affine-data-copy-generate" using namespace mlir; +#define GEN_PASS_DEF_AffineDataCopyGeneration +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { /// Replaces all loads and stores on memref's living in 'slowMemorySpace' by diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineLoopInvariantCodeMotion.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineLoopInvariantCodeMotion.cpp --- 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,7 +19,6 @@ #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/IR/AffineExpr.h" @@ -32,10 +32,15 @@ #include "llvm/Support/Debug.h" #include "llvm/Support/raw_ostream.h" +#include "PassDetail.h" + #define DEBUG_TYPE "licm" using namespace mlir; +#define GEN_PASS_DEF_AffineLoopInvariantCodeMotion +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { /// Loop invariant code motion (LICM) pass. diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp --- a/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp @@ -10,13 +10,18 @@ // //===----------------------------------------------------------------------===// -#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 "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_AffineLoopNormalize +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { /// Normalize affine.parallel ops so that lower bounds are 0 and steps are 1. diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineParallelize.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineParallelize.cpp --- 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,16 +20,20 @@ #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 "llvm/Support/Debug.h" #include +#include "PassDetail.h" + #define DEBUG_TYPE "affine-parallel" using namespace mlir; +#define GEN_PASS_DEF_AffineParallelize +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { /// Convert all parallel affine.for op into 1-D affine.parallel op. struct AffineParallelize : public AffineParallelizeBase { diff --git a/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp b/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp --- a/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp @@ -16,16 +16,20 @@ #include "mlir/Dialect/Affine/Passes.h" -#include "PassDetail.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/IR/Dominance.h" #include "mlir/Support/LogicalResult.h" #include +#include "PassDetail.h" + #define DEBUG_TYPE "affine-scalrep" using namespace mlir; +#define GEN_PASS_DEF_AffineScalarReplacement +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { struct AffineScalarReplacement : public AffineScalarReplacementBase { diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp --- a/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp @@ -6,7 +6,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/LoopUtils.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" @@ -16,11 +15,16 @@ #include "mlir/Transforms/RegionUtils.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + #define PASS_NAME "loop-coalescing" #define DEBUG_TYPE PASS_NAME using namespace mlir; +#define GEN_PASS_DEF_LoopCoalescing +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { struct LoopCoalescingPass : public LoopCoalescingBase { diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopFusion.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopFusion.cpp --- 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" @@ -33,10 +34,16 @@ #include "llvm/Support/raw_ostream.h" #include #include + +#include "PassDetail.h" + #define DEBUG_TYPE "affine-loop-fusion" using namespace mlir; +#define GEN_PASS_DEF_AffineLoopFusion +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { /// Loop fusion pass. This pass currently supports a greedy fusion policy, /// which fuses loop nests with single-writer/single-reader memref dependences diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopTiling.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopTiling.cpp --- 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,14 +19,19 @@ #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/IR/BlockAndValueMapping.h" #include "mlir/IR/Builders.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" + +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_AffineLoopTiling +#include "mlir/Dialect/Affine/Passes.cpp.inc" + #define DEBUG_TYPE "affine-loop-tile" namespace { diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp --- a/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp @@ -9,11 +9,12 @@ // 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/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/Builders.h" @@ -21,8 +22,13 @@ #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_AffineLoopUnroll +#include "mlir/Dialect/Affine/Passes.cpp.inc" + #define DEBUG_TYPE "affine-loop-unroll" namespace { diff --git a/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp b/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp --- a/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp @@ -33,12 +33,12 @@ // 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/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/BlockAndValueMapping.h" @@ -46,8 +46,13 @@ #include "llvm/ADT/DenseMap.h" #include "llvm/Support/CommandLine.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_AffineLoopUnrollAndJam +#include "mlir/Dialect/Affine/Passes.cpp.inc" + #define DEBUG_TYPE "affine-loop-unroll-jam" namespace { diff --git a/mlir/lib/Dialect/Affine/Transforms/PassDetail.h b/mlir/lib/Dialect/Affine/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Affine/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Affine/Transforms/PassDetail.h @@ -34,9 +34,6 @@ 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 --- 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" @@ -24,10 +25,15 @@ #include "llvm/ADT/DenseMap.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + #define DEBUG_TYPE "affine-pipeline-data-transfer" using namespace mlir; +#define GEN_PASS_DEF_AffinePipelineDataTransfer +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { struct PipelineDataTransfer : public AffinePipelineDataTransferBase { diff --git a/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp b/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp --- a/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp +++ b/mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp @@ -10,18 +10,23 @@ // //===----------------------------------------------------------------------===// -#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/IR/IntegerSet.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "PassDetail.h" + #define DEBUG_TYPE "simplify-affine-structure" using namespace mlir; +#define GEN_PASS_DEF_SimplifyAffineStructures +#include "mlir/Dialect/Affine/Passes.cpp.inc" + namespace { /// Simplifies affine maps and sets appearing in the operations of the Function. diff --git a/mlir/lib/Dialect/Affine/Transforms/SuperVectorize.cpp b/mlir/lib/Dialect/Affine/Transforms/SuperVectorize.cpp --- 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" @@ -26,8 +27,13 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + using namespace mlir; -using namespace vector; +using namespace mlir::vector; + +#define GEN_PASS_DEF_AffineVectorize +#include "mlir/Dialect/Affine/Passes.cpp.inc" /// /// Implements a high-level vectorization strategy on a Function. diff --git a/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp --- a/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Arithmetic/Transforms/Bufferize.cpp @@ -6,18 +6,23 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Bufferization/Transforms/Bufferize.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" +#include "PassDetail.h" + using namespace mlir; -using namespace bufferization; +using namespace mlir::arith; +using namespace mlir::bufferization; + +#define GEN_PASS_DEF_ArithmeticBufferize +#include "mlir/Dialect/Arithmetic/Transforms/Passes.cpp.inc" namespace { /// Pass to bufferize Arithmetic ops. diff --git a/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp b/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp --- a/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp +++ b/mlir/lib/Dialect/Arithmetic/Transforms/ExpandOps.cpp @@ -6,13 +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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::arith; + +#define GEN_PASS_DEF_ArithmeticExpandOps +#include "mlir/Dialect/Arithmetic/Transforms/Passes.cpp.inc" /// Create an integer or index constant. static Value createConst(Location loc, Type type, int value, diff --git a/mlir/lib/Dialect/Arithmetic/Transforms/PassDetail.h b/mlir/lib/Dialect/Arithmetic/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Arithmetic/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Arithmetic/Transforms/PassDetail.h @@ -26,9 +26,6 @@ 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 --- a/mlir/lib/Dialect/Arithmetic/Transforms/UnsignedWhenEquivalent.cpp +++ b/mlir/lib/Dialect/Arithmetic/Transforms/UnsignedWhenEquivalent.cpp @@ -8,17 +8,22 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::arith; using namespace mlir::dataflow; +#define GEN_PASS_DEF_ArithmeticUnsignedWhenEquivalent +#include "mlir/Dialect/Arithmetic/Transforms/Passes.cpp.inc" + /// Succeeds when a value is statically non-negative in that it has a lower /// bound on its value (if it is treated as signed) and that bound is /// non-negative. diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp --- a/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp @@ -10,12 +10,10 @@ // //===----------------------------------------------------------------------===// -#include +#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,10 +24,16 @@ #include "mlir/Support/LLVM.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/RegionUtils.h" +#include + +#include "PassDetail.h" using namespace mlir; using namespace mlir::async; +#define GEN_PASS_DEF_AsyncParallelFor +#include "mlir/Dialect/Async/Passes.cpp.inc" + #define DEBUG_TYPE "async-parallel-for" namespace { diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCounting.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCounting.cpp --- 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,9 +22,15 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/SmallSet.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::async; +#define GEN_PASS_DEF_AsyncRuntimePolicyBasedRefCounting +#define GEN_PASS_DEF_AsyncRuntimeRefCounting +#include "mlir/Dialect/Async/Passes.cpp.inc" + #define DEBUG_TYPE "async-runtime-ref-counting" //===----------------------------------------------------------------------===// diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp --- a/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncRuntimeRefCountingOpt.cpp @@ -10,16 +10,20 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Async/IR/Async.h" #include "mlir/Dialect/Async/Passes.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "llvm/ADT/SmallSet.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::async; +#define GEN_PASS_DEF_AsyncRuntimeRefCountingOpt +#include "mlir/Dialect/Async/Passes.cpp.inc" + #define DEBUG_TYPE "async-ref-counting" namespace { diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp --- a/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp @@ -11,11 +11,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Async/Passes.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,9 +27,14 @@ #include "llvm/ADT/SetVector.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::async; +#define GEN_PASS_DEF_AsyncToAsyncRuntime +#include "mlir/Dialect/Async/Passes.cpp.inc" + #define DEBUG_TYPE "async-to-async-runtime" // Prefix for functions outlined from `async.execute` op regions. static constexpr const char kAsyncFnPrefix[] = "async_execute_fn"; diff --git a/mlir/lib/Dialect/Async/Transforms/PassDetail.h b/mlir/lib/Dialect/Async/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Async/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Async/Transforms/PassDetail.h @@ -27,9 +27,6 @@ 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 --- a/mlir/lib/Dialect/Bufferization/Transforms/AllocTensorElimination.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/AllocTensorElimination.cpp @@ -6,20 +6,24 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::bufferization; +#define GEN_PASS_DEF_AllocTensorElimination +#include "mlir/Dialect/Bufferization/Transforms/Passes.cpp.inc" + /// Return true if all `neededValues` are in scope at the given /// `insertionPoint`. static bool diff --git a/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp b/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp --- a/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/BufferDeallocation.cpp @@ -50,18 +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/MemRef/IR/MemRef.h" #include "llvm/ADT/SetOperations.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::bufferization; +#define GEN_PASS_DEF_BufferDeallocation +#include "mlir/Dialect/Bufferization/Transforms/Passes.cpp.inc" + /// Walks over all immediate return-like terminators in the given region. static LogicalResult walkReturnOperations(Region *region, diff --git a/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp b/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp --- a/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/BufferOptimizations.cpp @@ -11,17 +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/MemRef/IR/MemRef.h" #include "mlir/IR/Operation.h" #include "mlir/Interfaces/LoopLikeInterface.h" #include "mlir/Pass/Pass.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::bufferization; +#define GEN_PASS_DEF_BufferHoisting +#define GEN_PASS_DEF_BufferLoopHoisting +#define GEN_PASS_DEF_PromoteBuffersToStack +#include "mlir/Dialect/Bufferization/Transforms/Passes.cpp.inc" + /// Returns true if the given operation implements a known high-level region- /// based control-flow interface. static bool isKnownControlFlowInterface(Operation *op) { diff --git a/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp b/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp --- a/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/BufferResultsToOutParams.cpp @@ -6,14 +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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::bufferization; + +#define GEN_PASS_DEF_BufferResultsToOutParams +#include "mlir/Dialect/Bufferization/Transforms/Passes.cpp.inc" /// Return `true` if the given MemRef type has a fully dynamic layout. static bool hasFullyDynamicLayoutMap(MemRefType type) { diff --git a/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp --- 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,9 +21,16 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/Passes.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::bufferization; +#define GEN_PASS_DEF_BufferizationBufferize +#define GEN_PASS_DEF_FinalizingBufferize +#define GEN_PASS_DEF_OneShotBufferize +#include "mlir/Dialect/Bufferization/Transforms/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // BufferizeTypeConverter //===----------------------------------------------------------------------===// diff --git a/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp b/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp --- a/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/DropEquivalentBufferResults.cpp @@ -27,14 +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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::bufferization; + +#define GEN_PASS_DEF_DropEquivalentBufferResults +#include "mlir/Dialect/Bufferization/Transforms/Passes.cpp.inc" /// Return the unique ReturnOp that terminates `funcOp`. /// Return nullptr if there is no such unique ReturnOp. diff --git a/mlir/lib/Dialect/Bufferization/Transforms/PassDetail.h b/mlir/lib/Dialect/Bufferization/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Bufferization/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Bufferization/Transforms/PassDetail.h @@ -27,9 +27,6 @@ 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 --- a/mlir/lib/Dialect/Bufferization/Transforms/TensorCopyInsertion.cpp +++ b/mlir/lib/Dialect/Bufferization/Transforms/TensorCopyInsertion.cpp @@ -8,8 +8,6 @@ #include "mlir/Dialect/Bufferization/Transforms/TensorCopyInsertion.h" -#include "PassDetail.h" - #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" @@ -17,9 +15,14 @@ #include "mlir/Dialect/Bufferization/Transforms/OneShotModuleBufferize.h" #include "mlir/Dialect/Bufferization/Transforms/Passes.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::bufferization; +#define GEN_PASS_DEF_TensorCopyInsertion +#include "mlir/Dialect/Bufferization/Transforms/Passes.cpp.inc" + /// Resolve all operands that are also used inside of repetitive regions of the /// same op. Such cases are not fully supported by One-Shot Bufferize. /// diff --git a/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp b/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp --- a/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp +++ b/mlir/lib/Dialect/Func/Transforms/FuncBufferize.cpp @@ -10,18 +10,23 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::func; +#define GEN_PASS_DEF_FuncBufferize +#include "mlir/Dialect/Func/Transforms/Passes.cpp.inc" + namespace { struct FuncBufferizePass : public FuncBufferizeBase { using FuncBufferizeBase::FuncBufferizeBase; diff --git a/mlir/lib/Dialect/Func/Transforms/PassDetail.h b/mlir/lib/Dialect/Func/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Func/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Func/Transforms/PassDetail.h @@ -25,9 +25,6 @@ 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 --- a/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp +++ b/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp @@ -11,10 +11,10 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/GPU/Transforms/Passes.h" + #include "mlir/Dialect/Async/IR/Async.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,7 +24,13 @@ #include "mlir/Transforms/RegionUtils.h" #include "llvm/ADT/TypeSwitch.h" +#include "PassDetail.h" + using namespace mlir; + +#define GEN_PASS_DEF_GpuAsyncRegionPass +#include "mlir/Dialect/GPU/Transforms/Passes.cpp.inc" + namespace { class GpuAsyncRegionPass : public GpuAsyncRegionPassBase { struct ThreadTokenCallback; diff --git a/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp --- 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,8 +27,14 @@ #include "mlir/Support/LLVM.h" #include "mlir/Transforms/RegionUtils.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_GpuKernelOutlining +#define GEN_PASS_DEF_GpuLaunchSinkIndexComputations +#include "mlir/Dialect/GPU/Transforms/Passes.cpp.inc" + template static void createForAllDimensions(OpBuilder &builder, Location loc, SmallVectorImpl &values) { diff --git a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp --- a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp +++ b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp @@ -13,12 +13,18 @@ #include "mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h" -#include "PassDetail.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/GPU/Transforms/Passes.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/IR/AffineMap.h" +#include "PassDetail.h" + +using namespace mlir; + +#define GEN_PASS_DEF_GpuMapParallelLoopsPass +#include "mlir/Dialect/GPU/Transforms/Passes.cpp.inc" + namespace mlir { using scf::ParallelOp; diff --git a/mlir/lib/Dialect/GPU/Transforms/PassDetail.h b/mlir/lib/Dialect/GPU/Transforms/PassDetail.h --- a/mlir/lib/Dialect/GPU/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/GPU/Transforms/PassDetail.h @@ -17,9 +17,6 @@ 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 --- a/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp @@ -7,13 +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 "PassDetail.h" + using namespace mlir; +using namespace mlir::LLVM; + +#define GEN_PASS_DEF_LLVMLegalizeForExport +#include "mlir/Dialect/LLVMIR/Transforms/Passes.cpp.inc" /// If the given block has the same successor with different arguments, /// introduce dummy successor blocks so that all successors of the given block diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp --- a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp @@ -7,13 +7,19 @@ //===----------------------------------------------------------------------===// #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/Transforms/GreedyPatternRewriteDriver.h" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::NVVM; + +#define GEN_PASS_DEF_NVVMOptimizeForTarget +#include "mlir/Dialect/LLVMIR/Transforms/Passes.cpp.inc" namespace { // Replaces fdiv on fp16 with fp32 multiplication with reciprocal plus one diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h b/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h --- a/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h @@ -16,9 +16,6 @@ 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 --- a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp @@ -7,11 +7,17 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h" -#include "PassDetail.h" + #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::LLVM; + +#define GEN_PASS_DEF_LLVMRequestCWrappers +#include "mlir/Dialect/LLVMIR/Transforms/Passes.cpp.inc" namespace { class RequestCWrappersPass diff --git a/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp @@ -6,10 +6,9 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" -#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/Passes.h" #include "mlir/Dialect/Linalg/Transforms/BufferizableOpInterfaceImpl.h" @@ -19,8 +18,13 @@ #include "mlir/IR/Operation.h" #include "mlir/Pass/Pass.h" +#include "PassDetail.h" + using namespace mlir; -using namespace bufferization; +using namespace mlir::bufferization; + +#define GEN_PASS_DEF_LinalgBufferize +#include "mlir/Dialect/Linalg/Passes.cpp.inc" namespace { /// Converts Linalg operations that work on tensor-type operands or results to diff --git a/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp b/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Detensorize.cpp @@ -6,11 +6,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" + #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.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,9 +19,14 @@ #include #include +#include "PassDetail.h" + using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgDetensorize +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + static Value sourceMaterializationCallback(OpBuilder &builder, Type type, ValueRange inputs, Location loc) { assert(inputs.size() == 1); diff --git a/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp b/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp @@ -12,10 +12,10 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.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,11 +27,16 @@ #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + #define DEBUG_TYPE "linalg-drop-unit-dims" using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgFoldUnitExtentDims +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + /// Implements a pass that canonicalizes the uses of unit-extent dimensions for /// broadcasting. For example, /// diff --git a/mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp b/mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp --- 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 -#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,10 +23,16 @@ #include "mlir/IR/PatternMatch.h" #include "mlir/Support/LLVM.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include + +#include "PassDetail.h" using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgElementwiseOpFusion +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + //===---------------------------------------------------------------------===// // Methods and patterns that fuse elementwise `linalg.generic` operations. //===---------------------------------------------------------------------===// diff --git a/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp b/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/ElementwiseToLinalg.cpp @@ -17,6 +17,9 @@ using namespace mlir; +#define GEN_PASS_DEF_ConvertElementwiseToLinalg +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + static bool isElementwiseMappableOpOnRankedTensors(Operation *op) { if (!OpTrait::hasElementwiseMappableTraits(op)) return false; diff --git a/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp b/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Generalization.cpp @@ -11,10 +11,10 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Transforms/Transforms.h" + #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/Passes.h" -#include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" @@ -24,11 +24,16 @@ #include "llvm/ADT/SmallVector.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + #define DEBUG_TYPE "linalg-generalization" using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgGeneralization +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + static LogicalResult generalizeNamedOpPrecondition(LinalgOp linalgOp) { // Check if the operation is a LinalgOp but not a GenericOp. if (isa(linalgOp)) diff --git a/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp b/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/InitTensorToAllocTensor.cpp @@ -6,17 +6,21 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::bufferization; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgInitTensorToAllocTensor +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + namespace { struct InitTensorLoweringPattern : public OpRewritePattern { using OpRewritePattern::OpRewritePattern; diff --git a/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp b/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/InlineScalarOperands.cpp @@ -12,18 +12,23 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.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/IR/AffineExpr.h" #include "mlir/IR/AffineMap.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgInlineScalarOperands +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + namespace { struct InlineScalarOperands : public OpRewritePattern { using OpRewritePattern::OpRewritePattern; diff --git a/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp b/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/LinalgStrategyPasses.cpp @@ -11,15 +11,13 @@ // //===----------------------------------------------------------------------===// -#include +#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/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,11 +31,27 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/LoopInvariantCodeMotionUtils.h" #include "mlir/Transforms/Passes.h" +#include + +#include "PassDetail.h" using namespace mlir; using namespace mlir::vector; using namespace linalg; +#define GEN_PASS_DEF_LinalgStrategyDecomposePass +#define GEN_PASS_DEF_LinalgStrategyEnablePass +#define GEN_PASS_DEF_LinalgStrategyGeneralizePass +#define GEN_PASS_DEF_LinalgStrategyInterchangePass +#define GEN_PASS_DEF_LinalgStrategyLowerVectorsPass +#define GEN_PASS_DEF_LinalgStrategyPadPass +#define GEN_PASS_DEF_LinalgStrategyPeelPass +#define GEN_PASS_DEF_LinalgStrategyRemoveMarkersPass +#define GEN_PASS_DEF_LinalgStrategyTilePass +#define GEN_PASS_DEF_LinalgStrategyTileAndFusePass +#define GEN_PASS_DEF_LinalgStrategyVectorizePass +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + namespace { /// Configurable pass to apply pattern-based tiling and fusion. diff --git a/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp b/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Loops.cpp @@ -6,11 +6,11 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Linalg/Passes.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Arithmetic/Utils/Utils.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,9 +24,16 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/TypeSwitch.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgLowerToAffineLoops +#define GEN_PASS_DEF_LinalgLowerToLoops +#define GEN_PASS_DEF_LinalgLowerToParallelLoops +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + static SmallVector makeCanonicalAffineApplies(OpBuilder &b, Location loc, AffineMap map, ArrayRef vals) { diff --git a/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp b/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/NamedOpConversions.cpp @@ -10,18 +10,23 @@ // 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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::linalg; +#define GEN_PASS_DEF_LinalgNamedOpConversion +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + static llvm::SmallVector getIndicesVector(int start, int end) { return llvm::to_vector<2>(llvm::seq(start, end)); } diff --git a/mlir/lib/Dialect/Linalg/Transforms/PassDetail.h b/mlir/lib/Dialect/Linalg/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Linalg/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Linalg/Transforms/PassDetail.h @@ -44,9 +44,6 @@ 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/Tiling.cpp b/mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp @@ -10,14 +10,12 @@ // //===----------------------------------------------------------------------===// -#include +#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/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,13 +26,18 @@ #include "mlir/IR/AffineMap.h" #include "mlir/Transforms/FoldUtils.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" - #include "llvm/Support/CommandLine.h" +#include + +#include "PassDetail.h" using namespace mlir; using namespace mlir::linalg; using namespace mlir::scf; +#define GEN_PASS_DEF_LinalgTiling +#include "mlir/Dialect/Linalg/Passes.cpp.inc" + #define DEBUG_TYPE "linalg-tiling" static bool isZero(OpFoldResult v) { diff --git a/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp b/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp --- a/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/ExpandOps.cpp @@ -12,16 +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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::memref; + +#define GEN_PASS_DEF_ExpandOps +#include "mlir/Dialect/MemRef/Transforms/Passes.cpp.inc" namespace { diff --git a/mlir/lib/Dialect/MemRef/Transforms/FoldSubViewOps.cpp b/mlir/lib/Dialect/MemRef/Transforms/FoldSubViewOps.cpp --- a/mlir/lib/Dialect/MemRef/Transforms/FoldSubViewOps.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/FoldSubViewOps.cpp @@ -11,17 +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/Vector/IR/VectorOps.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/SmallBitVector.h" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::memref; + +#define GEN_PASS_DEF_FoldSubViewOps +#include "mlir/Dialect/MemRef/Transforms/Passes.cpp.inc" //===----------------------------------------------------------------------===// // Utility functions diff --git a/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp b/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp --- a/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/NormalizeMemRefs.cpp @@ -11,18 +11,24 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/MemRef/Transforms/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" -#include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "llvm/ADT/SmallSet.h" #include "llvm/Support/Debug.h" +#include "PassDetail.h" + #define DEBUG_TYPE "normalize-memrefs" using namespace mlir; +using namespace mlir::memref; + +#define GEN_PASS_DEF_NormalizeMemRefs +#include "mlir/Dialect/MemRef/Transforms/Passes.cpp.inc" namespace { diff --git a/mlir/lib/Dialect/MemRef/Transforms/PassDetail.h b/mlir/lib/Dialect/MemRef/Transforms/PassDetail.h --- a/mlir/lib/Dialect/MemRef/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/MemRef/Transforms/PassDetail.h @@ -41,9 +41,6 @@ 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 --- a/mlir/lib/Dialect/MemRef/Transforms/ResolveShapedTypeResultDims.cpp +++ b/mlir/lib/Dialect/MemRef/Transforms/ResolveShapedTypeResultDims.cpp @@ -11,16 +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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::memref; + +#define GEN_PASS_DEF_ResolveRankedShapeTypeResultDims +#define GEN_PASS_DEF_ResolveShapedTypeResultDims +#include "mlir/Dialect/MemRef/Transforms/Passes.cpp.inc" namespace { /// Fold dim of an operation that implements the InferShapedTypeOpInterface diff --git a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp --- a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp +++ b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp @@ -9,7 +9,7 @@ // This file implements transforms to optimize accesses to shared memory. // //===----------------------------------------------------------------------===// -#include "PassDetail.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -22,9 +22,14 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/Support/MathExtras.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::nvgpu; +#define GEN_PASS_DEF_OptimizeSharedMemory +#include "mlir/Dialect/NVGPU/Passes.cpp.inc" + /// The size of a shared memory line according to NV documentation. constexpr int64_t kSharedMemoryLineSizeBytes = 128; /// We optimize for 128bit accesses, but this can be made an argument in the diff --git a/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp b/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp --- a/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/Bufferize.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" -#include "PassDetail.h" + #include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" @@ -15,9 +15,14 @@ #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/Transforms/DialectConversion.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::scf; +#define GEN_PASS_DEF_SCFBufferize +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + namespace { struct SCFBufferizePass : public SCFBufferizeBase { void runOnOperation() override { diff --git a/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp b/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp --- a/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ForToWhile.cpp @@ -10,19 +10,24 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Transforms.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" +#include "PassDetail.h" + using namespace llvm; using namespace mlir; using scf::ForOp; using scf::WhileOp; +#define GEN_PASS_DEF_SCFForToWhileLoop +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + namespace { struct ForLoopLoweringPattern : public OpRewritePattern { diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp --- a/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/LoopCanonicalization.cpp @@ -11,7 +11,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" @@ -23,9 +22,14 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/TypeSwitch.h" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::scf; +#define GEN_PASS_DEF_SCFForLoopCanonicalization +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + /// A simple, conservative analysis to determine if the loop is shape /// conserving. I.e., the type of the arg-th yielded value is the same as the /// type of the corresponding basic block argument of the loop. diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp --- a/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp @@ -10,17 +10,22 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/SCF/Transforms/Transforms.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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::scf; +#define GEN_PASS_DEF_SCFForLoopRangeFolding +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + namespace { struct ForLoopRangeFolding : public SCFForLoopRangeFoldingBase { diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopSpecialization.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopSpecialization.cpp --- 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,10 +26,17 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "llvm/ADT/DenseMap.h" +#include "PassDetail.h" + using namespace mlir; using scf::ForOp; using scf::ParallelOp; +#define GEN_PASS_DEF_SCFForLoopPeeling +#define GEN_PASS_DEF_SCFForLoopSpecialization +#define GEN_PASS_DEF_SCFParallelLoopSpecialization +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + /// Rewrite a parallel loop with bounds defined by an affine.min with a constant /// into 2 loops after checking if the bounds are equal to that constant. This /// is beneficial if the loop will almost always have the constant bound and diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp --- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopCollapsing.cpp @@ -6,18 +6,23 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + #define DEBUG_TYPE "parallel-loop-collapsing" using namespace mlir; +#define GEN_PASS_DEF_SCFParallelLoopCollapsing +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + namespace { struct ParallelLoopCollapsing : public SCFParallelLoopCollapsingBase { diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp --- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp @@ -10,18 +10,23 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::scf; +#define GEN_PASS_DEF_SCFParallelLoopFusion +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + /// Verify there are no nested ParallelOps. static bool hasNestedParallelOp(ParallelOp ploop) { auto walkResult = diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp --- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp @@ -10,17 +10,22 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::scf; +#define GEN_PASS_DEF_SCFParallelLoopTiling +#include "mlir/Dialect/SCF/Transforms/Passes.cpp.inc" + /// Tile a parallel loop of the form /// scf.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3) /// step (%arg4, %arg5) diff --git a/mlir/lib/Dialect/SCF/Transforms/PassDetail.h b/mlir/lib/Dialect/SCF/Transforms/PassDetail.h --- a/mlir/lib/Dialect/SCF/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/SCF/Transforms/PassDetail.h @@ -35,9 +35,6 @@ 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/SPIRV/Transforms/CanonicalizeGLPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/CanonicalizeGLPass.cpp --- a/mlir/lib/Dialect/SPIRV/Transforms/CanonicalizeGLPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/CanonicalizeGLPass.cpp @@ -6,14 +6,20 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::spirv; + +#define GEN_PASS_DEF_SPIRVCanonicalizeGL +#include "mlir/Dialect/SPIRV/Transforms/Passes.cpp.inc" namespace { class CanonicalizeGLPass final diff --git a/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp --- a/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/DecorateCompositeTypeLayoutPass.cpp @@ -13,15 +13,21 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::spirv; + +#define GEN_PASS_DEF_SPIRVCompositeTypeLayout +#include "mlir/Dialect/SPIRV/Transforms/Passes.cpp.inc" namespace { class SPIRVGlobalVariableOpLayoutInfoDecoration @@ -139,7 +145,7 @@ signalPassFailure(); } -std::unique_ptr> +std::unique_ptr> mlir::spirv::createDecorateSPIRVCompositeTypeLayoutPass() { return std::make_unique(); } diff --git a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp --- a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp @@ -11,16 +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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::spirv; + +#define GEN_PASS_DEF_SPIRVLowerABIAttributes +#include "mlir/Dialect/SPIRV/Transforms/Passes.cpp.inc" /// Creates a global variable for an argument based on the ABI info. static spirv::GlobalVariableOp diff --git a/mlir/lib/Dialect/SPIRV/Transforms/PassDetail.h b/mlir/lib/Dialect/SPIRV/Transforms/PassDetail.h --- a/mlir/lib/Dialect/SPIRV/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/SPIRV/Transforms/PassDetail.h @@ -18,9 +18,6 @@ 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 --- a/mlir/lib/Dialect/SPIRV/Transforms/RewriteInsertsPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/RewriteInsertsPass.cpp @@ -12,13 +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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::spirv; + +#define GEN_PASS_DEF_SPIRVRewriteInsertsPass +#include "mlir/Dialect/SPIRV/Transforms/Passes.cpp.inc" namespace { diff --git a/mlir/lib/Dialect/SPIRV/Transforms/UnifyAliasedResourcePass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/UnifyAliasedResourcePass.cpp --- 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,9 +27,15 @@ #include #include +#include "PassDetail.h" + #define DEBUG_TYPE "spirv-unify-aliased-resource" using namespace mlir; +using namespace mlir::spirv; + +#define GEN_PASS_DEF_SPIRVUnifyAliasedResourcePass +#include "mlir/Dialect/SPIRV/Transforms/Passes.cpp.inc" //===----------------------------------------------------------------------===// // Utility functions diff --git a/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp --- a/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/UpdateVCEPass.cpp @@ -11,19 +11,25 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; +using namespace mlir::spirv; + +#define GEN_PASS_DEF_SPIRVUpdateVCE +#include "mlir/Dialect/SPIRV/Transforms/Passes.cpp.inc" namespace { /// Pass to deduce minimal version/extension/capability requirements for a diff --git a/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp --- a/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Shape/Transforms/Bufferize.cpp @@ -6,19 +6,24 @@ // //===----------------------------------------------------------------------===// -#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/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" +#include "PassDetail.h" + using namespace mlir; using namespace bufferization; +#define GEN_PASS_DEF_ShapeBufferize +#include "mlir/Dialect/Shape/Transforms/Passes.cpp.inc" + namespace { struct ShapeBufferizePass : public ShapeBufferizeBase { void runOnOperation() override { diff --git a/mlir/lib/Dialect/Shape/Transforms/PassDetail.h b/mlir/lib/Dialect/Shape/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Shape/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Shape/Transforms/PassDetail.h @@ -22,9 +22,6 @@ 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 --- a/mlir/lib/Dialect/Shape/Transforms/RemoveShapeConstraints.cpp +++ b/mlir/lib/Dialect/Shape/Transforms/RemoveShapeConstraints.cpp @@ -6,14 +6,19 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/Dialect/Shape/IR/Shape.h" #include "mlir/Dialect/Shape/Transforms/Passes.h" + +#include "mlir/Dialect/Shape/IR/Shape.h" #include "mlir/Transforms/DialectConversion.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_RemoveShapeConstraints +#include "mlir/Dialect/Shape/Transforms/Passes.cpp.inc" + namespace { /// Removal patterns. class RemoveCstrBroadcastableOp diff --git a/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp b/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp --- a/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp +++ b/mlir/lib/Dialect/Shape/Transforms/ShapeToShapeLowering.cpp @@ -6,18 +6,23 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Dialect/Shape/Transforms/Passes.h" + #include "mlir/Dialect/Arithmetic/IR/Arithmetic.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" +#include "PassDetail.h" + using namespace mlir; using namespace mlir::shape; +#define GEN_PASS_DEF_ShapeToShapeLowering +#include "mlir/Dialect/Shape/Transforms/Passes.cpp.inc" + namespace { /// Converts `shape.num_elements` to `shape.reduce`. struct NumElementsOpConverter : public OpRewritePattern { diff --git a/mlir/lib/Dialect/SparseTensor/Pipelines/SparseTensorPipelines.cpp b/mlir/lib/Dialect/SparseTensor/Pipelines/SparseTensorPipelines.cpp --- a/mlir/lib/Dialect/SparseTensor/Pipelines/SparseTensorPipelines.cpp +++ b/mlir/lib/Dialect/SparseTensor/Pipelines/SparseTensorPipelines.cpp @@ -69,10 +69,20 @@ // it to this pipeline. pm.addNestedPass(createConvertLinalgToLoopsPass()); pm.addNestedPass(createConvertVectorToSCFPass()); - pm.addNestedPass(createConvertSCFToCFPass()); + pm.addNestedPass(createConvertSCFToControlFlowPass()); pm.addPass(createLowerAffinePass()); - pm.addPass(createConvertVectorToLLVMPass(options.lowerVectorToLLVMOptions())); - pm.addPass(createMemRefToLLVMPass()); + + ConvertVectorToLLVMPassOptions vectorToLLVMOptions; + vectorToLLVMOptions.reassociateFPReductions = options.reassociateFPReductions; + vectorToLLVMOptions.force32BitVectorIndices = options.indexOptimizations; + vectorToLLVMOptions.armNeon = options.armNeon; + vectorToLLVMOptions.armSVE = options.armSVE; + vectorToLLVMOptions.amx = options.amx; + vectorToLLVMOptions.x86Vector = options.x86Vector; + + pm.addPass(createConvertVectorToLLVMPass(vectorToLLVMOptions)); + + pm.addPass(createConvertMemRefToLLVMPass()); pm.addNestedPass(createConvertComplexToStandardPass()); pm.addNestedPass( mlir::arith::createArithmeticExpandOpsPass()); diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp --- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorPasses.cpp @@ -6,6 +6,8 @@ // //===----------------------------------------------------------------------===// +#include "mlir/Dialect/SparseTensor/Transforms/Passes.h" + #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/Bufferization/IR/Bufferization.h" @@ -16,7 +18,6 @@ #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/SCF/Transforms/Transforms.h" #include "mlir/Dialect/SparseTensor/IR/SparseTensor.h" -#include "mlir/Dialect/SparseTensor/Transforms/Passes.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" @@ -29,8 +30,9 @@ // Passes declaration. //===----------------------------------------------------------------------===// -#define GEN_PASS_CLASSES -#include "mlir/Dialect/SparseTensor/Transforms/Passes.h.inc" +#define GEN_PASS_DEF_SparseTensorConversion +#define GEN_PASS_DEF_Sparsification +#include "mlir/Dialect/SparseTensor/Transforms/Passes.cpp.inc" //===----------------------------------------------------------------------===// // Passes implementation. diff --git a/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp --- a/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Tensor/Transforms/Bufferize.cpp @@ -11,7 +11,7 @@ //===----------------------------------------------------------------------===// #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" @@ -23,9 +23,14 @@ #include "mlir/IR/ImplicitLocOpBuilder.h" #include "mlir/Transforms/DialectConversion.h" +#include "PassDetail.h" + using namespace mlir; using namespace bufferization; +#define GEN_PASS_DEF_TensorBufferize +#include "mlir/Dialect/Tensor/Transforms/Passes.cpp.inc" + namespace { struct TensorBufferizePass : public TensorBufferizeBase { void runOnOperation() override { diff --git a/mlir/lib/Dialect/Tensor/Transforms/PassDetail.h b/mlir/lib/Dialect/Tensor/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Tensor/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Tensor/Transforms/PassDetail.h @@ -26,9 +26,6 @@ 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 --- a/mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp @@ -29,6 +29,9 @@ using namespace mlir; using namespace mlir::tosa; +#define GEN_PASS_DEF_TosaInferShapes +#include "mlir/Dialect/Tosa/Transforms/Passes.cpp.inc" + namespace { void propagateShapesInRegion(Region ®ion); diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp --- a/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaLayerwiseConstantFoldPass.cpp @@ -19,6 +19,9 @@ using namespace mlir; using namespace mlir::tosa; +#define GEN_PASS_DEF_TosaLayerwiseConstantFoldPass +#include "mlir/Dialect/Tosa/Transforms/Passes.cpp.inc" + namespace { template diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp --- a/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaMakeBroadcastable.cpp @@ -10,10 +10,11 @@ // //===----------------------------------------------------------------------===// +#include "mlir/Dialect/Tosa/Transforms/Passes.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/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" @@ -21,6 +22,9 @@ using namespace mlir; using namespace mlir::tosa; +#define GEN_PASS_DEF_TosaMakeBroadcastable +#include "mlir/Dialect/Tosa/Transforms/Passes.cpp.inc" + /// There are two potential ways implementing broadcast: /// a. https://www.tensorflow.org/xla/broadcasting#formal_definition /// b. https://numpy.org/doc/stable/user/basics.broadcasting.html diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp --- a/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaOptionalDecompositions.cpp @@ -12,13 +12,18 @@ // //===----------------------------------------------------------------------===// +#include "mlir/Dialect/Tosa/Transforms/Passes.h" + #include "mlir/Dialect/Tosa/IR/TosaOps.h" #include "mlir/Dialect/Tosa/Transforms/PassDetail.h" -#include "mlir/Dialect/Tosa/Transforms/Passes.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" using namespace mlir; +using namespace mlir::tosa; + +#define GEN_PASS_DEF_TosaOptionalDecompositions +#include "mlir/Dialect/Tosa/Transforms/Passes.cpp.inc" namespace { diff --git a/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp b/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp --- a/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp +++ b/mlir/lib/Dialect/Transform/Transforms/CheckUses.cpp @@ -18,6 +18,10 @@ #include "llvm/ADT/SetOperations.h" using namespace mlir; +using namespace mlir::transform; + +#define GEN_PASS_DEF_CheckUses +#include "mlir/Dialect/Transform/Transforms/Passes.cpp.inc" namespace { @@ -361,6 +365,9 @@ DenseMap> reachableFromCache; }; +#define GEN_PASS_OPTIONS +#include "mlir/Dialect/Transform/Transforms/Passes.h.inc" + #define GEN_PASS_CLASSES #include "mlir/Dialect/Transform/Transforms/Passes.h.inc" diff --git a/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp --- a/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp +++ b/mlir/lib/Dialect/Vector/Transforms/Bufferize.cpp @@ -11,7 +11,7 @@ //===----------------------------------------------------------------------===// #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/MemRef/IR/MemRef.h" @@ -20,8 +20,14 @@ #include "mlir/Dialect/Vector/Transforms/BufferizableOpInterfaceImpl.h" #include "mlir/Dialect/Vector/Transforms/Passes.h" +#include "PassDetail.h" + using namespace mlir; -using namespace bufferization; +using namespace mlir::bufferization; +using namespace mlir::vector; + +#define GEN_PASS_DEF_VectorBufferize +#include "mlir/Dialect/Vector/Transforms/Passes.cpp.inc" namespace { struct VectorBufferizePass : public VectorBufferizeBase { diff --git a/mlir/lib/Dialect/Vector/Transforms/PassDetail.h b/mlir/lib/Dialect/Vector/Transforms/PassDetail.h --- a/mlir/lib/Dialect/Vector/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/Vector/Transforms/PassDetail.h @@ -22,9 +22,6 @@ 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 --- a/mlir/lib/Reducer/OptReductionPass.cpp +++ b/mlir/lib/Reducer/OptReductionPass.cpp @@ -23,6 +23,9 @@ using namespace mlir; +#define GEN_PASS_DEF_OptReduction +#include "mlir/Reducer/Passes.cpp.inc" + namespace { class OptReductionPass : public OptReductionBase { diff --git a/mlir/lib/Reducer/ReductionTreePass.cpp b/mlir/lib/Reducer/ReductionTreePass.cpp --- a/mlir/lib/Reducer/ReductionTreePass.cpp +++ b/mlir/lib/Reducer/ReductionTreePass.cpp @@ -23,7 +23,6 @@ #include "mlir/Reducer/Tester.h" #include "mlir/Rewrite/FrozenRewritePatternSet.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" - #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/SmallVector.h" #include "llvm/Support/Allocator.h" @@ -31,6 +30,9 @@ using namespace mlir; +#define GEN_PASS_DEF_ReductionTree +#include "mlir/Reducer/Passes.cpp.inc" + /// We implicitly number each operation in the region and if an operation's /// number falls into rangeToKeep, we need to keep it and apply the given /// rewrite patterns on it. diff --git a/mlir/lib/TableGen/Pass.cpp b/mlir/lib/TableGen/Pass.cpp --- a/mlir/lib/TableGen/Pass.cpp +++ b/mlir/lib/TableGen/Pass.cpp @@ -87,9 +87,10 @@ return def->getValueAsString("description"); } -StringRef Pass::getConstructor() const { - return def->getValueAsString("constructor"); +Optional Pass::getConstructor() const { + return def->getValueAsOptionalString("constructor"); } + ArrayRef Pass::getDependentDialects() const { return dependentDialects; } diff --git a/mlir/lib/Transforms/CSE.cpp b/mlir/lib/Transforms/CSE.cpp --- a/mlir/lib/Transforms/CSE.cpp +++ b/mlir/lib/Transforms/CSE.cpp @@ -11,7 +11,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/IR/Dominance.h" #include "mlir/Interfaces/SideEffectInterfaces.h" #include "mlir/Pass/Pass.h" @@ -23,8 +22,13 @@ #include "llvm/Support/RecyclingAllocator.h" #include +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_CSE +#include "mlir/Transforms/Passes.cpp.inc" + namespace { struct SimpleOperationInfo : public llvm::DenseMapInfo { static unsigned getHashValue(const Operation *opC) { diff --git a/mlir/lib/Transforms/Canonicalizer.cpp b/mlir/lib/Transforms/Canonicalizer.cpp --- a/mlir/lib/Transforms/Canonicalizer.cpp +++ b/mlir/lib/Transforms/Canonicalizer.cpp @@ -11,13 +11,17 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" #include "mlir/Transforms/Passes.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_Canonicalizer +#include "mlir/Transforms/Passes.cpp.inc" + namespace { /// Canonicalize operations in nested regions. struct Canonicalizer : public CanonicalizerBase { diff --git a/mlir/lib/Transforms/ControlFlowSink.cpp b/mlir/lib/Transforms/ControlFlowSink.cpp --- a/mlir/lib/Transforms/ControlFlowSink.cpp +++ b/mlir/lib/Transforms/ControlFlowSink.cpp @@ -13,7 +13,6 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/IR/Dominance.h" #include "mlir/Interfaces/ControlFlowInterfaces.h" #include "mlir/Interfaces/SideEffectInterfaces.h" @@ -21,8 +20,13 @@ #include "mlir/Transforms/Passes.h" #include "mlir/Transforms/SideEffectUtils.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_ControlFlowSink +#include "mlir/Transforms/Passes.cpp.inc" + namespace { /// A control-flow sink pass. struct ControlFlowSink : public ControlFlowSinkBase { diff --git a/mlir/lib/Transforms/Inliner.cpp b/mlir/lib/Transforms/Inliner.cpp --- 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,14 +22,18 @@ #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" +#include "PassDetail.h" + #define DEBUG_TYPE "inlining" using namespace mlir; +#define GEN_PASS_DEF_Inliner +#include "mlir/Transforms/Passes.cpp.inc" + /// This function implements the default inliner optimization pipeline. static void defaultInlinerOptPipeline(OpPassManager &pm) { pm.addPass(createCanonicalizerPass()); diff --git a/mlir/lib/Transforms/LocationSnapshot.cpp b/mlir/lib/Transforms/LocationSnapshot.cpp --- a/mlir/lib/Transforms/LocationSnapshot.cpp +++ b/mlir/lib/Transforms/LocationSnapshot.cpp @@ -7,15 +7,19 @@ //===----------------------------------------------------------------------===// #include "mlir/Transforms/LocationSnapshot.h" -#include "PassDetail.h" #include "mlir/IR/AsmState.h" #include "mlir/IR/Builders.h" #include "mlir/Support/FileUtilities.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/ToolOutputFile.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_LocationSnapshot +#include "mlir/Transforms/Passes.cpp.inc" + /// This function generates new locations from the given IR by snapshotting the /// IR to the given stream, and using the printed locations within that stream. /// If a 'tag' is non-empty, the generated locations are represented as a diff --git a/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp b/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp --- a/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp +++ b/mlir/lib/Transforms/LoopInvariantCodeMotion.cpp @@ -10,14 +10,18 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/Interfaces/LoopLikeInterface.h" #include "mlir/Transforms/LoopInvariantCodeMotionUtils.h" #include "mlir/Transforms/Passes.h" #include "mlir/Transforms/SideEffectUtils.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_LoopInvariantCodeMotion +#include "mlir/Transforms/Passes.cpp.inc" + namespace { /// Loop invariant code motion (LICM) pass. struct LoopInvariantCodeMotion diff --git a/mlir/lib/Transforms/OpStats.cpp b/mlir/lib/Transforms/OpStats.cpp --- a/mlir/lib/Transforms/OpStats.cpp +++ b/mlir/lib/Transforms/OpStats.cpp @@ -6,17 +6,22 @@ // //===----------------------------------------------------------------------===// -#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" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_PrintOpStats +#include "mlir/Transforms/Passes.cpp.inc" + namespace { struct PrintOpStatsPass : public PrintOpStatsBase { explicit PrintOpStatsPass(raw_ostream &os) : os(os) {} diff --git a/mlir/lib/Transforms/PassDetail.h b/mlir/lib/Transforms/PassDetail.h --- a/mlir/lib/Transforms/PassDetail.h +++ b/mlir/lib/Transforms/PassDetail.h @@ -14,8 +14,7 @@ #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 --- 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,11 +23,15 @@ #include "mlir/Interfaces/SideEffectInterfaces.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/FoldUtils.h" -#include "mlir/Transforms/Passes.h" + +#include "PassDetail.h" using namespace mlir; using namespace mlir::dataflow; +#define GEN_PASS_DEF_SCCP +#include "mlir/Transforms/Passes.cpp.inc" + //===----------------------------------------------------------------------===// // SCCP Rewrites //===----------------------------------------------------------------------===// diff --git a/mlir/lib/Transforms/StripDebugInfo.cpp b/mlir/lib/Transforms/StripDebugInfo.cpp --- a/mlir/lib/Transforms/StripDebugInfo.cpp +++ b/mlir/lib/Transforms/StripDebugInfo.cpp @@ -6,14 +6,18 @@ // //===----------------------------------------------------------------------===// +#include "mlir/Transforms/Passes.h" + #include "PassDetail.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/Operation.h" #include "mlir/Pass/Pass.h" -#include "mlir/Transforms/Passes.h" using namespace mlir; +#define GEN_PASS_DEF_StripDebugInfo +#include "mlir/Transforms/Passes.cpp.inc" + namespace { struct StripDebugInfo : public StripDebugInfoBase { void runOnOperation() override; diff --git a/mlir/lib/Transforms/SymbolDCE.cpp b/mlir/lib/Transforms/SymbolDCE.cpp --- a/mlir/lib/Transforms/SymbolDCE.cpp +++ b/mlir/lib/Transforms/SymbolDCE.cpp @@ -11,12 +11,16 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" #include "mlir/IR/SymbolTable.h" #include "mlir/Transforms/Passes.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_SymbolDCE +#include "mlir/Transforms/Passes.cpp.inc" + namespace { struct SymbolDCE : public SymbolDCEBase { void runOnOperation() override; diff --git a/mlir/lib/Transforms/SymbolPrivatize.cpp b/mlir/lib/Transforms/SymbolPrivatize.cpp --- a/mlir/lib/Transforms/SymbolPrivatize.cpp +++ b/mlir/lib/Transforms/SymbolPrivatize.cpp @@ -11,12 +11,17 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" -#include "mlir/IR/SymbolTable.h" #include "mlir/Transforms/Passes.h" +#include "mlir/IR/SymbolTable.h" + +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_SymbolPrivatize +#include "mlir/Transforms/Passes.cpp.inc" + namespace { struct SymbolPrivatize : public SymbolPrivatizeBase { explicit SymbolPrivatize(ArrayRef excludeSymbols); diff --git a/mlir/lib/Transforms/TopologicalSort.cpp b/mlir/lib/Transforms/TopologicalSort.cpp --- a/mlir/lib/Transforms/TopologicalSort.cpp +++ b/mlir/lib/Transforms/TopologicalSort.cpp @@ -6,12 +6,18 @@ // //===----------------------------------------------------------------------===// -#include "PassDetail.h" +#include "mlir/Transforms/Passes.h" + #include "mlir/IR/RegionKindInterface.h" #include "mlir/Transforms/TopologicalSortUtils.h" +#include "PassDetail.h" + using namespace mlir; +#define GEN_PASS_DEF_TopologicalSort +#include "mlir/Transforms/Passes.cpp.inc" + namespace { struct TopologicalSortPass : public TopologicalSortBase { void runOnOperation() override { diff --git a/mlir/lib/Transforms/ViewOpGraph.cpp b/mlir/lib/Transforms/ViewOpGraph.cpp --- a/mlir/lib/Transforms/ViewOpGraph.cpp +++ b/mlir/lib/Transforms/ViewOpGraph.cpp @@ -18,6 +18,9 @@ using namespace mlir; +#define GEN_PASS_DEF_ViewOpGraph +#include "mlir/Transforms/Passes.cpp.inc" + static const StringRef kLineStyleControlFlow = "dashed"; static const StringRef kLineStyleDataFlow = "solid"; static const StringRef kShapeNode = "ellipse"; diff --git a/mlir/tools/mlir-tblgen/PassCAPIGen.cpp b/mlir/tools/mlir-tblgen/PassCAPIGen.cpp --- a/mlir/tools/mlir-tblgen/PassCAPIGen.cpp +++ b/mlir/tools/mlir-tblgen/PassCAPIGen.cpp @@ -90,6 +90,14 @@ } )"; +static void emitConstructorCall(const Pass &pass, raw_ostream &os) { + if (auto constructor = pass.getConstructor()) { + os << *constructor; + } else { + os << llvm::formatv("create{0}Pass()", pass.getDef()->getName()); + } +} + static bool emitCAPIImpl(const llvm::RecordKeeper &records, raw_ostream &os) { os << "/* Autogenerated by mlir-tblgen; don't manually edit. */"; os << llvm::formatv(passGroupRegistrationCode, groupName); @@ -97,8 +105,12 @@ for (const auto *def : records.getAllDerivedDefinitions("PassBase")) { Pass pass(def); StringRef defName = pass.getDef()->getName(); - os << llvm::formatv(passCreateDef, groupName, defName, - pass.getConstructor()); + + std::string constructorCall; + llvm::raw_string_ostream constructorCallOs(constructorCall); + emitConstructorCall(pass, constructorCallOs); + + os << llvm::formatv(passCreateDef, groupName, defName, constructorCall); } return false; } diff --git a/mlir/tools/mlir-tblgen/PassGen.cpp b/mlir/tools/mlir-tblgen/PassGen.cpp --- a/mlir/tools/mlir-tblgen/PassGen.cpp +++ b/mlir/tools/mlir-tblgen/PassGen.cpp @@ -27,6 +27,12 @@ groupName("name", llvm::cl::desc("The name of this group of passes"), llvm::cl::cat(passGenCat)); +const char *const passHeader = R"( +//===----------------------------------------------------------------------===// +// {0} +//===----------------------------------------------------------------------===// +)"; + //===----------------------------------------------------------------------===// // GEN: Pass base class generation //===----------------------------------------------------------------------===// @@ -38,10 +44,6 @@ /// {2): The command line argument for the pass. /// {3}: The dependent dialects registration. const char *const passDeclBegin = R"( -//===----------------------------------------------------------------------===// -// {0} -//===----------------------------------------------------------------------===// - template class {0}Base : public {1} { public: @@ -79,6 +81,10 @@ {4} } + {0}Base(const {0}PassOptions &options) : {0}Base() {{ + {5} + } + /// Explicitly declare the TypeID for this class. We declare an explicit private /// instantiation because Pass classes should only be visible by the current /// library. @@ -119,33 +125,6 @@ } } -static void emitPassDecl(const Pass &pass, raw_ostream &os) { - StringRef defName = pass.getDef()->getName(); - std::string dependentDialectRegistrations; - { - llvm::raw_string_ostream dialectsOs(dependentDialectRegistrations); - for (StringRef dependentDialect : pass.getDependentDialects()) - dialectsOs << llvm::formatv(dialectRegistrationTemplate, - dependentDialect); - } - os << llvm::formatv(passDeclBegin, defName, pass.getBaseClass(), - pass.getArgument(), pass.getSummary(), - dependentDialectRegistrations); - emitPassOptionDecls(pass, os); - emitPassStatisticDecls(pass, os); - os << "};\n"; -} - -/// Emit the code for registering each of the given passes with the global -/// PassRegistry. -static void emitPassDecls(ArrayRef passes, raw_ostream &os) { - os << "#ifdef GEN_PASS_CLASSES\n"; - for (const Pass &pass : passes) - emitPassDecl(pass, os); - os << "#undef GEN_PASS_CLASSES\n"; - os << "#endif // GEN_PASS_CLASSES\n"; -} - //===----------------------------------------------------------------------===// // GEN: Pass registration generation //===----------------------------------------------------------------------===// @@ -178,40 +157,147 @@ inline void register{0}Passes() {{ )"; +std::vector getPasses(const llvm::RecordKeeper &recordKeeper) { + std::vector passes; + + for (const auto *def : recordKeeper.getAllDerivedDefinitions("PassBase")) + passes.emplace_back(def); + + return passes; +} + +static void emitPassOptionsStruct(const Pass &pass, raw_ostream &os) { + StringRef passName = pass.getDef()->getName(); + ArrayRef options = pass.getOptions(); + + os << llvm::formatv("struct {0}PassOptions {{\n", passName); + + for (const PassOption &opt : options) { + os.indent(2) << llvm::formatv("{0} {1}", opt.getType(), + opt.getCppVariableName()); + + if (Optional defaultVal = opt.getDefaultValue()) + os << " = " << defaultVal; + + os << ";\n"; + } + + os << "};\n"; +} + +static void emitPassDecls(const Pass &pass, raw_ostream &os) { + StringRef passName = pass.getDef()->getName(); + + os << "#ifdef GEN_PASS_DECL_" << passName << "\n"; + os << llvm::formatv(passHeader, passName); + + emitPassOptionsStruct(pass, os); + + if (!pass.getConstructor()) { + // Default constructor declaration. + os << "std::unique_ptr<::mlir::Pass> create" << passName << "Pass();\n"; + + // Declaration of the constructor with options. + os << llvm::formatv("std::unique_ptr<::mlir::Pass> create{0}Pass(const {0}PassOptions &options);\n", passName); + } + + os << "#undef GEN_PASS_DECL_" << passName << "\n"; + os << "#endif // GEN_PASS_DECL_" << passName << "\n"; +} + +static void emitConstructorCall(const Pass &pass, raw_ostream &os) { + if (auto constructor = pass.getConstructor()) { + os << *constructor; + } else { + os << llvm::formatv("create{0}Pass()", pass.getDef()->getName()); + } +} + /// Emit the code for registering each of the given passes with the global /// PassRegistry. -static void emitRegistration(ArrayRef passes, raw_ostream &os) { +static void emitRegistrations(llvm::ArrayRef passes, raw_ostream &os) { os << "#ifdef GEN_PASS_REGISTRATION\n"; + for (const Pass &pass : passes) { + std::string constructorCall; + llvm::raw_string_ostream constructorCallOs(constructorCall); + emitConstructorCall(pass, constructorCallOs); + os << llvm::formatv(passRegistrationCode, pass.getDef()->getName(), - pass.getConstructor()); + constructorCall); } os << llvm::formatv(passGroupRegistrationCode, groupName); + for (const Pass &pass : passes) os << " register" << pass.getDef()->getName() << "Pass();\n"; + os << "}\n"; os << "#undef GEN_PASS_REGISTRATION\n"; os << "#endif // GEN_PASS_REGISTRATION\n"; } -//===----------------------------------------------------------------------===// -// GEN: Registration hooks -//===----------------------------------------------------------------------===// - static void emitDecls(const llvm::RecordKeeper &recordKeeper, raw_ostream &os) { + std::vector passes = getPasses(recordKeeper); os << "/* Autogenerated by mlir-tblgen; don't manually edit */\n"; - std::vector passes; - for (const auto *def : recordKeeper.getAllDerivedDefinitions("PassBase")) - passes.emplace_back(def); - emitPassDecls(passes, os); - emitRegistration(passes, os); + for (const Pass &pass : passes) + emitPassDecls(pass, os); + + emitRegistrations(passes, os); +} + +static void emitPassDefs(const Pass &pass, raw_ostream &os) { + StringRef defName = pass.getDef()->getName(); + + os << "#ifdef GEN_PASS_DEF_" << defName << "\n"; + os << llvm::formatv(passHeader, defName); + + std::string dependentDialectRegistrations; + { + llvm::raw_string_ostream dialectsOs(dependentDialectRegistrations); + for (StringRef dependentDialect : pass.getDependentDialects()) + dialectsOs << llvm::formatv(dialectRegistrationTemplate, + dependentDialect); + } + + std::string optionsConstructor; + { + llvm::raw_string_ostream optionsConstructorOs(optionsConstructor); + for (const PassOption &opt : pass.getOptions()) + optionsConstructorOs << llvm::formatv("{0} = options.{0};\n", opt.getCppVariableName()); + } + + os << llvm::formatv(passDeclBegin, defName, pass.getBaseClass(), + pass.getArgument(), pass.getSummary(), + dependentDialectRegistrations, optionsConstructor); + + emitPassOptionDecls(pass, os); + emitPassStatisticDecls(pass, os); + os << "};\n"; + + os << "#undef GEN_PASS_DEF_" << defName << "\n"; + os << "#endif // GEN_PASS_DEF_" << defName << "\n"; +} + +static void emitDefs(const llvm::RecordKeeper &recordKeeper, raw_ostream &os) { + std::vector passes = getPasses(recordKeeper); + os << "/* Autogenerated by mlir-tblgen; don't manually edit */\n"; + + for (const Pass &pass : passes) + emitPassDefs(pass, os); } static mlir::GenRegistration - genRegister("gen-pass-decls", "Generate pass declarations", + genPassDecls("gen-pass-decls", "Generate pass declarations", + [](const llvm::RecordKeeper &records, raw_ostream &os) { + emitDecls(records, os); + return false; + }); + +static mlir::GenRegistration + genPassDefs("gen-pass-defs", "Generate pass definitions", [](const llvm::RecordKeeper &records, raw_ostream &os) { - emitDecls(records, os); + emitDefs(records, os); return false; }); diff --git a/mlir/unittests/ExecutionEngine/Invoke.cpp b/mlir/unittests/ExecutionEngine/Invoke.cpp --- a/mlir/unittests/ExecutionEngine/Invoke.cpp +++ b/mlir/unittests/ExecutionEngine/Invoke.cpp @@ -46,9 +46,8 @@ /// dialects lowering to LLVM Dialect. static LogicalResult lowerToLLVMDialect(ModuleOp module) { PassManager pm(module.getContext()); - pm.addPass(mlir::createMemRefToLLVMPass()); - pm.addNestedPass( - mlir::arith::createConvertArithmeticToLLVMPass()); + pm.addPass(mlir::createConvertMemRefToLLVMPass()); + pm.addNestedPass(mlir::createConvertArithmeticToLLVMPass()); pm.addPass(mlir::createConvertFuncToLLVMPass()); pm.addPass(mlir::createReconcileUnrealizedCastsPass()); return pm.run(module); diff --git a/mlir/unittests/TableGen/CMakeLists.txt b/mlir/unittests/TableGen/CMakeLists.txt --- a/mlir/unittests/TableGen/CMakeLists.txt +++ b/mlir/unittests/TableGen/CMakeLists.txt @@ -5,6 +5,7 @@ set(LLVM_TARGET_DEFINITIONS passes.td) mlir_tablegen(PassGenTest.h.inc -gen-pass-decls -name TableGenTest) +mlir_tablegen(PassGenTest.cpp.inc -gen-pass-defs -name TableGenTest) add_public_tablegen_target(MLIRTableGenTestPassIncGen) add_mlir_unittest(MLIRTableGenTests diff --git a/mlir/unittests/TableGen/PassGenTest.cpp b/mlir/unittests/TableGen/PassGenTest.cpp --- a/mlir/unittests/TableGen/PassGenTest.cpp +++ b/mlir/unittests/TableGen/PassGenTest.cpp @@ -12,21 +12,27 @@ std::unique_ptr createTestPass(int v = 0); +#define GEN_PASS_DECL_TestPass #define GEN_PASS_REGISTRATION #include "PassGenTest.h.inc" -#define GEN_PASS_CLASSES -#include "PassGenTest.h.inc" +#define GEN_PASS_DEF_TestPass +#include "PassGenTest.cpp.inc" struct TestPass : public TestPassBase { explicit TestPass(int v) : extraVal(v) {} + TestPass(int v, const TestPassPassOptions &options) + : TestPassBase(options), extraVal(v) {} + void runOnOperation() override {} std::unique_ptr clone() const { return TestPassBase::clone(); } + unsigned getTestOption() const { return testOption; } + int extraVal; }; @@ -34,6 +40,11 @@ return std::make_unique(v); } +std::unique_ptr createTestPass(int v, + const TestPassPassOptions &options) { + return std::make_unique(v, options); +} + TEST(PassGenTest, PassClone) { mlir::MLIRContext context; @@ -46,3 +57,18 @@ EXPECT_EQ(unwrap(origPass)->extraVal, unwrap(clonePass)->extraVal); } + +TEST(PassGenTest, PassOptions) { + mlir::MLIRContext context; + + TestPassPassOptions options; + options.testOption = 57; + + const auto unwrap = [](const std::unique_ptr &pass) { + return static_cast(pass.get()); + }; + + const auto pass = createTestPass(10, options); + + EXPECT_EQ(unwrap(pass)->getTestOption(), 57); +} diff --git a/mlir/unittests/TableGen/passes.td b/mlir/unittests/TableGen/passes.td --- a/mlir/unittests/TableGen/passes.td +++ b/mlir/unittests/TableGen/passes.td @@ -15,5 +15,7 @@ let constructor = "::createTestPass()"; - let options = RewritePassUtils.options; + let options = [ + Option<"testOption", "testOption", "unsigned", "0", "Test option"> + ]; }