-
Notifications
You must be signed in to change notification settings - Fork 74k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
This change is to canonicalize boundary values in the IR replacing -I…
…nf/Inf with MIN/MAX float value. PiperOrigin-RevId: 630372176
- Loading branch information
1 parent
c662838
commit b59bae9
Showing
18 changed files
with
309 additions
and
13 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
79 changes: 79 additions & 0 deletions
79
tensorflow/compiler/mlir/lite/tests/canonicalize_boundary_value.mlir
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,79 @@ | ||
| // RUN: tf-opt %s --canonicalize-boundary-value --split-input-file | FileCheck %s | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_neg_inf_f32() -> tensor<f32> { | ||
| // CHECK: %[[CONST:.*]] = stablehlo.constant dense<-3.40282347E+38> : tensor<f32> | ||
| // CHECK: return %[[CONST]] : tensor<f32> | ||
|
|
||
| func.func @clamp_neg_inf_f32() -> tensor<f32> { | ||
| %ret = stablehlo.constant dense<0xFF800000> : tensor<f32> | ||
| return %ret : tensor<f32> | ||
| } | ||
|
|
||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_pos_inf_f32() -> tensor<f32> { | ||
| // CHECK: %[[CONST:.*]] = stablehlo.constant dense<3.40282347E+38> : tensor<f32> | ||
| // CHECK: return %[[CONST]] : tensor<f32> | ||
| func.func @clamp_pos_inf_f32() -> tensor<f32> { | ||
| %ret = stablehlo.constant dense<0x7F800000> : tensor<f32> | ||
| return %ret : tensor<f32> | ||
| } | ||
|
|
||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_pos_inf_f32_tensor() -> tensor<1x4xf32> { | ||
| // CHECK: %[[CONST:.*]] = stablehlo.constant dense<{{\[\[}}3.40282347E+38, 1.000000e+01, 2.000000e+01, -3.40282347E+38]]> : tensor<1x4xf32> | ||
| // CHECK: return %[[CONST]] : tensor<1x4xf32> | ||
| func.func @clamp_pos_inf_f32_tensor() -> tensor<1x4xf32> { | ||
| %ret = stablehlo.constant dense<[[0x7F800000, 10.0, 20.0, 0xFF800000]]> : tensor<1x4xf32> | ||
| return %ret : tensor<1x4xf32> | ||
| } | ||
|
|
||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_neg_inf_f16() -> tensor<f16> { | ||
| // CHECK: %[[CONST:.*]] = stablehlo.constant dense<-6.550400e+04> : tensor<f16> | ||
| // CHECK: return %[[CONST]] : tensor<f16> | ||
| func.func @clamp_neg_inf_f16() -> tensor<f16> { | ||
| %ret = stablehlo.constant dense<0xFC00> : tensor<f16> | ||
| return %ret : tensor<f16> | ||
| } | ||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_neg_inf_bf16() -> tensor<bf16> { | ||
| // CHECK: %[[CONST:.*]] = stablehlo.constant dense<-1.038460e+34> : tensor<bf16> | ||
| // CHECK: return %[[CONST]] : tensor<bf16> | ||
| func.func @clamp_neg_inf_bf16() -> tensor<bf16> { | ||
| %ret = stablehlo.constant dense<0xF800> : tensor<bf16> | ||
| return %ret : tensor<bf16> | ||
| } | ||
|
|
||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_pos_inf_f16() -> tensor<f16> { | ||
| // CHECK: %[[CONST:.*]] = stablehlo.constant dense<6.550400e+04> : tensor<f16> | ||
| // CHECK: return %[[CONST]] : tensor<f16> | ||
| func.func @clamp_pos_inf_f16() -> tensor<f16> { | ||
| %ret = stablehlo.constant dense<0x7C00> : tensor<f16> | ||
| return %ret : tensor<f16> | ||
| } | ||
|
|
||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_pos_inf_f16_tensor() -> tensor<1x4xf16> { | ||
| // CHECK: %[[CONST:.*]] = stablehlo.constant dense<{{\[\[}}6.550400e+04, 1.000000e+01, 2.000000e+01, -6.550400e+04]]> : tensor<1x4xf16> | ||
| // CHECK: return %[[CONST]] : tensor<1x4xf16> | ||
| func.func @clamp_pos_inf_f16_tensor() -> tensor<1x4xf16> { | ||
| %ret = stablehlo.constant dense<[[0x7C00, 10.0, 20.0, 0xFC00]]> : tensor<1x4xf16> | ||
| return %ret : tensor<1x4xf16> | ||
| } | ||
|
|
||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_pos_inf_f16_tensor_tf_const() -> tensor<3xf16> { | ||
| // CHECK: %[[CONST:.*]] = "tf.Const"() <{value = dense<6.550400e+04> : tensor<3xf16>}> : () -> tensor<3xf16> | ||
| // CHECK: return %[[CONST]] : tensor<3xf16> | ||
| func.func @clamp_pos_inf_f16_tensor_tf_const() -> tensor<3xf16> { | ||
| %ret = "tf.Const"() <{value = dense<0x7C00> : tensor<3xf16>}> : () -> tensor<3xf16> | ||
| return %ret : tensor<3xf16> | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
104 changes: 104 additions & 0 deletions
104
tensorflow/compiler/mlir/lite/transforms/canonicalize_boundary_value_pass.cc
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,104 @@ | ||
| /* Copyright 2024 The TensorFlow Authors. All Rights Reserved. | ||
| Licensed under the Apache License, Version 2.0 (the "License"); | ||
| you may not use this file except in compliance with the License. | ||
| You may obtain a copy of the License at | ||
| http://www.apache.org/licenses/LICENSE-2.0 | ||
| Unless required by applicable law or agreed to in writing, software | ||
| distributed under the License is distributed on an "AS IS" BASIS, | ||
| WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| See the License for the specific language governing permissions and | ||
| limitations under the License. | ||
| ==============================================================================*/ | ||
|
|
||
| #include <memory> | ||
| #include <utility> | ||
|
|
||
| #include "mlir/Dialect/Arith/IR/Arith.h" | ||
| #include "llvm/ADT/STLExtras.h" | ||
| #include "mlir/Dialect/Func/IR/FuncOps.h" // from @llvm-project | ||
| #include "mlir/IR/Builders.h" // from @llvm-project | ||
| #include "mlir/IR/BuiltinAttributeInterfaces.h" // from @llvm-project | ||
| #include "mlir/IR/BuiltinAttributes.h" // from @llvm-project | ||
| #include "mlir/IR/BuiltinOps.h" // from @llvm-project | ||
| #include "mlir/IR/BuiltinTypeInterfaces.h" // from @llvm-project | ||
| #include "mlir/IR/BuiltinTypes.h" // from @llvm-project | ||
| #include "mlir/IR/Diagnostics.h" // from @llvm-project | ||
| #include "mlir/IR/PatternMatch.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 "stablehlo/dialect/StablehloOps.h" // from @stablehlo | ||
| #include "tensorflow/compiler/mlir/lite/transforms/passes.h" | ||
| #include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h" | ||
|
|
||
| namespace mlir { | ||
| namespace TFL { | ||
| namespace { | ||
|
|
||
| #define DEBUG_TYPE "canonicalize-boundary-value" | ||
|
|
||
| #define GEN_PASS_DEF_CANONICALIZEBOUNDARYVALUEPASS | ||
| #include "tensorflow/compiler/mlir/lite/transforms/passes.h.inc" | ||
|
|
||
| class CanonicalizeBoundaryValuePass | ||
| : public impl::CanonicalizeBoundaryValuePassBase< | ||
| CanonicalizeBoundaryValuePass> { | ||
| void runOnOperation() override; | ||
| }; | ||
|
|
||
| // Clamp constant -Inf/Inf to MIN/MAX float value. | ||
| template <typename OpTy> | ||
| struct ClampInfToMinMaxFloat : public OpRewritePattern<OpTy> { | ||
| using OpRewritePattern<OpTy>::OpRewritePattern; | ||
|
|
||
| LogicalResult matchAndRewrite(OpTy const_op, | ||
| PatternRewriter& rewriter) const override { | ||
| ElementsAttr tensor_attr = | ||
| const_op.getValueAttr().template cast<ElementsAttr>(); | ||
| ShapedType tensor_type = tensor_attr.getShapedType(); | ||
| auto float_type = dyn_cast<mlir::FloatType>(tensor_type.getElementType()); | ||
| if (!float_type) return failure(); | ||
|
|
||
| auto vals_orig = tensor_attr.getValues<APFloat>(); | ||
| // If all values are finite, no need to rewrite. | ||
| if (llvm::all_of(vals_orig, [&](APFloat val) { return !val.isInfinity(); })) | ||
| return failure(); | ||
|
|
||
| SmallVector<APFloat> vals_new(llvm::map_range(vals_orig, [&](APFloat val) { | ||
| return val.isInfinity() | ||
| ? APFloat::getLargest(float_type.getFloatSemantics(), | ||
| val.isNegative()) | ||
| : val; | ||
| })); | ||
| rewriter.replaceOpWithNewOp<OpTy>( | ||
| const_op, DenseElementsAttr::get(tensor_type, vals_new)); | ||
| return success(); | ||
| } | ||
| }; | ||
|
|
||
| void CanonicalizeBoundaryValuePass::runOnOperation() { | ||
| auto* ctx = &getContext(); | ||
|
|
||
| RewritePatternSet patterns(ctx); | ||
| patterns.add<ClampInfToMinMaxFloat<stablehlo::ConstantOp>, | ||
| ClampInfToMinMaxFloat<TF::ConstOp>, | ||
| ClampInfToMinMaxFloat<arith::ConstantOp>>(ctx); | ||
| if (failed( | ||
| applyPatternsAndFoldGreedily(getOperation(), std::move(patterns)))) { | ||
| return signalPassFailure(); | ||
| } | ||
| } | ||
|
|
||
| } // end namespace | ||
|
|
||
| std::unique_ptr<OperationPass<ModuleOp>> CreateCanonicalizeBoundaryValuePass() { | ||
| return std::make_unique<CanonicalizeBoundaryValuePass>(); | ||
| } | ||
|
|
||
| } // end namespace TFL | ||
| } // end namespace mlir |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.