-
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 provide users with the option `canonicalizing_inf_a…
…s_min_max_float` to canonicalize boundary values in the IR replacing -Inf/Inf with MIN/MAX float value. An option is provided to the user to enable this option. PiperOrigin-RevId: 630372176
- Loading branch information
1 parent
f8fdab0
commit 9933d53
Showing
19 changed files
with
338 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
91 changes: 91 additions & 0 deletions
91
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,91 @@ | ||
| // 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> | ||
| } | ||
|
|
||
|
|
||
| // ----- | ||
|
|
||
| // CHECK-LABEL: func.func @clamp_pos_inf_f16_arith_op() -> f16 { | ||
| // CHECK: %[[CONST:.*]] = arith.constant 6.550400e+04 : f16 | ||
| // CHECK: return %[[CONST]] : f16 | ||
| func.func @clamp_pos_inf_f16_arith_op() -> f16 { | ||
| %ret = arith.constant 0x7C00 : f16 | ||
| return %ret : f16 | ||
| } | ||
|
|
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
120 changes: 120 additions & 0 deletions
120
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,120 @@ | ||
| /* 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 "llvm/ADT/STLExtras.h" | ||
| #include "mlir/Dialect/Arith/IR/Arith.h" // from @llvm-project | ||
| #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 { | ||
| Attribute attr = const_op.getValueAttr(); | ||
| if (auto float_attr = llvm::dyn_cast<FloatAttr>(attr)) { | ||
| if (float_attr.getValue().isInfinity()) { | ||
| FloatType float_type = llvm::dyn_cast<FloatType>(const_op.getType()); | ||
| if (!float_type) return failure(); | ||
| rewriter.replaceOpWithNewOp<OpTy>( | ||
| const_op, rewriter.getFloatAttr( | ||
| float_type, APFloat::getLargest( | ||
| float_type.getFloatSemantics(), | ||
| float_attr.getValue().isNegative()))); | ||
| return success(); | ||
| } | ||
| } | ||
|
|
||
| ElementsAttr tensor_attr = llvm::dyn_cast<ElementsAttr>(attr); | ||
| if (!tensor_attr) return failure(); | ||
|
|
||
| Type type = tensor_attr.getType(); | ||
| ShapedType tensor_type = llvm::cast<ShapedType>(type); | ||
| auto float_type = dyn_cast<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
Oops, something went wrong.