|
| 1 | +/* Copyright 2025 The TensorFlow Authors. All Rights Reserved. |
| 2 | +
|
| 3 | +Licensed under the Apache License, Version 2.0 (the "License"); |
| 4 | +you may not use this file except in compliance with the License. |
| 5 | +You may obtain a copy of the License at |
| 6 | +
|
| 7 | + http://www.apache.org/licenses/LICENSE-2.0 |
| 8 | +
|
| 9 | +Unless required by applicable law or agreed to in writing, software |
| 10 | +distributed under the License is distributed on an "AS IS" BASIS, |
| 11 | +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| 12 | +See the License for the specific language governing permissions and |
| 13 | +limitations under the License. |
| 14 | +==============================================================================*/ |
| 15 | + |
| 16 | +#include "tflite/converter/stablehlo/transforms/legalize_hlo_conversions/case.h" |
| 17 | + |
| 18 | +#include "mlir/Dialect/Arith/IR/Arith.h" // from @llvm-project |
| 19 | +#include "mlir/IR/BuiltinTypeInterfaces.h" // from @llvm-project |
| 20 | +#include "mlir/IR/PatternMatch.h" // from @llvm-project |
| 21 | +#include "mlir/Support/LLVM.h" // from @llvm-project |
| 22 | +#include "mlir/Transforms/DialectConversion.h" // from @llvm-project |
| 23 | +#include "tflite/converter/ir/tfl_ops.h" |
| 24 | +#include "tflite/converter/stablehlo/transforms/legalize_hlo_conversions/util.h" |
| 25 | +#include "xla/mlir_hlo/mhlo/IR/hlo_ops.h" |
| 26 | + |
| 27 | +namespace mlir::odml { |
| 28 | +namespace { |
| 29 | + |
| 30 | +// Legalizes mhlo.case op to tfl.if op. |
| 31 | +// This pattern only supports mhlo.case ops with exactly two branches. |
| 32 | +class LegalizeCaseOp : public OpConversionPattern<mhlo::CaseOp> { |
| 33 | + public: |
| 34 | + using OpConversionPattern<mhlo::CaseOp>::OpConversionPattern; |
| 35 | + |
| 36 | + LogicalResult matchAndRewrite( |
| 37 | + mhlo::CaseOp case_op, OpAdaptor adaptor, |
| 38 | + ConversionPatternRewriter& rewriter) const final { |
| 39 | + // mhlo.case can have N branches, but tfl.if only supports two. |
| 40 | + if (case_op.getBranches().size() != 2) { |
| 41 | + return rewriter.notifyMatchFailure( |
| 42 | + case_op, "can only convert mhlo.case with 2 branches"); |
| 43 | + } |
| 44 | + |
| 45 | + // `mhlo.case` takes an index, `tfl.if` takes a boolean predicate. |
| 46 | + // For a 2-branch `mhlo.case` (branch 0 and branch 1), we need to map |
| 47 | + // the index to a boolean. |
| 48 | + // According to the mhlo.case spec, an out-of-bounds index defaults to the |
| 49 | + // index of the last branch, which is 1 in this case. |
| 50 | + // So, index 0 maps to branch 0, and any other index (1, or out of bounds) |
| 51 | + // maps to branch 1. |
| 52 | + // This can be expressed as a predicate `index != 0` for branch 1. |
| 53 | + |
| 54 | + auto loc = case_op->getLoc(); |
| 55 | + auto index = case_op.getIndex(); |
| 56 | + auto index_type = mlir::cast<ShapedType>(index.getType()); |
| 57 | + |
| 58 | + // Create a constant tensor of the same shape as the index, filled with |
| 59 | + // zeros. |
| 60 | + auto const_zero = arith::ConstantOp::create( |
| 61 | + rewriter, loc, rewriter.getZeroAttr(index_type)); |
| 62 | + |
| 63 | + // Create the predicate `index != 0`. |
| 64 | + auto pred_type = index_type.clone(rewriter.getI1Type()); |
| 65 | + auto pred = mhlo::CompareOp::create( |
| 66 | + rewriter, loc, pred_type, index, const_zero, |
| 67 | + mhlo::ComparisonDirectionAttr::get(rewriter.getContext(), |
| 68 | + mhlo::ComparisonDirection::NE), |
| 69 | + mhlo::ComparisonTypeAttr{}); // Default comparison type is fine for |
| 70 | + // integers. |
| 71 | + |
| 72 | + // Create the tfl.if op. |
| 73 | + auto tfl_if = |
| 74 | + TFL::IfOp::create(rewriter, loc, case_op.getResultTypes(), pred); |
| 75 | + |
| 76 | + // Branch 1 of mhlo.case becomes the `then_region` of tfl.if. |
| 77 | + tfl_if.getThenRegion().takeBody(case_op.getBranches()[1]); |
| 78 | + ReplaceTerminatorWithYield(tfl_if.getThenRegion(), rewriter); |
| 79 | + |
| 80 | + // Branch 0 of mhlo.case becomes the `else_region` of tfl.if. |
| 81 | + tfl_if.getElseRegion().takeBody(case_op.getBranches()[0]); |
| 82 | + ReplaceTerminatorWithYield(tfl_if.getElseRegion(), rewriter); |
| 83 | + |
| 84 | + rewriter.replaceOp(case_op, tfl_if.getResults()); |
| 85 | + return success(); |
| 86 | + } |
| 87 | +}; |
| 88 | + |
| 89 | +} // namespace |
| 90 | + |
| 91 | +void PopulateCasePatterns(MLIRContext* context, RewritePatternSet& patterns, |
| 92 | + ConversionTarget& target) { |
| 93 | + patterns.add<LegalizeCaseOp>(context); |
| 94 | + // Mark mhlo.case as dynamically legal: it's legal if it does NOT have |
| 95 | + // exactly 2 branches, as those are the ones we want to convert. |
| 96 | + target.addDynamicallyLegalOp<mhlo::CaseOp>( |
| 97 | + [](mhlo::CaseOp op) { return op.getBranches().size() != 2; }); |
| 98 | +} |
| 99 | + |
| 100 | +} // namespace mlir::odml |
0 commit comments