From 30994be76cfc11bbb91bce5819dac1fbc38a212a Mon Sep 17 00:00:00 2001 From: Thomas Joerg Date: Tue, 27 Oct 2020 06:55:28 -0700 Subject: [PATCH] Integrate LLVM at llvm/llvm-project@0fc1aa22ee6a Updates LLVM usage to match [0fc1aa22ee6a](https://github.com/llvm/llvm-project/commit/0fc1aa22ee6a) PiperOrigin-RevId: 339239851 Change-Id: I78072452756354de403b60c19f2441b250a63f53 --- tensorflow/compiler/mlir/hlo/BUILD | 2 + .../transforms/chlo_legalize_to_hlo_pass.cc | 2 +- .../mhlo/transforms/hlo_legalize_to_lhlo.cc | 3 +- .../legalize_gather_to_torch_index_select.cc | 4 +- .../mhlo/transforms/legalize_to_linalg.cc | 4 +- .../mhlo/transforms/legalize_to_standard.cc | 4 +- ...legalize_trigonometric_to_approximation.cc | 4 +- .../mhlo/transforms/lhlo_fuse_linalg.cc | 5 ++- .../transforms/lhlo_legalize_to_affine.cc | 6 +-- .../mhlo/transforms/lhlo_legalize_to_gpu.cc | 2 +- .../transforms/lhlo_legalize_to_llvm_pass.cc | 2 +- .../lhlo_legalize_to_parallel_loops.cc | 2 +- .../Dialect/mhlo/transforms/lower_complex.cc | 4 +- .../mhlo/transforms/lower_general_dot.cc | 4 +- .../transforms/materialize_broadcasts_pass.cc | 2 +- .../mhlo/transforms/optimize_mhlo_pass.cc | 4 +- .../transforms/test_infer_shaped_type_pass.cc | 4 +- .../mhlo/transforms/transform_unranked_hlo.cc | 3 +- .../mhlo/transforms/unfuse_batch_norm_pass.cc | 4 +- tensorflow/compiler/mlir/lite/BUILD | 2 + .../mlir/lite/quantization/tensorflow/BUILD | 4 +- .../quantization/tensorflow/tf_to_quant.cc | 4 +- .../mlir/lite/transforms/dilated_conv.cc | 4 +- .../mlir/lite/transforms/legalize_tf.cc | 3 +- .../transforms/lower_static_tensor_list.cc | 2 +- .../compiler/mlir/lite/transforms/optimize.cc | 16 ++++--- .../transforms/optimize_functional_ops.cc | 4 +- .../mlir/lite/transforms/post_quantize.cc | 11 +++-- .../mlir/lite/transforms/prepare_quantize.cc | 4 +- .../mlir/lite/transforms/prepare_tf.cc | 24 +++++----- .../compiler/mlir/lite/transforms/quantize.cc | 3 +- tensorflow/compiler/mlir/tensorflow/BUILD | 2 + .../transforms/batchmatmul_to_einsum.cc | 3 +- .../transforms/contraction_fusion.cc | 4 +- .../transforms/decompose_resource_ops_pass.cc | 3 +- .../mlir/tensorflow/transforms/einsum.cc | 3 +- .../tensorflow/transforms/fold_broadcast.cc | 4 +- .../transforms/fused_kernel_matcher.cc | 3 +- .../mlir/tensorflow/transforms/gpu_fusion.cc | 3 +- .../transforms/init_text_file_to_import.cc | 3 +- .../tensorflow/transforms/legalize_hlo.cc | 3 +- .../tensorflow/transforms/lower_tf_pass.cc | 3 +- .../mlir/tensorflow/transforms/optimize.cc | 3 +- .../tensor_device_copy_conversion.cc | 3 +- .../transforms/tf_data_optimization_pass.cc | 3 +- .../transforms/unroll_batch_matmul.cc | 3 +- tensorflow/compiler/mlir/tfjs/BUILD | 1 + .../compiler/mlir/tfjs/transforms/optimize.cc | 3 +- .../compiler/mlir/tfr/passes/decompose.cc | 4 +- .../compiler/mlir/tfr/passes/raise_to_tf.cc | 4 +- .../mlir/tools/kernel_gen/kernel_creator.cc | 1 + .../kernel_gen/transforms/bufferize_pass.cc | 2 +- .../transforms/embed_tf_framework_pass.cc | 2 +- .../transforms/materialize_broadcasts_pass.cc | 2 +- .../parallel_loops_to_sequential.cc | 3 +- .../transforms/shape_to_descriptors_pass.cc | 2 +- .../transforms/tf_kernel_to_llvm_pass.cc | 2 +- .../transforms/unfuse_batch_norm_pass.cc | 3 +- tensorflow/compiler/mlir/xla/BUILD | 1 + .../mlir/xla/transforms/legalize_tf.cc | 6 +-- .../xla/transforms/legalize_tf_with_tf2xla.cc | 4 +- .../xla/service/mlir_gpu/kernel_lowering.cc | 16 ++++--- tensorflow/workspace.bzl | 4 +- third_party/mlir/BUILD | 45 +++++++++++++++++++ 64 files changed, 192 insertions(+), 105 deletions(-) diff --git a/tensorflow/compiler/mlir/hlo/BUILD b/tensorflow/compiler/mlir/hlo/BUILD index 1636bbb89ee..a719f303d3d 100644 --- a/tensorflow/compiler/mlir/hlo/BUILD +++ b/tensorflow/compiler/mlir/hlo/BUILD @@ -538,6 +538,7 @@ cc_library( "@llvm-project//mlir:IR", "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) @@ -764,6 +765,7 @@ cc_library( "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo_pass.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo_pass.cc index d2f415d91f9..ce63f6b59a0 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo_pass.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo_pass.cc @@ -49,7 +49,7 @@ struct ChloLegalizeToHloPass chlo::PopulateLegalizeChloToHloPatterns(&getContext(), &conversionPatterns); if (failed(applyPartialConversion(getFunction(), conversionTarget, - conversionPatterns))) { + std::move(conversionPatterns)))) { return signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc index 1ce93b977c6..1cf39150368 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc @@ -555,7 +555,8 @@ struct HloLegalizeToLhlo &context, converter, patterns); populateShapeStructuralTypeConversionsAndLegality(&context, converter, patterns, target); - if (failed(applyPartialConversion(getOperation(), target, patterns))) + if (failed(applyPartialConversion(getOperation(), target, + std::move(patterns)))) signalPassFailure(); } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_gather_to_torch_index_select.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_gather_to_torch_index_select.cc index 59cd3381133..b3722e3bf47 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_gather_to_torch_index_select.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_gather_to_torch_index_select.cc @@ -17,8 +17,8 @@ limitations under the License. #include "mlir-hlo/Dialect/mhlo/transforms/passes.h" #include "mlir-hlo/Dialect/mhlo/transforms/rewriters.h" #include "mlir/IR/Function.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/Pass/Pass.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { @@ -133,7 +133,7 @@ struct LegalizeGatherToTorchIndexSelectPass void runOnFunction() override { OwningRewritePatternList patterns; PopulateGatherToTorchIndexSelectPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } }; } // namespace diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_linalg.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_linalg.cc index 834384e37f3..80d089ec4b9 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_linalg.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_linalg.cc @@ -893,7 +893,7 @@ struct LhloLegalizeToLinalgPass auto func = getFunction(); populateLHLOToLinalgConversionPattern(func.getContext(), &patterns); - if (failed(applyPartialConversion(func, target, patterns, nullptr))) { + if (failed(applyPartialConversion(func, target, std::move(patterns)))) { signalPassFailure(); } } @@ -912,7 +912,7 @@ struct HloLegalizeToLinalgPass auto func = getFunction(); mhlo::populateHLOToLinalgConversionPattern(func.getContext(), &patterns); - if (failed(applyPartialConversion(func, target, patterns, nullptr))) { + if (failed(applyPartialConversion(func, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_standard.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_standard.cc index 84255c2810e..0173bcdd643 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_standard.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_to_standard.cc @@ -21,8 +21,8 @@ limitations under the License. #include "mlir-hlo/Dialect/mhlo/transforms/rewriters.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/Function.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/Pass/Pass.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { namespace { @@ -201,7 +201,7 @@ void PopulateMhloToStdPatterns(OwningRewritePatternList *patterns, void LegalizeToStandardPass::runOnFunction() { OwningRewritePatternList patterns; mlir::mhlo::PopulateMhloToStdPatterns(&patterns, &getContext()); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } } // end namespace mhlo diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_trigonometric_to_approximation.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_trigonometric_to_approximation.cc index 10030866d0f..b9f513f30d1 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_trigonometric_to_approximation.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_trigonometric_to_approximation.cc @@ -20,8 +20,8 @@ limitations under the License. #include "mlir-hlo/Dialect/mhlo/transforms/rewriters.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/Function.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/Pass/Pass.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { namespace mhlo { @@ -259,7 +259,7 @@ struct LegalizeTrigonometricToApproximationPass void runOnFunction() override { OwningRewritePatternList patterns; PopulateTrigonometricToApproximationPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_fuse_linalg.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_fuse_linalg.cc index 232948775c0..da760935391 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_fuse_linalg.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_fuse_linalg.cc @@ -27,6 +27,7 @@ limitations under the License. #include "mlir/Interfaces/ViewLikeInterface.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/FoldUtils.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { namespace lmhlo { @@ -111,7 +112,7 @@ class LhloFuseLinalgPass } }); auto patterns = linalg::getLinalgTilingCanonicalizationPatterns(ctx); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); // Fuse producers of tiled linalg ops. llvm::SmallDenseSet erase_set; @@ -132,7 +133,7 @@ class LhloFuseLinalgPass } auto patterns = linalg::getLinalgTilingCanonicalizationPatterns(ctx); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } for (auto* e : erase_set) e->erase(); } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_affine.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_affine.cc index 2041d22c62b..5bd27f0dbdc 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_affine.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_affine.cc @@ -19,12 +19,10 @@ limitations under the License. #include "mlir-hlo/Dialect/mhlo/transforms/map_lmhlo_to_scalar_op.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" -#include "mlir/IR/Attributes.h" #include "mlir/IR/Location.h" -#include "mlir/IR/MLIRContext.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/IR/StandardTypes.h" #include "mlir/Pass/Pass.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { namespace lmhlo { @@ -160,7 +158,7 @@ struct LhloLegalizeToAffinePass OwningRewritePatternList patterns; auto func = getFunction(); populateLHLOToAffineConversionPattern(func.getContext(), &patterns); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_gpu.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_gpu.cc index fbade8f7387..ea3eea0af61 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_gpu.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_gpu.cc @@ -184,7 +184,7 @@ struct LhloLegalizeToGpuPass target.addIllegalOp(); auto func = getFunction(); patterns.insert(func.getContext()); - if (failed(applyPartialConversion(func, target, patterns))) { + if (failed(applyPartialConversion(func, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_llvm_pass.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_llvm_pass.cc index 3d49027bb50..6b9286a6e89 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_llvm_pass.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_llvm_pass.cc @@ -47,7 +47,7 @@ class TestLhloToLLVMPass target.addLegalOp(); target.addIllegalDialect(); - if (failed(applyFullConversion(m, target, patterns))) { + if (failed(applyFullConversion(m, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_parallel_loops.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_parallel_loops.cc index d9a2d993496..78d681bdd06 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_parallel_loops.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lhlo_legalize_to_parallel_loops.cc @@ -713,7 +713,7 @@ struct LhloLegalizeToParallelLoopsPass target.addIllegalOp(); - if (failed(applyPartialConversion(func, target, patterns))) { + if (failed(applyPartialConversion(func, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_complex.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_complex.cc index 491f1c01cf7..7bea88fe5c9 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_complex.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_complex.cc @@ -29,11 +29,11 @@ limitations under the License. #include "mlir/IR/Attributes.h" #include "mlir/IR/MLIRContext.h" #include "mlir/IR/Operation.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/IR/TypeUtilities.h" #include "mlir/IR/Types.h" #include "mlir/Pass/Pass.h" #include "mlir/Pass/PassRegistry.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" using mlir::FunctionPass; using mlir::OwningRewritePatternList; @@ -70,7 +70,7 @@ void LowerComplexPass::runOnFunction() { OwningRewritePatternList patterns; mlir::mhlo::PopulateComplexLoweringPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } std::unique_ptr mlir::mhlo::createLowerComplexPass() { diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_general_dot.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_general_dot.cc index ada30a289a4..bc7007470f4 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_general_dot.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/lower_general_dot.cc @@ -25,10 +25,10 @@ limitations under the License. #include "mlir/IR/Function.h" #include "mlir/IR/Location.h" #include "mlir/IR/Operation.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/IR/StandardTypes.h" #include "mlir/IR/TypeUtilities.h" #include "mlir/Pass/Pass.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" using mlir::DenseIntElementsAttr; using mlir::ElementsAttr; @@ -182,7 +182,7 @@ struct LegalizeGeneralDotPass void runOnFunction() override { OwningRewritePatternList patterns; mlir::mhlo::PopulateGeneralDotOpLoweringPatterns(&patterns, &getContext()); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/materialize_broadcasts_pass.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/materialize_broadcasts_pass.cc index 3909f046007..d410a26ced0 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/materialize_broadcasts_pass.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/materialize_broadcasts_pass.cc @@ -42,7 +42,7 @@ struct TestMaterializeBroadcastsPass PopulateMaterializeBroadcastsPatterns(&getContext(), &conversionPatterns); if (failed(applyPartialConversion(getFunction(), conversionTarget, - conversionPatterns))) { + std::move(conversionPatterns)))) { return signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/optimize_mhlo_pass.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/optimize_mhlo_pass.cc index febd4423bf2..62fde5ec337 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/optimize_mhlo_pass.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/optimize_mhlo_pass.cc @@ -19,9 +19,9 @@ limitations under the License. #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/MLIRContext.h" #include "mlir/IR/Operation.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" using mlir::FunctionPass; using mlir::PassWrapper; @@ -42,7 +42,7 @@ void OptimizeMhloPass::runOnFunction() { mlir::OwningRewritePatternList patterns; mlir::mhlo::PopulateOptimizeMHLOPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } std::unique_ptr mlir::mhlo::createOptimizeMhloPass() { diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/test_infer_shaped_type_pass.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/test_infer_shaped_type_pass.cc index 35e5a184472..fed7f70c102 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/test_infer_shaped_type_pass.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/test_infer_shaped_type_pass.cc @@ -17,9 +17,9 @@ limitations under the License. #include "mlir/IR/Identifier.h" #include "mlir/IR/MLIRContext.h" #include "mlir/IR/OperationSupport.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/Interfaces/InferTypeOpInterface.h" #include "mlir/Pass/Pass.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { namespace mhlo { @@ -87,7 +87,7 @@ struct TestInferShapedTypeMethodsPass OwningRewritePatternList patterns; patterns.insert(&getContext()); patterns.insert(&getContext()); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/transform_unranked_hlo.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/transform_unranked_hlo.cc index 7c01fa22372..0dc9f1be929 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/transform_unranked_hlo.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/transform_unranked_hlo.cc @@ -154,7 +154,8 @@ struct TransformUnrankedHloPass PopulateTransformUnrankedHloPatterns(&ctx, &patterns); // Apply transformation. - if (failed(applyPartialConversion(getFunction(), target, patterns))) + if (failed( + applyPartialConversion(getFunction(), target, std::move(patterns)))) return signalPassFailure(); } }; diff --git a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/unfuse_batch_norm_pass.cc b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/unfuse_batch_norm_pass.cc index f187a7470cf..19261d49d1f 100644 --- a/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/unfuse_batch_norm_pass.cc +++ b/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/unfuse_batch_norm_pass.cc @@ -18,9 +18,9 @@ limitations under the License. #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/MLIRContext.h" #include "mlir/IR/Operation.h" -#include "mlir/IR/PatternMatch.h" #include "mlir/Pass/Pass.h" #include "mlir/Transforms/DialectConversion.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { namespace mhlo { @@ -32,7 +32,7 @@ struct TestUnfuseBatchNormPass void runOnOperation() override { OwningRewritePatternList patterns; PopulateUnfuseBatchNormPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getOperation(), patterns); + applyPatternsAndFoldGreedily(getOperation(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/lite/BUILD b/tensorflow/compiler/mlir/lite/BUILD index 4376eb55045..cbe15e83762 100644 --- a/tensorflow/compiler/mlir/lite/BUILD +++ b/tensorflow/compiler/mlir/lite/BUILD @@ -468,6 +468,7 @@ cc_library( "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) @@ -501,6 +502,7 @@ cc_library( "@llvm-project//mlir:QuantOps", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) diff --git a/tensorflow/compiler/mlir/lite/quantization/tensorflow/BUILD b/tensorflow/compiler/mlir/lite/quantization/tensorflow/BUILD index 76fd75e18ea..52beceaf084 100644 --- a/tensorflow/compiler/mlir/lite/quantization/tensorflow/BUILD +++ b/tensorflow/compiler/mlir/lite/quantization/tensorflow/BUILD @@ -25,14 +25,12 @@ cc_library( "passes.h", ], deps = [ - "//tensorflow/compiler/mlir/lite/quantization:quantization_config", "//tensorflow/compiler/mlir/lite/quantization:quantization_lib", "//tensorflow/compiler/mlir/tensorflow", - "@com_google_absl//absl/strings", - "@llvm-project//llvm:Support", "@llvm-project//mlir:IR", "@llvm-project//mlir:Pass", "@llvm-project//mlir:QuantOps", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) diff --git a/tensorflow/compiler/mlir/lite/quantization/tensorflow/tf_to_quant.cc b/tensorflow/compiler/mlir/lite/quantization/tensorflow/tf_to_quant.cc index b043834188c..e9935499f38 100644 --- a/tensorflow/compiler/mlir/lite/quantization/tensorflow/tf_to_quant.cc +++ b/tensorflow/compiler/mlir/lite/quantization/tensorflow/tf_to_quant.cc @@ -13,8 +13,8 @@ See the License for the specific language governing permissions and limitations under the License. ==============================================================================*/ #include "mlir/Dialect/Quant/QuantOps.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/lite/quantization/quantization_utils.h" #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" @@ -145,7 +145,7 @@ void LegalizeTFToQuant::runOnFunction() { auto func = getFunction(); auto *ctx = func.getContext(); patterns.insert(ctx); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/lite/transforms/dilated_conv.cc b/tensorflow/compiler/mlir/lite/transforms/dilated_conv.cc index b85b3de989a..d98d963c91d 100644 --- a/tensorflow/compiler/mlir/lite/transforms/dilated_conv.cc +++ b/tensorflow/compiler/mlir/lite/transforms/dilated_conv.cc @@ -14,6 +14,8 @@ limitations under the License. ==============================================================================*/ #include "tensorflow/compiler/mlir/lite/transforms/dilated_conv.h" +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project + namespace mlir { namespace TFL { namespace { @@ -30,7 +32,7 @@ void IdentifyDilatedConvPass::runOnFunction() { patterns.insert, ConvertTFDilatedConvOp>( &getContext()); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/lite/transforms/legalize_tf.cc b/tensorflow/compiler/mlir/lite/transforms/legalize_tf.cc index ca889e776d5..16a3db6c503 100644 --- a/tensorflow/compiler/mlir/lite/transforms/legalize_tf.cc +++ b/tensorflow/compiler/mlir/lite/transforms/legalize_tf.cc @@ -677,6 +677,7 @@ void LegalizeTF::runOnFunction() { // Ophint python converter converted tf node pattern. patterns.insert(context); + FrozenRewritePatternList frozenPatterns(std::move(patterns)); ConversionTarget target(*context); // It is legal to have TF ops in the graph still which can be @@ -716,7 +717,7 @@ void LegalizeTF::runOnFunction() { // Currently unit-test doesn't do multiple tries, so we need this. const int max_iterations = 15; for (int i = 0; i < max_iterations; ++i) { - if (failed(applyPartialConversion(func, target, patterns))) { + if (failed(applyPartialConversion(func, target, frozenPatterns))) { return; } } diff --git a/tensorflow/compiler/mlir/lite/transforms/lower_static_tensor_list.cc b/tensorflow/compiler/mlir/lite/transforms/lower_static_tensor_list.cc index c0a7ea9337b..04ee875689d 100644 --- a/tensorflow/compiler/mlir/lite/transforms/lower_static_tensor_list.cc +++ b/tensorflow/compiler/mlir/lite/transforms/lower_static_tensor_list.cc @@ -899,7 +899,7 @@ LogicalResult LowerStaticTensorListPass::RewriteFunction( ConvertTensorListSetItem, ConvertTensorListStack, ConvertTensorListResize, ConvertWhile, ConvertWhileRegion>( context); - return applyPartialConversion(func, target, patterns); + return applyPartialConversion(func, target, std::move(patterns)); } void LowerStaticTensorListPass::runOnOperation() { diff --git a/tensorflow/compiler/mlir/lite/transforms/optimize.cc b/tensorflow/compiler/mlir/lite/transforms/optimize.cc index 3c11fe2b610..e78dcc25ff3 100644 --- a/tensorflow/compiler/mlir/lite/transforms/optimize.cc +++ b/tensorflow/compiler/mlir/lite/transforms/optimize.cc @@ -35,13 +35,13 @@ limitations under the License. #include "mlir/Dialect/StandardOps/IR/Ops.h" // from @llvm-project #include "mlir/IR/Attributes.h" // from @llvm-project #include "mlir/IR/Matchers.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/StandardTypes.h" // from @llvm-project #include "mlir/IR/TypeUtilities.h" // from @llvm-project #include "mlir/IR/Value.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/lite/ir/tfl_ops.h" #include "tensorflow/compiler/mlir/lite/quantization/quantization_utils.h" #include "tensorflow/compiler/mlir/lite/transforms/passes.h" @@ -804,15 +804,21 @@ void Optimize::runOnFunction() { FuseFullyConnectedAndReluX, FuseFullyConnectedAndReluX, FuseFullyConnectedAndMul>(ctx); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); // Fuse the binary ops with the following ops. - patterns.insert< - FuseBinaryOpToFollowingConv2D, FuseBinaryOpToFollowingDepthwiseConv2D, + OwningRewritePatternList phase_2_patterns; + TFL::populateWithGenerated(ctx, phase_2_patterns); + phase_2_patterns.insert< + FuseFullyConnectedAndAdd, FuseFullyConnectedAndReluX, + FuseFullyConnectedAndReluX, + FuseFullyConnectedAndReluX, + FuseFullyConnectedAndMul, FuseBinaryOpToFollowingConv2D, + FuseBinaryOpToFollowingDepthwiseConv2D, FuseBinaryOpToFollowingFullyConnected, FuseConv2DAndMulWithQDQs, FuseDepthwiseConv2DAndMulWithQDQs, ConvertTrivialTransposeOpToReshapeOp>( ctx); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(phase_2_patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/lite/transforms/optimize_functional_ops.cc b/tensorflow/compiler/mlir/lite/transforms/optimize_functional_ops.cc index f1ea837446b..d0cedcc5051 100644 --- a/tensorflow/compiler/mlir/lite/transforms/optimize_functional_ops.cc +++ b/tensorflow/compiler/mlir/lite/transforms/optimize_functional_ops.cc @@ -21,11 +21,11 @@ limitations under the License. #include "mlir/IR/BlockAndValueMapping.h" // from @llvm-project #include "mlir/IR/MLIRContext.h" // from @llvm-project #include "mlir/IR/Module.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/StandardTypes.h" // from @llvm-project #include "mlir/IR/TypeUtilities.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" namespace mlir { @@ -150,7 +150,7 @@ void OptimizeFunctionalOpsPass::runOnOperation() { patterns.insert(&getContext()); ModuleOp module = getOperation(); - applyPatternsAndFoldGreedily(module, patterns); + applyPatternsAndFoldGreedily(module, std::move(patterns)); } PassRegistration pass( diff --git a/tensorflow/compiler/mlir/lite/transforms/post_quantize.cc b/tensorflow/compiler/mlir/lite/transforms/post_quantize.cc index ca30b2f1fcf..5afbfe18320 100644 --- a/tensorflow/compiler/mlir/lite/transforms/post_quantize.cc +++ b/tensorflow/compiler/mlir/lite/transforms/post_quantize.cc @@ -19,6 +19,7 @@ limitations under the License. #include "mlir/IR/MLIRContext.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/lite/ir/tfl_ops.h" #include "tensorflow/compiler/mlir/lite/quantization/quantization_utils.h" #include "tensorflow/compiler/mlir/lite/transforms/passes.h" @@ -146,14 +147,18 @@ void PostQuantizePass::runOnFunction() { auto* ctx = func.getContext(); TFL::populateWithGenerated(ctx, patterns); patterns.insert>(ctx); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); if (!emit_quant_adaptor_ops_) { RemoveQuantizationAdaptorOps(getFunction()); } - patterns.insert(ctx); - applyPatternsAndFoldGreedily(func, patterns); + OwningRewritePatternList phase_2_patterns; + TFL::populateWithGenerated(ctx, phase_2_patterns); + phase_2_patterns + .insert, RemoveVolatileOps>( + ctx); + applyPatternsAndFoldGreedily(func, std::move(phase_2_patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/lite/transforms/prepare_quantize.cc b/tensorflow/compiler/mlir/lite/transforms/prepare_quantize.cc index 783f21fce21..78fac7f4e11 100644 --- a/tensorflow/compiler/mlir/lite/transforms/prepare_quantize.cc +++ b/tensorflow/compiler/mlir/lite/transforms/prepare_quantize.cc @@ -25,9 +25,9 @@ limitations under the License. #include "mlir/Dialect/Quant/QuantOps.h" // from @llvm-project #include "mlir/IR/Function.h" // from @llvm-project #include "mlir/IR/MLIRContext.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/Value.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/lite/ir/tfl_ops.h" #include "tensorflow/compiler/mlir/lite/quantization/lite/tfl_to_std.h" #include "tensorflow/compiler/mlir/lite/quantization/quantization_config.h" @@ -337,7 +337,7 @@ void PrepareQuantizePass::runOnFunction() { // Currently, only activation stats are imported, so narrow_range = false. patterns.insert(bit_width, false, false, ctx); } - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); SanityCheckAndAdjustment(func); diff --git a/tensorflow/compiler/mlir/lite/transforms/prepare_tf.cc b/tensorflow/compiler/mlir/lite/transforms/prepare_tf.cc index c4f30c22be3..ef127659504 100644 --- a/tensorflow/compiler/mlir/lite/transforms/prepare_tf.cc +++ b/tensorflow/compiler/mlir/lite/transforms/prepare_tf.cc @@ -46,12 +46,12 @@ limitations under the License. #include "mlir/IR/Attributes.h" // from @llvm-project #include "mlir/IR/Function.h" // from @llvm-project #include "mlir/IR/MLIRContext.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/StandardTypes.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project #include "mlir/Transforms/DialectConversion.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.h" #include "tensorflow/compiler/mlir/lite/ir/tfl_ops.h" #include "tensorflow/compiler/mlir/lite/quantization/quantization_utils.h" @@ -1045,7 +1045,7 @@ LogicalResult ConvertTf2XlaOps(FuncOp func, MLIRContext *context) { TF::PopulateLegalizeHloToTfPatterns(&patterns, context); mhlo::GatherOp::getCanonicalizationPatterns(patterns, context); - return applyPartialConversion(func, target, patterns); + return applyPartialConversion(func, target, std::move(patterns)); } // Convert rfft to rfft2d. @@ -1145,7 +1145,7 @@ struct ConvertRfftToRfft2d : public RewritePattern { }; void PrepareTFPass::runOnFunction() { - OwningRewritePatternList patterns; + OwningRewritePatternList patterns, phase_2_patterns; auto func = getFunction(); MLIRContext *ctx = &getContext(); @@ -1183,20 +1183,20 @@ void PrepareTFPass::runOnFunction() { // This will allow optimizing any TF_Mul->TF_Conv in the graph // and any expanded from FusedBatchNorm. We need to do this // before converting TF_Conv to TFL_Conv - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); // Load the generated pattern again, so new quantization pass-through // will be applied. - patterns.clear(); - TFL::populateWithGenerated(ctx, patterns); + TFL::populateWithGenerated(ctx, phase_2_patterns); if (unfold_batch_matmul_) { - patterns.insert, - TF::ConvertTFBatchMatMulOp>(ctx); + phase_2_patterns.insert, + TF::ConvertTFBatchMatMulOp>( + ctx); } - patterns.insert(ctx); - applyPatternsAndFoldGreedily(func, patterns); + phase_2_patterns.insert(ctx); + applyPatternsAndFoldGreedily(func, std::move(phase_2_patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/lite/transforms/quantize.cc b/tensorflow/compiler/mlir/lite/transforms/quantize.cc index 529e57780c3..65ecaacaea5 100644 --- a/tensorflow/compiler/mlir/lite/transforms/quantize.cc +++ b/tensorflow/compiler/mlir/lite/transforms/quantize.cc @@ -29,6 +29,7 @@ limitations under the License. #include "mlir/IR/OperationSupport.h" // from @llvm-project #include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/lite/ir/tfl_ops.h" #include "tensorflow/compiler/mlir/lite/quantization/quantization_utils.h" #include "tensorflow/compiler/mlir/lite/transforms/passes.h" @@ -87,7 +88,7 @@ void QuantizePass::runOnFunction() { TFL::populateWithGenerated(ctx, patterns); patterns.insert( ctx, enable_numeric_verify, error_tolerance, enable_single_layer_verify); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tensorflow/BUILD b/tensorflow/compiler/mlir/tensorflow/BUILD index 77615f65d9e..0b323af1563 100644 --- a/tensorflow/compiler/mlir/tensorflow/BUILD +++ b/tensorflow/compiler/mlir/tensorflow/BUILD @@ -740,6 +740,7 @@ cc_library( "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TransformUtils", ], ) @@ -998,6 +999,7 @@ cc_library( "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/batchmatmul_to_einsum.cc b/tensorflow/compiler/mlir/tensorflow/transforms/batchmatmul_to_einsum.cc index fe0c5bea44e..2a607065cdb 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/batchmatmul_to_einsum.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/batchmatmul_to_einsum.cc @@ -32,6 +32,7 @@ limitations under the License. #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/core/util/matmul_bcast.h" @@ -93,7 +94,7 @@ void BatchMatMulToEinsumPass::runOnFunction() { patterns.insert, ConvertTFBatchMatMulToEinsumOp>( &getContext()); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } PassRegistration pass( diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/contraction_fusion.cc b/tensorflow/compiler/mlir/tensorflow/transforms/contraction_fusion.cc index b5d09f7a794..09712f9fca2 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/contraction_fusion.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/contraction_fusion.cc @@ -17,10 +17,10 @@ limitations under the License. #include "llvm/ADT/SmallVector.h" #include "mlir/IR/Attributes.h" // from @llvm-project #include "mlir/IR/MLIRContext.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/UseDefLists.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" namespace mlir { @@ -144,7 +144,7 @@ void ContractionFusionPass::runOnFunction() { OwningRewritePatternList patterns; patterns.insert(); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/decompose_resource_ops_pass.cc b/tensorflow/compiler/mlir/tensorflow/transforms/decompose_resource_ops_pass.cc index dcd0b9af5e1..2856cf65dbd 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/decompose_resource_ops_pass.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/decompose_resource_ops_pass.cc @@ -15,6 +15,7 @@ limitations under the License. #include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/transforms/decompose_resource_ops.h" #include "tensorflow/compiler/mlir/tensorflow/transforms/passes.h" @@ -45,7 +46,7 @@ struct DecomposeResourceOps OwningRewritePatternList patterns; mlir::TF::PopulateDecomposeResourceOpsPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/einsum.cc b/tensorflow/compiler/mlir/tensorflow/transforms/einsum.cc index c3d43c27ac5..1fea0fd3c7d 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/einsum.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/einsum.cc @@ -39,6 +39,7 @@ limitations under the License. #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/core/util/matmul_bcast.h" @@ -364,7 +365,7 @@ void TransformEinsumPass::runOnFunction() { auto func = getFunction(); patterns.insert(&getContext()); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } static PassRegistration pass( diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/fold_broadcast.cc b/tensorflow/compiler/mlir/tensorflow/transforms/fold_broadcast.cc index 66311101cee..c2760000e82 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/fold_broadcast.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/fold_broadcast.cc @@ -22,10 +22,10 @@ limitations under the License. #include "mlir/Dialect/Traits.h" // from @llvm-project #include "mlir/IR/Function.h" // from @llvm-project #include "mlir/IR/Operation.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/StandardTypes.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" namespace mlir { @@ -100,7 +100,7 @@ void BroadcastFoldPass::runOnFunction() { auto func = getFunction(); patterns.insert(); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/fused_kernel_matcher.cc b/tensorflow/compiler/mlir/tensorflow/transforms/fused_kernel_matcher.cc index 21f4581f76a..01ae4c3dd1e 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/fused_kernel_matcher.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/fused_kernel_matcher.cc @@ -24,6 +24,7 @@ limitations under the License. #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" namespace mlir { @@ -206,7 +207,7 @@ void FusedKernelMatcherPass::runOnFunction() { auto func = getFunction(); patterns.insert(&getContext()); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/gpu_fusion.cc b/tensorflow/compiler/mlir/tensorflow/transforms/gpu_fusion.cc index fbe0524ce8b..acc05f3a434 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/gpu_fusion.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/gpu_fusion.cc @@ -21,6 +21,7 @@ limitations under the License. #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Pass/PassManager.h" // from @llvm-project #include "mlir/Pass/PassRegistry.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "mlir/Transforms/Passes.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/compiler/mlir/tensorflow/transforms/passes.h" @@ -118,7 +119,7 @@ void GpuOpFusionPass::runOnFunction() { FuncOp func = getFunction(); OwningRewritePatternList patterns; patterns.insert(&getContext()); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/init_text_file_to_import.cc b/tensorflow/compiler/mlir/tensorflow/transforms/init_text_file_to_import.cc index 615ca26012e..49b88492be9 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/init_text_file_to_import.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/init_text_file_to_import.cc @@ -22,6 +22,7 @@ limitations under the License. #include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/FileUtilities.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" namespace mlir { @@ -115,7 +116,7 @@ void InitTextFileToImportPass::runOnFunction() { FuncOp func = getFunction(); patterns.insert(context); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/legalize_hlo.cc b/tensorflow/compiler/mlir/tensorflow/transforms/legalize_hlo.cc index 8ab348c1e5b..4e353252c4e 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/legalize_hlo.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/legalize_hlo.cc @@ -808,7 +808,8 @@ void LegalizeHloToTf::runOnFunction() { ConversionTarget target(context); target.addLegalDialect(); target.addLegalOp(); - if (failed(applyPartialConversion(getFunction(), target, patterns))) { + if (failed( + applyPartialConversion(getFunction(), target, std::move(patterns)))) { getFunction().emitError("mhlo to TF legalization failed."); signalPassFailure(); } diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/lower_tf_pass.cc b/tensorflow/compiler/mlir/tensorflow/transforms/lower_tf_pass.cc index 340b965cdd7..78be2c9a419 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/lower_tf_pass.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/lower_tf_pass.cc @@ -15,6 +15,7 @@ limitations under the License. #include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/transforms/lower_tf.h" namespace mlir { @@ -29,7 +30,7 @@ struct LowerTF : public PassWrapper { OwningRewritePatternList patterns; mlir::TF::PopulateLoweringTFPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/optimize.cc b/tensorflow/compiler/mlir/tensorflow/transforms/optimize.cc index bbb6cb44096..805a345ebe1 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/optimize.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/optimize.cc @@ -21,6 +21,7 @@ limitations under the License. #include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Pass/PassManager.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "mlir/Transforms/Passes.h" // from @llvm-project #include "tensorflow/compiler/mlir/lite/utils/validators.h" #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" @@ -129,7 +130,7 @@ struct TFOptimizePass : public PassWrapper { auto func = getFunction(); populateWithGenerated(&getContext(), patterns); patterns.insert(&getContext()); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/tensor_device_copy_conversion.cc b/tensorflow/compiler/mlir/tensorflow/transforms/tensor_device_copy_conversion.cc index f14efeb91ce..397a6f56fa0 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/tensor_device_copy_conversion.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/tensor_device_copy_conversion.cc @@ -59,7 +59,8 @@ class TensorDeviceCopyConversionPass patterns.insert, PassThroughConversion>(&getContext()); - if (failed(applyPartialConversion(getFunction(), target, patterns))) { + if (failed(applyPartialConversion(getFunction(), target, + std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/tf_data_optimization_pass.cc b/tensorflow/compiler/mlir/tensorflow/transforms/tf_data_optimization_pass.cc index 5be69bddb11..ae74c73d22f 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/tf_data_optimization_pass.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/tf_data_optimization_pass.cc @@ -15,6 +15,7 @@ limitations under the License. #include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/transforms/tf_data_optimization.h" namespace mlir { @@ -28,7 +29,7 @@ struct TFDataOptimization OwningRewritePatternList patterns; mlir::TF::PopulateTFDataOptimizationPatterns(&getContext(), &patterns); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/tensorflow/transforms/unroll_batch_matmul.cc b/tensorflow/compiler/mlir/tensorflow/transforms/unroll_batch_matmul.cc index ceb2d86899b..330e76884d6 100644 --- a/tensorflow/compiler/mlir/tensorflow/transforms/unroll_batch_matmul.cc +++ b/tensorflow/compiler/mlir/tensorflow/transforms/unroll_batch_matmul.cc @@ -33,6 +33,7 @@ limitations under the License. #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/core/util/matmul_bcast.h" @@ -54,7 +55,7 @@ void UnrollBatchMatMulPass::runOnFunction() { patterns.insert, ConvertTFBatchMatMulOp>(&getContext()); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tfjs/BUILD b/tensorflow/compiler/mlir/tfjs/BUILD index 34686cc0f68..68fc5c3be6a 100644 --- a/tensorflow/compiler/mlir/tfjs/BUILD +++ b/tensorflow/compiler/mlir/tfjs/BUILD @@ -111,6 +111,7 @@ cc_library( "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) diff --git a/tensorflow/compiler/mlir/tfjs/transforms/optimize.cc b/tensorflow/compiler/mlir/tfjs/transforms/optimize.cc index 5d3ee121577..04811ff8ede 100644 --- a/tensorflow/compiler/mlir/tfjs/transforms/optimize.cc +++ b/tensorflow/compiler/mlir/tfjs/transforms/optimize.cc @@ -25,6 +25,7 @@ limitations under the License. #include "mlir/IR/StandardTypes.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LLVM.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/compiler/mlir/tfjs/ir/tfjs_ops.h" @@ -51,7 +52,7 @@ void Optimize::runOnFunction() { auto func = getFunction(); populateWithGenerated(ctx, patterns); - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tfr/passes/decompose.cc b/tensorflow/compiler/mlir/tfr/passes/decompose.cc index e1a9ae8c2e6..796ad688a5b 100644 --- a/tensorflow/compiler/mlir/tfr/passes/decompose.cc +++ b/tensorflow/compiler/mlir/tfr/passes/decompose.cc @@ -36,7 +36,6 @@ limitations under the License. #include "mlir/IR/Builders.h" // from @llvm-project #include "mlir/IR/MLIRContext.h" // from @llvm-project #include "mlir/IR/Module.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/StandardTypes.h" // from @llvm-project #include "mlir/IR/SymbolTable.h" // from @llvm-project #include "mlir/IR/Value.h" // from @llvm-project @@ -45,6 +44,7 @@ limitations under the License. #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project #include "mlir/Transforms/DialectConversion.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "mlir/Transforms/InliningUtils.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/compiler/mlir/tfr/ir/tfr_ops.h" @@ -107,7 +107,7 @@ void DecomposeTFOpsPass::ApplyCanonicalization() { } populateSCFOpsCanonicalizationPatterns(patterns, context); - applyPatternsAndFoldGreedily(getFunction(), patterns); + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)); } LogicalResult DecomposeTFOpsPass::RewriteUnregisteredTFOps() { diff --git a/tensorflow/compiler/mlir/tfr/passes/raise_to_tf.cc b/tensorflow/compiler/mlir/tfr/passes/raise_to_tf.cc index f3fe9618c62..6448cafcb6a 100644 --- a/tensorflow/compiler/mlir/tfr/passes/raise_to_tf.cc +++ b/tensorflow/compiler/mlir/tfr/passes/raise_to_tf.cc @@ -40,7 +40,6 @@ limitations under the License. #include "mlir/IR/Matchers.h" // from @llvm-project #include "mlir/IR/Module.h" // from @llvm-project #include "mlir/IR/OperationSupport.h" // from @llvm-project -#include "mlir/IR/PatternMatch.h" // from @llvm-project #include "mlir/IR/StandardTypes.h" // from @llvm-project #include "mlir/IR/SymbolTable.h" // from @llvm-project #include "mlir/IR/Types.h" // from @llvm-project @@ -50,6 +49,7 @@ limitations under the License. #include "mlir/Support/LLVM.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project #include "mlir/Transforms/DialectConversion.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "mlir/Transforms/InliningUtils.h" // from @llvm-project #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/compiler/mlir/tfr/ir/tfr_ops.h" @@ -455,7 +455,7 @@ void RaiseToTFOpsPass::runOnFunction() { op->getCanonicalizationPatterns(patterns, ctx); } - applyPatternsAndFoldGreedily(func, patterns); + applyPatternsAndFoldGreedily(func, std::move(patterns)); } } // namespace diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/kernel_creator.cc b/tensorflow/compiler/mlir/tools/kernel_gen/kernel_creator.cc index 5a0fa9e2296..c5ba1b3ac8b 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/kernel_creator.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/kernel_creator.cc @@ -44,6 +44,7 @@ limitations under the License. #include "mlir/Pass/PassManager.h" // from @llvm-project #include "mlir/Transforms/Bufferize.h" // from @llvm-project #include "mlir/Transforms/DialectConversion.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "mlir/Transforms/Passes.h" // from @llvm-project #include "tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/transforms/passes.h" #include "tensorflow/compiler/mlir/tensorflow/dialect_registration.h" diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/bufferize_pass.cc b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/bufferize_pass.cc index bba32fc56dc..33276fd37af 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/bufferize_pass.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/bufferize_pass.cc @@ -110,7 +110,7 @@ struct BufferizePass : public BufferizePassBase { patterns.insert(&context); auto module = getOperation(); - if (failed(applyPartialConversion(module, target, patterns))) { + if (failed(applyPartialConversion(module, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/embed_tf_framework_pass.cc b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/embed_tf_framework_pass.cc index 6aea4d9c619..098c3cf34bb 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/embed_tf_framework_pass.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/embed_tf_framework_pass.cc @@ -64,7 +64,7 @@ class EmbedTFFrameworkPass return !op->getParentOfType().getAttrOfType(kTFEntry); }); - if (failed(applyPartialConversion(m, target, patterns))) { + if (failed(applyPartialConversion(m, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/materialize_broadcasts_pass.cc b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/materialize_broadcasts_pass.cc index dd3f32e2b3c..e0c21f0b2e4 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/materialize_broadcasts_pass.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/materialize_broadcasts_pass.cc @@ -44,7 +44,7 @@ struct MaterializeBroadcastsPass &conversionPatterns); if (failed(applyPartialConversion(getFunction(), conversionTarget, - conversionPatterns))) { + std::move(conversionPatterns)))) { return signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/parallel_loops_to_sequential.cc b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/parallel_loops_to_sequential.cc index 7981dbe5534..a428c618b7e 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/parallel_loops_to_sequential.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/parallel_loops_to_sequential.cc @@ -36,7 +36,8 @@ struct ParallelLoopsToSequentialPass target.addIllegalOp(); target.addLegalOp(); target.markUnknownOpDynamicallyLegal([](mlir::Operation*) { return true; }); - if (failed(applyPartialConversion(getOperation(), target, patterns))) + if (failed(applyPartialConversion(getOperation(), target, + std::move(patterns)))) signalPassFailure(); } }; diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/shape_to_descriptors_pass.cc b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/shape_to_descriptors_pass.cc index f5d01808c1b..d4a9baf17b9 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/shape_to_descriptors_pass.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/shape_to_descriptors_pass.cc @@ -61,7 +61,7 @@ struct ShapeToDescriptorsPass // Apply conversion. auto module = getOperation(); - if (failed(applyPartialConversion(module, target, patterns))) + if (failed(applyPartialConversion(module, target, std::move(patterns)))) signalPassFailure(); } }; diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/tf_kernel_to_llvm_pass.cc b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/tf_kernel_to_llvm_pass.cc index b2fcc424a50..10c8f6ccc3b 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/tf_kernel_to_llvm_pass.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/tf_kernel_to_llvm_pass.cc @@ -62,7 +62,7 @@ class TFKernelToLLVMPass : public TFKernelToLLVMPassBase { .addIllegalDialect(); target.addIllegalOp(); - if (failed(applyPartialConversion(m, target, patterns))) { + if (failed(applyPartialConversion(m, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/unfuse_batch_norm_pass.cc b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/unfuse_batch_norm_pass.cc index d2773d91b07..5c347f471b1 100644 --- a/tensorflow/compiler/mlir/tools/kernel_gen/transforms/unfuse_batch_norm_pass.cc +++ b/tensorflow/compiler/mlir/tools/kernel_gen/transforms/unfuse_batch_norm_pass.cc @@ -13,6 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. ==============================================================================*/ +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/transforms/rewriters.h" #include "tensorflow/compiler/mlir/tools/kernel_gen/transforms/passes.h" @@ -29,7 +30,7 @@ struct UnfuseBatchNormPass void runOnFunction() override { mlir::OwningRewritePatternList patterns; mlir::mhlo::PopulateUnfuseBatchNormPatterns(&getContext(), &patterns); - mlir::applyPatternsAndFoldGreedily(getOperation(), patterns); + mlir::applyPatternsAndFoldGreedily(getOperation(), std::move(patterns)); } }; diff --git a/tensorflow/compiler/mlir/xla/BUILD b/tensorflow/compiler/mlir/xla/BUILD index 1e664ff40f4..aa45d34543a 100644 --- a/tensorflow/compiler/mlir/xla/BUILD +++ b/tensorflow/compiler/mlir/xla/BUILD @@ -124,6 +124,7 @@ cc_library( "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TransformUtils", ], alwayslink = 1, ) diff --git a/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc b/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc index 1703dadeb4c..29d76f2b5a6 100644 --- a/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc +++ b/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc @@ -6125,8 +6125,8 @@ LogicalResult legalizeTF( // Fully qualify ReturnOp here as mhlo dialect also defines a ReturnOp. target.addLegalOp(); DenseSet nonlegalized_ops; - LogicalResult result = - applyPartialConversion(op, target, patterns, &nonlegalized_ops); + LogicalResult result = applyPartialConversion( + op, target, std::move(patterns), &nonlegalized_ops); // In order to enforce that the conversion result is fully converted, // fail if there are any nonlegalized ops in the set. if (failed(result) || !nonlegalized_ops.empty()) { @@ -6136,7 +6136,7 @@ LogicalResult legalizeTF( return result; } - return applyPartialConversion(op, target, patterns); + return applyPartialConversion(op, target, std::move(patterns)); } void PopulateLegalizeTfPatterns(MLIRContext *context, diff --git a/tensorflow/compiler/mlir/xla/transforms/legalize_tf_with_tf2xla.cc b/tensorflow/compiler/mlir/xla/transforms/legalize_tf_with_tf2xla.cc index 9bf4d546edf..82bac4ecb14 100644 --- a/tensorflow/compiler/mlir/xla/transforms/legalize_tf_with_tf2xla.cc +++ b/tensorflow/compiler/mlir/xla/transforms/legalize_tf_with_tf2xla.cc @@ -38,6 +38,7 @@ limitations under the License. #include "mlir/IR/Value.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Support/LogicalResult.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "tensorflow/compiler/mlir/op_or_arg_name_mapper.h" #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops_a_m.h" @@ -573,7 +574,8 @@ class LegalizeTF : public PassWrapper { void runOnFunction() override { OwningRewritePatternList patterns; patterns.insert(device_type_); - if (failed(applyPatternsAndFoldGreedily(getFunction(), patterns))) + if (failed( + applyPatternsAndFoldGreedily(getFunction(), std::move(patterns)))) signalPassFailure(); } diff --git a/tensorflow/compiler/xla/service/mlir_gpu/kernel_lowering.cc b/tensorflow/compiler/xla/service/mlir_gpu/kernel_lowering.cc index 9194e44ba5f..4a84fead837 100644 --- a/tensorflow/compiler/xla/service/mlir_gpu/kernel_lowering.cc +++ b/tensorflow/compiler/xla/service/mlir_gpu/kernel_lowering.cc @@ -31,11 +31,11 @@ limitations under the License. #include "mlir/Dialect/Linalg/Passes.h" // from @llvm-project #include "mlir/Dialect/SCF/Passes.h" // from @llvm-project #include "mlir/Dialect/SCF/Transforms.h" // from @llvm-project -#include "mlir/IR/Builders.h" // from @llvm-project #include "mlir/IR/Dialect.h" // from @llvm-project #include "mlir/Pass/Pass.h" // from @llvm-project #include "mlir/Pass/PassManager.h" // from @llvm-project #include "mlir/Transforms/Bufferize.h" // from @llvm-project +#include "mlir/Transforms/GreedyPatternRewriteDriver.h" // from @llvm-project #include "mlir/Transforms/Passes.h" // from @llvm-project #include "tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.h" #include "tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/transforms/passes.h" @@ -170,7 +170,7 @@ class LowerToNVVMPass // TODO(csigg): Remove once we support replacing non-root ops. target.addLegalOp<::mlir::gpu::GPUModuleOp, ::mlir::gpu::ModuleEndOp, ::mlir::gpu::YieldOp>(); - if (failed(mlir::applyFullConversion(m, target, patterns))) { + if (failed(mlir::applyFullConversion(m, target, std::move(patterns)))) { signalPassFailure(); } } @@ -214,11 +214,13 @@ class LowerToROCDLPass void runOnOperation() override { ::mlir::gpu::GPUModuleOp m = getOperation(); - ::mlir::OwningRewritePatternList patterns; - ::mlir::populateGpuRewritePatterns(m.getContext(), patterns); - ::mlir::applyPatternsAndFoldGreedily(m, patterns); - patterns.clear(); + { + ::mlir::OwningRewritePatternList patterns; + ::mlir::populateGpuRewritePatterns(m.getContext(), patterns); + ::mlir::applyPatternsAndFoldGreedily(m, std::move(patterns)); + } + ::mlir::OwningRewritePatternList patterns; ::mlir::LLVMTypeConverter converter(m.getContext()); ::mlir::populateStdToLLVMConversionPatterns(converter, patterns); // TODO(b/145824979) Remove linalg once sliceop is in std. @@ -239,7 +241,7 @@ class LowerToROCDLPass // TODO(csigg): Remove once we support replacing non-root ops. target.addLegalOp<::mlir::gpu::GPUModuleOp, ::mlir::gpu::ModuleEndOp, ::mlir::gpu::YieldOp>(); - if (failed(mlir::applyFullConversion(m, target, patterns))) { + if (failed(mlir::applyFullConversion(m, target, std::move(patterns)))) { signalPassFailure(); } } diff --git a/tensorflow/workspace.bzl b/tensorflow/workspace.bzl index 952e786704a..410fe252c35 100755 --- a/tensorflow/workspace.bzl +++ b/tensorflow/workspace.bzl @@ -680,8 +680,8 @@ def tf_repositories(path_prefix = "", tf_repo_name = ""): ) # Check out LLVM and MLIR from llvm-project. - LLVM_COMMIT = "26750a1264b3df114a1efae7cde6f0784206b2ce" - LLVM_SHA256 = "eb360bbcd3e4b505689a21756c89fdf087064882899bab2f98258f8cf0546218" + LLVM_COMMIT = "0fc1aa22ee6ac337a5d51fa5666c9cd61da61b07" + LLVM_SHA256 = "e0682d8c3c29b99db0037913bae706bdfc6e2439d75132087d628e54d894f1e1" LLVM_URLS = [ "https://storage.googleapis.com/mirror.tensorflow.org/github.com/llvm/llvm-project/archive/{commit}.tar.gz".format(commit = LLVM_COMMIT), "https://github.com/llvm/llvm-project/archive/{commit}.tar.gz".format(commit = LLVM_COMMIT), diff --git a/third_party/mlir/BUILD b/third_party/mlir/BUILD index 04533ae5f24..c346a7f3f15 100644 --- a/third_party/mlir/BUILD +++ b/third_party/mlir/BUILD @@ -621,6 +621,7 @@ cc_library( ":Affine", ":IR", ":Support", + ":TransformUtils", "@llvm-project//llvm:Support", ], ) @@ -701,6 +702,7 @@ cc_library( ":LinalgToSPIRV", ":LinalgToStandard", ":OpenMPToLLVM", + ":PDLToPDLInterp", ":SCFToGPUPass", ":SCFToStandard", ":SPIRVToLLVM", @@ -1694,6 +1696,30 @@ cc_library( ], ) +cc_library( + name = "PDLToPDLInterp", + srcs = glob([ + "lib/Conversion/PDLToPDLInterp/*.cpp", + "lib/Conversion/PDLToPDLInterp/*.h", + ]) + [ + "lib/Conversion/PassDetail.h", + ], + hdrs = [ + "include/mlir/Conversion/PDLToPDLInterp/PDLToPDLInterp.h", + ], + includes = ["include"], + deps = [ + ":ConversionPassIncGen", + ":IR", + ":InferTypeOpInterface", + ":PDLDialect", + ":PDLInterpDialect", + ":Pass", + ":Support", + "@llvm-project//llvm:Support", + ], +) + cc_library( name = "SPIRVToLLVM", srcs = glob([ @@ -2348,6 +2374,22 @@ cc_library( ], ) +cc_library( + name = "Rewrite", + srcs = glob([ + "lib/Rewrite/*.cpp", + "lib/Rewrite/*.h", + ]), + hdrs = glob(["include/mlir/Rewrite/*.h"]), + includes = ["include"], + deps = [ + ":IR", + ":Pass", + ":Support", + "@llvm-project//llvm:Support", + ], +) + cc_library( name = "TransformUtils", srcs = glob([ @@ -2364,6 +2406,7 @@ cc_library( ":ControlFlowInterfaces", ":IR", ":Pass", + ":Rewrite", ":SCFDialect", ":SideEffectInterfaces", ":StandardOps", @@ -2523,6 +2566,7 @@ cc_library( ":LinalgOps", ":LoopLikeInterface", ":Pass", + ":Rewrite", ":SCFDialect", ":SideEffectInterfaces", ":StandardOps", @@ -3153,6 +3197,7 @@ cc_library( ":OpenMPToLLVM", ":PDLDialect", ":PDLInterpDialect", + ":PDLToPDLInterp", ":QuantOps", ":QuantPassIncGen", ":ROCDLDialect",