Remove mhlo.conv padding preprocessing.
diff --git a/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp b/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp
index 571ea7f..b20d1b5 100644
--- a/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp
+++ b/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp
@@ -79,71 +79,6 @@
.getResult();
}
-class ExtractConvOpPaddingAttributes : public OpRewritePattern<mhlo::ConvOp> {
- public:
- using OpRewritePattern<mhlo::ConvOp>::OpRewritePattern;
-
- LogicalResult matchAndRewrite(mhlo::ConvOp op,
- PatternRewriter &rewriter) const override {
- if (!hasPadding(op)) return failure();
- auto inputType = op.lhs().getType().cast<ShapedType>();
- int rank = inputType.getRank();
-
- // TODO(suderman): Add proper support for padding + dilation for codegen.
- // We can't extract padding if the left hand side has dilation.
- if (op.lhs_dilation().hasValue()) {
- for (auto val : op.lhs_dilation().getValue().getValues<APInt>()) {
- if (val != 1) {
- return failure();
- }
- }
- }
-
- SmallVector<int64_t, 4> paddingLow, paddingHigh, interiorPadding, shape;
- paddingLow.append(rank, 0);
- paddingHigh.append(rank, 0);
- interiorPadding.append(rank, 0);
- for (auto iter :
- llvm::enumerate(op.dimension_numbers().getInputSpatialDimensions())) {
- unsigned idx = iter.index();
- unsigned dim = iter.value();
- paddingLow[dim] = op.paddingAttr().getValues<int64_t>()[{idx, 0}];
- paddingHigh[dim] = op.paddingAttr().getValues<int64_t>()[{idx, 1}];
- }
- for (unsigned i = 0; i < rank; ++i) {
- // mhlo.pad doesn't support dynamic shape.
- if (inputType.isDynamicDim(i)) return failure();
- int size = inputType.getShape()[i];
- shape.push_back(size + paddingLow[i] + paddingHigh[i]);
- }
-
- auto toDenseAttr = [&rewriter](ArrayRef<int64_t> elements) {
- return DenseIntElementsAttr::get(
- RankedTensorType::get(elements.size(), rewriter.getIntegerType(64)),
- elements);
- };
-
- auto loc = op.getLoc();
- auto padResultType =
- RankedTensorType::get(shape, inputType.getElementType());
- Attribute zeroAttr = rewriter.getZeroAttr(
- RankedTensorType::get({}, inputType.getElementType()));
- auto zero = rewriter.create<arith::ConstantOp>(loc, zeroAttr);
- auto padOp = rewriter.create<mhlo::PadOp>(
- loc, padResultType, op.lhs(), zero, toDenseAttr(paddingLow),
- toDenseAttr(paddingHigh), toDenseAttr(interiorPadding));
- auto resultType = op.getResult().getType();
- auto newOp = rewriter.create<mhlo::ConvOp>(
- op.getLoc(), resultType, padOp.getResult(), op.rhs(),
- op.window_stridesAttr(), /*padding=*/nullptr, op.lhs_dilationAttr(),
- op.rhs_dilationAttr(), /*window_reversal=*/nullptr,
- op.dimension_numbersAttr(), op.feature_group_countAttr(),
- op.batch_group_countAttr(), op.precision_configAttr());
- rewriter.replaceOp(op, newOp.getResult());
- return success();
- }
-};
-
// Guarantee that the input dimensions are ordered batch, spatial_dims, feature
// dim.
class ReorderConvOpInputDimensions : public OpRewritePattern<mhlo::ConvOp> {
@@ -987,9 +922,6 @@
ReorderBroadcastInDimOpAndElementwiseOp<mhlo::AndOp>,
ReorderBroadcastInDimOpAndElementwiseOp<mhlo::OrOp>,
ReorderBroadcastInDimOpAndElementwiseOp<mhlo::XorOp>>(context);
- if (extractPadFromConv) {
- patterns.insert<ExtractConvOpPaddingAttributes>(context);
- }
if (orderConvFeatures) {
patterns.insert<ReorderConvOpInputDimensions>(context);
patterns.insert<ReorderConvOpKernelDimensions>(context);
diff --git a/iree/compiler/InputConversion/MHLO/Passes.td b/iree/compiler/InputConversion/MHLO/Passes.td
index aa633a8..4cf16ae 100644
--- a/iree/compiler/InputConversion/MHLO/Passes.td
+++ b/iree/compiler/InputConversion/MHLO/Passes.td
@@ -40,8 +40,6 @@
let summary = "Apply mhlo to mhlo transformations for some mhlo ops";
let constructor = "mlir::iree_compiler::MHLO::createMHLOToMHLOPreprocessingPass()";
let options = [
- Option<"extractPadFromConv", "extract-pad-from-conv", "bool", /*default=*/"true",
- "Extract padding attributes from conv op">,
Option<"orderConvFeatures", "order-conv-features", "bool", /*default=*/"true",
"Guarantees input/output features ordered from conv kernel">
];
diff --git a/iree/compiler/InputConversion/MHLO/test/BUILD b/iree/compiler/InputConversion/MHLO/test/BUILD
index d0fbcbf..8e14a4d 100644
--- a/iree/compiler/InputConversion/MHLO/test/BUILD
+++ b/iree/compiler/InputConversion/MHLO/test/BUILD
@@ -29,7 +29,6 @@
"legalize_input_types.mlir",
"mhlo_to_mhlo_preprocessing.mlir",
"mhlo_to_mhlo_preprocessing_canoncalize_dot_general.mlir",
- "mhlo_to_mhlo_preprocessing_extract_pad_from_conv.mlir",
"missing_legalizations.mlir",
"verify_compiler_mhlo_input_legality.mlir",
],
diff --git a/iree/compiler/InputConversion/MHLO/test/CMakeLists.txt b/iree/compiler/InputConversion/MHLO/test/CMakeLists.txt
index f95c73d..96e60e4 100644
--- a/iree/compiler/InputConversion/MHLO/test/CMakeLists.txt
+++ b/iree/compiler/InputConversion/MHLO/test/CMakeLists.txt
@@ -24,7 +24,6 @@
"legalize_input_types.mlir"
"mhlo_to_mhlo_preprocessing.mlir"
"mhlo_to_mhlo_preprocessing_canoncalize_dot_general.mlir"
- "mhlo_to_mhlo_preprocessing_extract_pad_from_conv.mlir"
"missing_legalizations.mlir"
"verify_compiler_mhlo_input_legality.mlir"
TOOLS
diff --git a/iree/compiler/InputConversion/MHLO/test/mhlo_to_mhlo_preprocessing_extract_pad_from_conv.mlir b/iree/compiler/InputConversion/MHLO/test/mhlo_to_mhlo_preprocessing_extract_pad_from_conv.mlir
deleted file mode 100644
index 427c37f..0000000
--- a/iree/compiler/InputConversion/MHLO/test/mhlo_to_mhlo_preprocessing_extract_pad_from_conv.mlir
+++ /dev/null
@@ -1,29 +0,0 @@
-// RUN: iree-opt -iree-mhlo-to-mhlo-preprocessing %s | FileCheck %s
-
-// CHECK-LABEL: @conv
-// CHECK: mhlo.pad
-// CHECK-SAME: edge_padding_high = dense<[0, 1, 1, 0]>
-// CHECK-SAME: edge_padding_low = dense<[0, 1, 0, 0]>
-// CHECK: mhlo.convolution
-// CHECK-NOT: padding
-func @conv(%inputs: tensor<1x4x5x2xf32>, %weights: tensor<3x2x2x1xf32>) -> tensor<1x4x5x1xf32> {
- %0 = "mhlo.convolution"(%inputs, %weights) {
- batch_group_count = 1 : i64,
- dimension_numbers = #mhlo.conv<raw
- input_batch_dimension = 0,
- input_feature_dimension = 3,
- input_spatial_dimensions = [1, 2],
- kernel_input_feature_dimension = 2,
- kernel_output_feature_dimension = 3,
- kernel_spatial_dimensions = [0, 1],
- output_batch_dimension = 0,
- output_feature_dimension = 3,
- output_spatial_dimensions = [1, 2]
- >,
- feature_group_count = 1 : i64,
- padding = dense<[[1, 1], [0, 1]]> : tensor<2x2xi64>,
- rhs_dilation = dense<1> : tensor<2xi64>,
- window_strides = dense<1> : tensor<2xi64>} :
- (tensor<1x4x5x2xf32>, tensor<3x2x2x1xf32>) -> tensor<1x4x5x1xf32>
- return %0 : tensor<1x4x5x1xf32>
-}
