From 5d37bd0350f0144632629c1aa2ebaef6ca76300b Mon Sep 17 00:00:00 2001 From: TensorFlow Jenkins <16359713+tensorflow-jenkins@users.noreply.github.com> Date: Wed, 30 Nov 2022 11:05:04 -0800 Subject: [PATCH 01/52] lite: update androidx annotation location for opensource (#58705) PiperOrigin-RevId: 482398723 Co-authored-by: Zichuan Wei --- tensorflow/lite/java/BUILD | 4 ++-- tensorflow/workspace2.bzl | 16 ++++++++++++++++ tensorflow/workspace3.bzl | 9 +++++++++ 3 files changed, 27 insertions(+), 2 deletions(-) diff --git a/tensorflow/lite/java/BUILD b/tensorflow/lite/java/BUILD index 08f4c12be52ab3..15e5ab01b0364b 100644 --- a/tensorflow/lite/java/BUILD +++ b/tensorflow/lite/java/BUILD @@ -303,7 +303,7 @@ android_library( deps = [ ":tensorflowlite_gpu_api", "//tensorflow/lite/java:tensorflowlite_api", - "//third_party/java/androidx/annotation", + "@maven//:androidx_annotation_annotation", ], ) @@ -315,7 +315,7 @@ android_library( deps = [ ":tensorflowlite_api", ":tensorflowlite_gpu_api", - "//third_party/java/androidx/annotation", + "@maven//:androidx_annotation_annotation", ], ) diff --git a/tensorflow/workspace2.bzl b/tensorflow/workspace2.bzl index 18b63ab47fb438..182741583951c1 100644 --- a/tensorflow/workspace2.bzl +++ b/tensorflow/workspace2.bzl @@ -50,6 +50,7 @@ load("@tf_runtime//:dependencies.bzl", "tfrt_dependencies") load("//tensorflow/tools/toolchains/remote_config:configs.bzl", "initialize_rbe_configs") load("//tensorflow/tools/toolchains/remote:configure.bzl", "remote_execution_configure") load("//tensorflow/tools/toolchains/clang6:repo.bzl", "clang6_configure") +load("@rules_jvm_external//:defs.bzl", "maven_install") def _initialize_third_party(): """ Load third party repositories. See above load() statements. """ @@ -919,6 +920,21 @@ def _tf_repositories(): ], ) + # used for adding androidx.annotation dependencies in tflite android jni. + maven_install( + artifacts = [ + "androidx.annotation:annotation:aar:1.1.0", + ], + repositories = [ + "https://jcenter.bintray.com", + "https://maven.google.com", + "https://dl.google.com/dl/android/maven2", + "https://repo1.maven.org/maven2", + ], + fetch_sources = True, + version_conflict_policy = "pinned", + ) + def workspace(): # Check the bazel version before executing any repository rules, in case # those rules rely on the version we require here. diff --git a/tensorflow/workspace3.bzl b/tensorflow/workspace3.bzl index a6c2c5c5835d7f..1d02b6fe0dc266 100644 --- a/tensorflow/workspace3.bzl +++ b/tensorflow/workspace3.bzl @@ -36,6 +36,15 @@ def workspace(): sha256 = "8a298e832762eda1830597d64fe7db58178aa84cd5926d76d5b744d6558941c2", ) + # Maven dependencies. + RULES_JVM_EXTERNAL_TAG = "3.2" + http_archive( + name = "rules_jvm_external", + strip_prefix = "rules_jvm_external-%s" % RULES_JVM_EXTERNAL_TAG, + sha256 = "82262ff4223c5fda6fb7ff8bd63db8131b51b413d26eb49e3131037e79e324af", + url = "https://github.com/bazelbuild/rules_jvm_external/archive/%s.zip" % RULES_JVM_EXTERNAL_TAG, + ) + # Load the raw llvm-project. llvm does not have build rules set up by default, # but provides a script for setting up build rules via overlays. llvm("llvm-raw") From f2680d9966ee972b63571a2df786435ef29ac6c7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tom=C3=A1s=20Longeri?= Date: Fri, 23 Dec 2022 14:56:53 -0800 Subject: [PATCH 02/52] [tfg] Fix out-of-bounds access due to mismatched integer type sizes in ValueMap::Manager::GetValueOrCreatePlaceholder PiperOrigin-RevId: 497419488 --- .../ir/importexport/functiondef_import.cc | 30 +++++++---- ...id_generic_function_named_edge_index.pbtxt | 52 +++++++++++++++++++ 2 files changed, 72 insertions(+), 10 deletions(-) create mode 100644 tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt diff --git a/tensorflow/core/ir/importexport/functiondef_import.cc b/tensorflow/core/ir/importexport/functiondef_import.cc index 35c2d9270ff239..27a1769f37f2fa 100644 --- a/tensorflow/core/ir/importexport/functiondef_import.cc +++ b/tensorflow/core/ir/importexport/functiondef_import.cc @@ -86,11 +86,11 @@ class ValueMapManager { return ::tensorflow::OkStatus(); } - Value GetValueOrCreatePlaceholder(StringRef full_name) { + tensorflow::StatusOr GetValueOrCreatePlaceholder(StringRef full_name) { StringRef node_name; StringRef output_name = ""; bool is_control_dep = full_name[0] == '^'; - int output_num = 0; + size_t output_num = 0; if (is_control_dep) full_name = full_name.drop_front(); { size_t colon_sep = full_name.find_first_of(':'); @@ -105,8 +105,16 @@ class ValueMapManager { // NOLINTNEXTLINE: type matching the API taking a reference. unsigned long long value; if (!llvm::getAsUnsignedInteger(output_name.drop_front(colon_sep + 1), - 10, value)) - output_num = value; + 10, value)) { + if (LLVM_LIKELY( + value <= + std::numeric_limits>::max() - + 1)) + output_num = value; + else + return InvalidArgument("Output index ", value, + " is invalid (too large)"); + } output_name = output_name.take_front(colon_sep); } } @@ -171,8 +179,9 @@ Status ImportNodes(ValueMapManager value_manager, for (const std::string& input : node.input()) { if (input.empty()) return InvalidArgument("Node '", node.name(), "' has an empty input"); - state.operands.push_back( - value_manager.GetValueOrCreatePlaceholder(input)); + TF_ASSIGN_OR_RETURN(Value value, + value_manager.GetValueOrCreatePlaceholder(input)); + state.operands.push_back(value); } // Retrieve the entry in the nodes_map for this node and infer the result // count from what was inferred during the first traversal above. @@ -472,8 +481,9 @@ Status ImportGenericFunction( return InvalidArgument("Function '", func.signature().name(), "' has empty result name"); } - ret_vals[position->second] = - value_manager.GetValueOrCreatePlaceholder(ret_val.second); + TF_ASSIGN_OR_RETURN( + ret_vals[position->second], + value_manager.GetValueOrCreatePlaceholder(ret_val.second)); } for (const auto& ret_val : func.control_ret()) { auto position = control_output_to_position.find(ret_val.first); @@ -487,8 +497,8 @@ Status ImportGenericFunction( return InvalidArgument("Function '", func.signature().name(), "' has empty control result name"); } - Value result = value_manager.GetValueOrCreatePlaceholder( - (Twine("^") + ret_val.second).str()); + TF_ASSIGN_OR_RETURN(Value result, value_manager.GetValueOrCreatePlaceholder( + (Twine("^") + ret_val.second).str())); if (!result.getType().isa()) return InvalidArgument("failed to map returned value ", ret_val.second, ", isn't a control output"); diff --git a/tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt b/tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt new file mode 100644 index 00000000000000..7275ac6f1e6ea1 --- /dev/null +++ b/tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt @@ -0,0 +1,52 @@ +# RUN: not tfg-translate -graphdef-to-mlir %s 2>&1 | FileCheck %s + +# CHECK: Output index 18446744073709551615 is invalid (too large) + +library { + function { + signature { + name: "foo" + attr { + name: "T" + type: "type" + } + } + node_def { + name: "two" + op: "Const" + attr { + key: "dtype" + value { + type: DT_INT64 + } + } + attr { + key: "value" + value { + tensor { + dtype: DT_INT64 + tensor_shape {} + int64_val: 2 + } + } + } + } + node_def { + name: "a" + op: "Cast" + input: "two:output:18446744073709551615" + attr { + key: "DstT" + value { + placeholder: "T" + } + } + attr { + key: "SrcT" + value { + type: DT_INT64 + } + } + } + } +} From ec5e54727d395bb57cedc738edf6329c88cff177 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tom=C3=A1s=20Longeri?= Date: Fri, 23 Dec 2022 14:56:53 -0800 Subject: [PATCH 03/52] [tfg] Fix out-of-bounds access due to mismatched integer type sizes in ValueMap::Manager::GetValueOrCreatePlaceholder PiperOrigin-RevId: 497419488 --- .../ir/importexport/functiondef_import.cc | 30 +++++++---- ...id_generic_function_named_edge_index.pbtxt | 52 +++++++++++++++++++ 2 files changed, 72 insertions(+), 10 deletions(-) create mode 100644 tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt diff --git a/tensorflow/core/ir/importexport/functiondef_import.cc b/tensorflow/core/ir/importexport/functiondef_import.cc index 35c2d9270ff239..27a1769f37f2fa 100644 --- a/tensorflow/core/ir/importexport/functiondef_import.cc +++ b/tensorflow/core/ir/importexport/functiondef_import.cc @@ -86,11 +86,11 @@ class ValueMapManager { return ::tensorflow::OkStatus(); } - Value GetValueOrCreatePlaceholder(StringRef full_name) { + tensorflow::StatusOr GetValueOrCreatePlaceholder(StringRef full_name) { StringRef node_name; StringRef output_name = ""; bool is_control_dep = full_name[0] == '^'; - int output_num = 0; + size_t output_num = 0; if (is_control_dep) full_name = full_name.drop_front(); { size_t colon_sep = full_name.find_first_of(':'); @@ -105,8 +105,16 @@ class ValueMapManager { // NOLINTNEXTLINE: type matching the API taking a reference. unsigned long long value; if (!llvm::getAsUnsignedInteger(output_name.drop_front(colon_sep + 1), - 10, value)) - output_num = value; + 10, value)) { + if (LLVM_LIKELY( + value <= + std::numeric_limits>::max() - + 1)) + output_num = value; + else + return InvalidArgument("Output index ", value, + " is invalid (too large)"); + } output_name = output_name.take_front(colon_sep); } } @@ -171,8 +179,9 @@ Status ImportNodes(ValueMapManager value_manager, for (const std::string& input : node.input()) { if (input.empty()) return InvalidArgument("Node '", node.name(), "' has an empty input"); - state.operands.push_back( - value_manager.GetValueOrCreatePlaceholder(input)); + TF_ASSIGN_OR_RETURN(Value value, + value_manager.GetValueOrCreatePlaceholder(input)); + state.operands.push_back(value); } // Retrieve the entry in the nodes_map for this node and infer the result // count from what was inferred during the first traversal above. @@ -472,8 +481,9 @@ Status ImportGenericFunction( return InvalidArgument("Function '", func.signature().name(), "' has empty result name"); } - ret_vals[position->second] = - value_manager.GetValueOrCreatePlaceholder(ret_val.second); + TF_ASSIGN_OR_RETURN( + ret_vals[position->second], + value_manager.GetValueOrCreatePlaceholder(ret_val.second)); } for (const auto& ret_val : func.control_ret()) { auto position = control_output_to_position.find(ret_val.first); @@ -487,8 +497,8 @@ Status ImportGenericFunction( return InvalidArgument("Function '", func.signature().name(), "' has empty control result name"); } - Value result = value_manager.GetValueOrCreatePlaceholder( - (Twine("^") + ret_val.second).str()); + TF_ASSIGN_OR_RETURN(Value result, value_manager.GetValueOrCreatePlaceholder( + (Twine("^") + ret_val.second).str())); if (!result.getType().isa()) return InvalidArgument("failed to map returned value ", ret_val.second, ", isn't a control output"); diff --git a/tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt b/tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt new file mode 100644 index 00000000000000..7275ac6f1e6ea1 --- /dev/null +++ b/tensorflow/core/ir/importexport/tests/graphdef_to_mlir/invalid_generic_function_named_edge_index.pbtxt @@ -0,0 +1,52 @@ +# RUN: not tfg-translate -graphdef-to-mlir %s 2>&1 | FileCheck %s + +# CHECK: Output index 18446744073709551615 is invalid (too large) + +library { + function { + signature { + name: "foo" + attr { + name: "T" + type: "type" + } + } + node_def { + name: "two" + op: "Const" + attr { + key: "dtype" + value { + type: DT_INT64 + } + } + attr { + key: "value" + value { + tensor { + dtype: DT_INT64 + tensor_shape {} + int64_val: 2 + } + } + } + } + node_def { + name: "a" + op: "Cast" + input: "two:output:18446744073709551615" + attr { + key: "DstT" + value { + placeholder: "T" + } + } + attr { + key: "SrcT" + value { + type: DT_INT64 + } + } + } + } +} From d7f4cdce9cae94f7df5729e7fe35617fc7f91bc8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?C=2E=20Antonio=20S=C3=A1nchez?= Date: Tue, 24 Jan 2023 13:26:05 -0800 Subject: [PATCH 04/52] Fix audio spectrogram FPE. (#59427) Do input validation in shape function. PiperOrigin-RevId: 503481241 --- tensorflow/core/kernels/BUILD | 1 + .../core/kernels/spectrogram_op_test.cc | 41 +++++++++++++++++++ tensorflow/core/ops/audio_ops.cc | 10 +++++ 3 files changed, 52 insertions(+) diff --git a/tensorflow/core/kernels/BUILD b/tensorflow/core/kernels/BUILD index a89b9fe5f8b4ff..60b53a94cfdf41 100644 --- a/tensorflow/core/kernels/BUILD +++ b/tensorflow/core/kernels/BUILD @@ -6118,6 +6118,7 @@ tf_cuda_cc_test( "//tensorflow/core:test", "//tensorflow/core:test_main", "//tensorflow/core:testlib", + "//tensorflow/core/platform:status_matchers", ], ) diff --git a/tensorflow/core/kernels/spectrogram_op_test.cc b/tensorflow/core/kernels/spectrogram_op_test.cc index 380c51e6508b59..3ffb83a83f472e 100644 --- a/tensorflow/core/kernels/spectrogram_op_test.cc +++ b/tensorflow/core/kernels/spectrogram_op_test.cc @@ -19,6 +19,8 @@ limitations under the License. #include #include +#include +#include #include "tensorflow/cc/client/client_session.h" #include "tensorflow/cc/ops/audio_ops.h" #include "tensorflow/cc/ops/const_op.h" @@ -29,6 +31,9 @@ limitations under the License. #include "tensorflow/core/kernels/ops_util.h" #include "tensorflow/core/lib/core/status_test_util.h" #include "tensorflow/core/platform/test.h" +#include "tensorflow/tsl/lib/core/status_test_util.h" +#include "tensorflow/tsl/platform/errors.h" +#include "tensorflow/tsl/platform/status_matchers.h" namespace tensorflow { namespace ops { @@ -140,6 +145,42 @@ TEST(SpectrogramOpTest, MultichannelTest) { } } +TEST(SpectrogramOpTest, InvalidWindowSize) { + Scope root = Scope::NewRootScope(); + const int audio_size = 8; + const int channel_size = 2; + Tensor audio_tensor(DT_FLOAT, TensorShape({audio_size, channel_size})); + test::FillValues( + &audio_tensor, {-1.0f, -1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, -1.0f, + -1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}); + Output audio_const_op = Const(root.WithOpName("audio_const_op"), + Input::Initializer(audio_tensor)); + AudioSpectrogram spectrogram_op = + AudioSpectrogram(root.WithOpName("spectrogram_op"), audio_const_op, + /*window_size=*/1, /*stride=*/1); + EXPECT_THAT(root.status(), + tsl::testing::StatusIs(tsl::error::Code::INVALID_ARGUMENT, + ::testing::ContainsRegex("window size"))); +} + +TEST(SpectrogramOpTest, InvalidStride) { + Scope root = Scope::NewRootScope(); + const int audio_size = 8; + const int channel_size = 2; + Tensor audio_tensor(DT_FLOAT, TensorShape({audio_size, channel_size})); + test::FillValues( + &audio_tensor, {-1.0f, -1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, -1.0f, + -1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}); + Output audio_const_op = Const(root.WithOpName("audio_const_op"), + Input::Initializer(audio_tensor)); + AudioSpectrogram spectrogram_op = + AudioSpectrogram(root.WithOpName("spectrogram_op"), audio_const_op, + /*window_size=*/2, /*stride=*/0); + EXPECT_THAT(root.status(), + tsl::testing::StatusIs(tsl::error::Code::INVALID_ARGUMENT, + ::testing::ContainsRegex("stride"))); +} + } // namespace } // namespace ops } // namespace tensorflow diff --git a/tensorflow/core/ops/audio_ops.cc b/tensorflow/core/ops/audio_ops.cc index 807cd6d9b48ae0..240567c53072d2 100644 --- a/tensorflow/core/ops/audio_ops.cc +++ b/tensorflow/core/ops/audio_ops.cc @@ -17,6 +17,7 @@ limitations under the License. #include "tensorflow/core/framework/op.h" #include "tensorflow/core/framework/shape_inference.h" #include "tensorflow/core/lib/core/bits.h" +#include "tensorflow/core/platform/errors.h" namespace tensorflow { @@ -72,8 +73,17 @@ Status SpectrogramShapeFn(InferenceContext* c) { TF_RETURN_IF_ERROR(c->WithRank(c->input(0), 2, &input)); int32_t window_size; TF_RETURN_IF_ERROR(c->GetAttr("window_size", &window_size)); + if (window_size <= 1) { + return errors::InvalidArgument("window size must be > 1, got ", + window_size); + } + int32_t stride; TF_RETURN_IF_ERROR(c->GetAttr("stride", &stride)); + if (stride <= 0) { + return errors::InvalidArgument("stride must be strictly positive, got ", + stride); + } DimensionHandle input_length = c->Dim(input, 0); DimensionHandle input_channels = c->Dim(input, 1); From 32023b8a610c2371e3d90764a91d0e85cd000e05 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?C=2E=20Antonio=20S=C3=A1nchez?= Date: Tue, 24 Jan 2023 13:29:46 -0800 Subject: [PATCH 05/52] Add validation checks in sparse binary ops. (#59426) Fixes a nullptr security vulnerability. PiperOrigin-RevId: 503501621 --- .../kernels/sparse_sparse_binary_op_shared.cc | 43 ++++--------------- .../sparse_ops/sparse_ops_test.py | 13 ++++++ 2 files changed, 21 insertions(+), 35 deletions(-) diff --git a/tensorflow/core/kernels/sparse_sparse_binary_op_shared.cc b/tensorflow/core/kernels/sparse_sparse_binary_op_shared.cc index 8d0642eb1d4130..0863e5e37949bb 100644 --- a/tensorflow/core/kernels/sparse_sparse_binary_op_shared.cc +++ b/tensorflow/core/kernels/sparse_sparse_binary_op_shared.cc @@ -41,6 +41,7 @@ limitations under the License. #include "tensorflow/core/framework/types.h" #include "tensorflow/core/kernels/cwise_ops.h" #include "tensorflow/core/kernels/cwise_ops_common.h" +#include "tensorflow/core/kernels/sparse_utils.h" #include "tensorflow/core/util/sparse/sparse_tensor.h" namespace tensorflow { @@ -131,22 +132,12 @@ class SparseSparseBinaryOpShared : public OpKernel { OP_REQUIRES_OK(ctx, ctx->input("b_shape", &b_shape_t)); // Validations. - OP_REQUIRES( - ctx, - TensorShapeUtils::IsMatrix(a_indices_t->shape()) && - TensorShapeUtils::IsMatrix(b_indices_t->shape()), - errors::InvalidArgument("Inputs a_indices and b_indices should be " - "matrices but received shapes: ", - a_indices_t->shape().DebugString(), ", ", - b_indices_t->shape().DebugString())); - OP_REQUIRES(ctx, - TensorShapeUtils::IsVector(a_values_t->shape()) && - TensorShapeUtils::IsVector(b_values_t->shape()), - errors::InvalidArgument( - "Inputs a_values and b_values should be vectors " - "but received shapes: ", - a_values_t->shape().DebugString(), " and ", - b_values_t->shape().DebugString())); + OP_REQUIRES_OK(ctx, sparse_utils::ValidateSparseTensor( + *a_indices_t, *a_values_t, *a_shape_t, + sparse_utils::IndexValidation::kUnordered)); + OP_REQUIRES_OK(ctx, sparse_utils::ValidateSparseTensor( + *b_indices_t, *b_values_t, *b_shape_t, + sparse_utils::IndexValidation::kUnordered)); const int64_t a_nnz = a_indices_t->dim_size(0); const int64_t b_nnz = b_indices_t->dim_size(0); @@ -154,25 +145,7 @@ class SparseSparseBinaryOpShared : public OpKernel { const auto a_values = a_values_t->vec(); const auto b_values = b_values_t->vec(); - OP_REQUIRES( - ctx, a_values.size() == a_nnz && b_values.size() == b_nnz, - errors::InvalidArgument("Expected ", a_nnz, " and ", b_nnz, - " non-empty input values, got ", - a_values.size(), " and ", b_values.size())); - - OP_REQUIRES(ctx, - TensorShapeUtils::IsVector(a_shape_t->shape()) && - TensorShapeUtils::IsVector(b_shape_t->shape()), - errors::InvalidArgument( - "Input shapes should be a vector but received shapes ", - a_shape_t->shape().DebugString(), " and ", - b_shape_t->shape().DebugString())); const int num_dims = a_indices_t->dim_size(1); - OP_REQUIRES( - ctx, a_shape_t->NumElements() == num_dims, - errors::InvalidArgument("Second dimension of a_indices and length of " - "a_shape must match, got ", - num_dims, " and ", a_shape_t->NumElements())); OP_REQUIRES(ctx, num_dims > 0, errors::InvalidArgument("Tensors must not be empty")); OP_REQUIRES(ctx, a_shape_t->IsSameSize(*b_shape_t), @@ -192,7 +165,7 @@ class SparseSparseBinaryOpShared : public OpKernel { const auto a_indices_mat = a_indices_t->matrix(); const auto b_indices_mat = b_indices_t->matrix(); std::vector a_augmented_values, b_augmented_values; - std::vector> entries_to_copy; // from_a?, idx + std::vector> entries_to_copy; // from_a?, idx UnionSparseIndicesAndValues(a_indices_mat, a_values, a_nnz, b_indices_mat, b_values, b_nnz, num_dims, &a_augmented_values, &b_augmented_values, &entries_to_copy); diff --git a/tensorflow/python/kernel_tests/sparse_ops/sparse_ops_test.py b/tensorflow/python/kernel_tests/sparse_ops/sparse_ops_test.py index 60f61cf4f14366..ee46404aedf1a5 100644 --- a/tensorflow/python/kernel_tests/sparse_ops/sparse_ops_test.py +++ b/tensorflow/python/kernel_tests/sparse_ops/sparse_ops_test.py @@ -26,6 +26,7 @@ from tensorflow.python.framework import tensor_spec from tensorflow.python.framework import test_util from tensorflow.python.ops import array_ops +from tensorflow.python.ops import gen_sparse_ops from tensorflow.python.ops import gradient_checker from tensorflow.python.ops import nn_ops from tensorflow.python.ops import sparse_ops @@ -1175,6 +1176,18 @@ def testBasic(self): self._assertSparseTensorValueEqual(expected, max_tf) self._assertSparseTensorValueEqual(expected, min_tf) + def testInvalidSparseInputs(self): + with test_util.force_cpu(): + with self.assertRaisesRegex( + (ValueError, errors.InvalidArgumentError), + ".*Index rank .* and shape rank .* do not match.*", + ): + self.evaluate( + gen_sparse_ops.sparse_sparse_maximum( + [[1]], [0], [2], [[]], [1], [2] + ) + ) + @test_util.run_deprecated_v1 def testRandom(self): np.random.seed(1618) From ade9725336ad0fe466760b7c8f5292957db9550b Mon Sep 17 00:00:00 2001 From: David Majnemer Date: Thu, 26 Jan 2023 09:41:46 -0800 Subject: [PATCH 06/52] [tf2xla] Validate that stride and window size are positive PiperOrigin-RevId: 504866231 --- tensorflow/compiler/tests/pooling_ops_test.py | 30 +++++ .../compiler/tf2xla/kernels/pooling_ops.cc | 108 ++++++++++++------ tensorflow/compiler/xla/client/padding.cc | 10 ++ 3 files changed, 112 insertions(+), 36 deletions(-) diff --git a/tensorflow/compiler/tests/pooling_ops_test.py b/tensorflow/compiler/tests/pooling_ops_test.py index 3d2695b15e9b0f..3a7e22c02e54c5 100644 --- a/tensorflow/compiler/tests/pooling_ops_test.py +++ b/tensorflow/compiler/tests/pooling_ops_test.py @@ -18,7 +18,9 @@ from tensorflow.compiler.tests import xla_test from tensorflow.python.framework import dtypes +from tensorflow.python.framework import errors from tensorflow.python.framework import ops +from tensorflow.python.framework import test_util from tensorflow.python.ops import array_ops from tensorflow.python.ops import gen_nn_ops from tensorflow.python.ops import nn_ops @@ -560,6 +562,34 @@ def AvgPoolGrad(inputs, outputs, output_gradients, ksize, strides, padding, self._TestPooling(nn_ops.avg_pool, AvgPoolGrad) + @test_util.disable_mlir_bridge( + "TODO(b/266613412): investigate FPE in AvgPoolGrad for TPU" + ) + def testAvgPoolGradSamePaddingZeroStrideZeroSize(self): + output_gradient_vals = np.array([0.39117979], dtype=np.float32) + output_gradient_vals = output_gradient_vals.reshape([1, 1, 1, 1]) + with self.session() as sess: + with self.test_scope(): + output_gradients = array_ops.placeholder( + dtypes.float32, shape=output_gradient_vals.shape + ) + t = gen_nn_ops.avg_pool_grad( + orig_input_shape=[1, 0, 0, 0], + grad=output_gradients, + ksize=[1, 0, 0, 0], + strides=[1, 0, 0, 0], + padding="SAME", + data_format="NCHW", + ) + with self.assertRaisesRegex( + errors.InvalidArgumentError, + ( + "Sliding window ksize field for dimension 1 must be positive but" + " is 0" + ), + ): + sess.run(t, {output_gradients: output_gradient_vals}) + # The CPU implementation of AvgPoolGrad doesn't accept kernels smaller than # the stride size, so we only run the following tests on MaxPoolGrad. diff --git a/tensorflow/compiler/tf2xla/kernels/pooling_ops.cc b/tensorflow/compiler/tf2xla/kernels/pooling_ops.cc index 7f0d712f13df1a..64351c6a7416b8 100644 --- a/tensorflow/compiler/tf2xla/kernels/pooling_ops.cc +++ b/tensorflow/compiler/tf2xla/kernels/pooling_ops.cc @@ -33,15 +33,41 @@ limitations under the License. #include "tensorflow/compiler/xla/util.h" #include "tensorflow/core/framework/bounds_check.h" #include "tensorflow/core/framework/op_kernel.h" +#include "tensorflow/core/framework/op_requires.h" #include "tensorflow/core/framework/register_types.h" #include "tensorflow/core/framework/tensor.h" #include "tensorflow/core/platform/errors.h" #include "tensorflow/core/util/determinism.h" #include "tensorflow/core/util/tensor_format.h" +#include "tensorflow/tsl/platform/errors.h" namespace tensorflow { namespace { +template +static Status ValidateKernelSizes(const T& ksizes) { + for (size_t i = 0; i < ksizes.size(); ++i) { + if (ksizes[i] <= 0) { + return errors::InvalidArgument( + "Sliding window ksize field for dimension ", i, + " must be positive but is ", ksizes[i]); + } + } + return OkStatus(); +} + +template +static Status ValidateStrides(const T& strides) { + for (size_t i = 0; i < strides.size(); ++i) { + if (strides[i] <= 0) { + return errors::InvalidArgument( + "Sliding window stride field for dimension ", i, + " must be positive but is ", strides[i]); + } + } + return OkStatus(); +} + // Superclass of pooling ops. class PoolingOp : public XlaOpKernel { public: @@ -83,50 +109,54 @@ class PoolingOp : public XlaOpKernel { protected: StatusOr> GetKernelSize(XlaOpKernelContext* ctx) { - if (ctx->num_inputs() == 1) { - return ksize_; - } - const TensorShape ksize_shape = ctx->InputShape(1); - // Validate input sizes. - if (!TensorShapeUtils::IsVector(ksize_shape)) { - return errors::InvalidArgument("ksize must be a vector, not shape ", - ksize_shape.DebugString()); - } - if (ksize_shape.num_elements() != num_dims()) { - return errors::InvalidArgument( - "Sliding window ksize field must " - "specify ", - num_dims(), " dimensions"); - } std::vector ksize; - auto status = ctx->ConstantInputAsIntVector(1, &ksize); - if (!status.ok()) { - return status; + if (ctx->num_inputs() == 1) { + ksize = ksize_; + } else { + const TensorShape ksize_shape = ctx->InputShape(1); + // Validate input sizes. + if (!TensorShapeUtils::IsVector(ksize_shape)) { + return errors::InvalidArgument("ksize must be a vector, not shape ", + ksize_shape.DebugString()); + } + if (ksize_shape.num_elements() != num_dims()) { + return errors::InvalidArgument( + "Sliding window ksize field must " + "specify ", + num_dims(), " dimensions"); + } + auto status = ctx->ConstantInputAsIntVector(1, &ksize); + if (!status.ok()) { + return status; + } } + TF_RETURN_IF_ERROR(ValidateKernelSizes(ksize)); return ksize; } StatusOr> GetStride(XlaOpKernelContext* ctx) { - if (ctx->num_inputs() == 1) { - return stride_; - } - const TensorShape stride_shape = ctx->InputShape(2); - // Validate input sizes. - if (!TensorShapeUtils::IsVector(stride_shape)) { - return errors::InvalidArgument("stride must be a vector, not shape ", - stride_shape.DebugString()); - } - if (stride_shape.num_elements() != num_dims()) { - return errors::InvalidArgument( - "Sliding window stride field must " - "specify ", - num_dims(), " dimensions"); - } std::vector stride; - auto status = ctx->ConstantInputAsIntVector(2, &stride); - if (!status.ok()) { - return status; + if (ctx->num_inputs() == 1) { + stride = stride_; + } else { + const TensorShape stride_shape = ctx->InputShape(2); + // Validate input sizes. + if (!TensorShapeUtils::IsVector(stride_shape)) { + return errors::InvalidArgument("stride must be a vector, not shape ", + stride_shape.DebugString()); + } + if (stride_shape.num_elements() != num_dims()) { + return errors::InvalidArgument( + "Sliding window stride field must " + "specify ", + num_dims(), " dimensions"); + } + auto status = ctx->ConstantInputAsIntVector(2, &stride); + if (!status.ok()) { + return status; + } } + TF_RETURN_IF_ERROR(ValidateStrides(stride)); return stride; } @@ -355,10 +385,12 @@ class MaxPoolGradOp : public XlaOpKernel { errors::InvalidArgument("Sliding window ksize field must " "specify ", num_dims(), " dimensions")); + OP_REQUIRES_OK(ctx, ValidateKernelSizes(ksize_)); OP_REQUIRES(ctx, stride_.size() == num_dims(), errors::InvalidArgument("Sliding window strides field must " "specify ", num_dims(), " dimensions")); + OP_REQUIRES_OK(ctx, ValidateStrides(stride_)); const TensorShape tensor_in_shape = ctx->InputShape(0); const TensorShape tensor_out_shape = ctx->InputShape(1); @@ -446,11 +478,13 @@ class AvgPoolGradOp : public XlaOpKernel { errors::InvalidArgument("Sliding window ksize field must " "specify ", num_dims(), " dimensions")); + OP_REQUIRES_OK(ctx, ValidateKernelSizes(ksize_)); OP_REQUIRES_OK(ctx, ctx->GetAttr("strides", &stride_)); OP_REQUIRES(ctx, stride_.size() == num_dims(), errors::InvalidArgument("Sliding window strides field must " "specify ", num_dims(), " dimensions")); + OP_REQUIRES_OK(ctx, ValidateStrides(stride_)); OP_REQUIRES_OK(ctx, ctx->GetAttr("padding", &padding_)); OP_REQUIRES(ctx, padding_ != EXPLICIT, errors::Unimplemented( @@ -579,10 +613,12 @@ class MaxPoolGradGradOp : public XlaOpKernel { errors::InvalidArgument("Sliding window ksize field must " "specify ", num_dims(), " dimensions")); + OP_REQUIRES_OK(ctx, ValidateKernelSizes(ksize_)); OP_REQUIRES(ctx, stride_.size() == num_dims(), errors::InvalidArgument("Sliding window strides field must " "specify ", num_dims(), " dimensions")); + OP_REQUIRES_OK(ctx, ValidateStrides(stride_)); const TensorShape tensor_in_shape = ctx->InputShape(0); const TensorShape tensor_out_shape = ctx->InputShape(1); diff --git a/tensorflow/compiler/xla/client/padding.cc b/tensorflow/compiler/xla/client/padding.cc index 1fe8d556d1d19f..7fec04e2ac5408 100644 --- a/tensorflow/compiler/xla/client/padding.cc +++ b/tensorflow/compiler/xla/client/padding.cc @@ -35,6 +35,16 @@ Status ValidatePaddingValues(absl::Span input_dimensions, input_dimensions.size(), window_dimensions.size(), window_strides.size()); } + for (size_t i = 0; i < input_dimensions.size(); ++i) { + if (window_dimensions[i] <= 0) { + return InvalidArgument("Window dimension %u has non-positive size %d", i, + window_dimensions[i]); + } + if (window_strides[i] <= 0) { + return InvalidArgument("Window dimension %u has non-positive stride %d", + i, window_strides[i]); + } + } return OkStatus(); } From 73ed953f5171bade7a32fdcc331ff640c4ed099e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?C=2E=20Antonio=20S=C3=A1nchez?= Date: Fri, 27 Jan 2023 11:25:49 -0800 Subject: [PATCH 07/52] 504645490 by A. Unique TensorFlower: (#59477) Fix input validation for GRU ops. PiperOrigin-RevId: 504645490 Co-authored-by: A. Unique TensorFlower --- tensorflow/core/kernels/rnn/gru_ops.cc | 88 ++++++++++++++++---------- 1 file changed, 55 insertions(+), 33 deletions(-) diff --git a/tensorflow/core/kernels/rnn/gru_ops.cc b/tensorflow/core/kernels/rnn/gru_ops.cc index 702e618804da50..93aa88ca06733e 100644 --- a/tensorflow/core/kernels/rnn/gru_ops.cc +++ b/tensorflow/core/kernels/rnn/gru_ops.cc @@ -49,61 +49,68 @@ class GRUCellBlockOp : public OpKernel { const Tensor* b_c_tensor = nullptr; OP_REQUIRES_OK(ctx, ctx->input("b_c", &b_c_tensor)); + // Sanity checks for input shapes. + + // Shape of 'x' must be [batch_size, input_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(x_tensor->shape()), + errors::InvalidArgument("Rank of x must be 2", x_tensor->dims(), + " vs. 2")); const int64_t batch_size = x_tensor->dim_size(0); const int64_t input_size = x_tensor->dim_size(1); - const int64_t cell_size = h_prev_tensor->dim_size(1); - - // Sanity checks for input shapes. // Shape of 'h' must be [batch_size, cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(h_prev_tensor->shape()), + errors::InvalidArgument("Rank of h_prev must be 2, got ", + h_prev_tensor->dims())); OP_REQUIRES(ctx, h_prev_tensor->dim_size(0) == batch_size, errors::InvalidArgument("h_prev.dims(0) != batch_size: ", h_prev_tensor->dim_size(0), " vs. ", batch_size)); - OP_REQUIRES(ctx, h_prev_tensor->dim_size(1) == cell_size, - errors::InvalidArgument( - "h_prev.dims(1) != cell_size: ", h_prev_tensor->dim_size(1), - " vs. ", cell_size)); + const int64_t cell_size = h_prev_tensor->dim_size(1); // Shape of 'w_ru' must be [input_size+cell_size, 2*cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(w_ru_tensor->shape()), + errors::InvalidArgument("Rank of w_ru_ must be 2, got ", + w_ru_tensor->dims())); OP_REQUIRES(ctx, w_ru_tensor->dim_size(0) == input_size + cell_size, errors::InvalidArgument( "w_ru.dim_size(0) != input_size + cell_size: ", w_ru_tensor->dim_size(0), " vs. ", input_size + cell_size)); - OP_REQUIRES(ctx, w_ru_tensor->dim_size(1) == cell_size * 2, errors::InvalidArgument("w_ru.dim_size(1) != cell_size * 2: ", w_ru_tensor->dim_size(1), " vs. ", cell_size * 2)); // Shape of 'w_c' must be [input_size+cell_size, cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(w_c_tensor->shape()), + errors::InvalidArgument("Rank of w_c must be 2, got ", + w_c_tensor->dims())); OP_REQUIRES(ctx, w_c_tensor->dim_size(0) == input_size + cell_size, errors::InvalidArgument( "w_c.dim_size(0) != input_size + cell_size: ", w_c_tensor->dim_size(0), " vs. ", input_size + cell_size)); - OP_REQUIRES(ctx, w_c_tensor->dim_size(1) == cell_size, errors::InvalidArgument( "w_c.dim_size(1) != cell_size: ", w_c_tensor->dim_size(1), " vs. ", cell_size)); // Shape of 'b_ru' must be [2*cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsVector(b_ru_tensor->shape()), + errors::InvalidArgument("Rank of b_ru must be 1, got ", + b_ru_tensor->dims())); OP_REQUIRES(ctx, b_ru_tensor->dim_size(0) == cell_size * 2, errors::InvalidArgument("b_ru.dim_size(0) != cell_size * 2: ", b_ru_tensor->dim_size(0), " vs. ", cell_size * 2)); - OP_REQUIRES(ctx, b_ru_tensor->dims() == 1, - errors::InvalidArgument("Rank of b_ru must be 1", - b_ru_tensor->dims(), " vs. 1", 1)); // Shape of 'b_c' must be [cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsVector(b_c_tensor->shape()), + errors::InvalidArgument("Rank of b_c must be 1, got ", + b_c_tensor->dims())); OP_REQUIRES(ctx, b_c_tensor->dim_size(0) == cell_size, errors::InvalidArgument( "b_c.dim_size(0) != cell_size: ", b_c_tensor->dim_size(0), " vs. ", cell_size)); - OP_REQUIRES(ctx, b_c_tensor->dims() == 1, - errors::InvalidArgument("Rank of b_c must be 1", - b_c_tensor->dims(), " vs. 1")); // Create output tensors. Tensor* r_tensor = nullptr; @@ -204,65 +211,71 @@ class GRUBlockCellGradOp : public OpKernel { const Tensor* d_h_tensor = nullptr; OP_REQUIRES_OK(ctx, ctx->input("d_h", &d_h_tensor)); + // Shape of 'x' must be [batch_size, input_size] + OP_REQUIRES( + ctx, TensorShapeUtils::IsMatrix(x_tensor->shape()), + errors::InvalidArgument("Rank of x must be 2, got ", x_tensor->dims())); const int64_t batch_size = x_tensor->dim_size(0); const int64_t input_size = x_tensor->dim_size(1); - const int64_t cell_size = h_prev_tensor->dim_size(1); - // Sanity checks for input shapes. - - // Shape of 'h_prev' must be [batch_size, cell_size] + // Shape of 'h' must be [batch_size, cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(h_prev_tensor->shape()), + errors::InvalidArgument("Rank of h_prev must be 2, got ", + h_prev_tensor->dims())); OP_REQUIRES(ctx, h_prev_tensor->dim_size(0) == batch_size, errors::InvalidArgument("h_prev.dims(0) != batch_size: ", h_prev_tensor->dim_size(0), " vs. ", batch_size)); - OP_REQUIRES(ctx, h_prev_tensor->dim_size(1) == cell_size, - errors::InvalidArgument( - "h_prev.dims(1) != cell_size: ", h_prev_tensor->dim_size(1), - " vs. ", cell_size)); + const int64_t cell_size = h_prev_tensor->dim_size(1); // Shape of 'w_ru' must be [input_size+cell_size, 2*cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(w_ru_tensor->shape()), + errors::InvalidArgument("Rank of w_ru_ must be 2, got ", + w_ru_tensor->dims())); OP_REQUIRES(ctx, w_ru_tensor->dim_size(0) == input_size + cell_size, errors::InvalidArgument( "w_ru.dim_size(0) != input_size + cell_size: ", w_ru_tensor->dim_size(0), " vs. ", input_size + cell_size)); - OP_REQUIRES(ctx, w_ru_tensor->dim_size(1) == cell_size * 2, errors::InvalidArgument("w_ru.dim_size(1) != cell_size * 2: ", w_ru_tensor->dim_size(1), " vs. ", cell_size * 2)); // Shape of 'w_c' must be [input_size+cell_size, cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(w_c_tensor->shape()), + errors::InvalidArgument("Rank of w_c must be 2, got ", + w_c_tensor->dims())); OP_REQUIRES(ctx, w_c_tensor->dim_size(0) == input_size + cell_size, errors::InvalidArgument( "w_c.dim_size(0) != input_size + cell_size: ", w_c_tensor->dim_size(0), " vs. ", input_size + cell_size)); - OP_REQUIRES(ctx, w_c_tensor->dim_size(1) == cell_size, errors::InvalidArgument( "w_c.dim_size(1) != cell_size: ", w_c_tensor->dim_size(1), " vs. ", cell_size)); // Shape of 'b_ru' must be [2*cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsVector(b_ru_tensor->shape()), + errors::InvalidArgument("Rank of b_ru must be 1, got ", + b_ru_tensor->dims())); OP_REQUIRES(ctx, b_ru_tensor->dim_size(0) == cell_size * 2, errors::InvalidArgument("b_ru.dim_size(0) != cell_size * 2: ", b_ru_tensor->dim_size(0), " vs. ", cell_size * 2)); - OP_REQUIRES(ctx, b_ru_tensor->dims() == 1, - errors::InvalidArgument("Rank of b_ru must be 1", - b_ru_tensor->dims(), " vs. 1")); - // Shape of 'b_c' must be [cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsVector(b_c_tensor->shape()), + errors::InvalidArgument("Rank of b_c must be 1, got ", + b_c_tensor->dims())); OP_REQUIRES(ctx, b_c_tensor->dim_size(0) == cell_size, errors::InvalidArgument( "b_c.dim_size(0) != cell_size: ", b_c_tensor->dim_size(0), " vs. ", cell_size)); - OP_REQUIRES(ctx, b_c_tensor->dims() == 1, - errors::InvalidArgument("Rank of b_c must be 1 ", - b_c_tensor->dims(), " vs. 1")); - // Shape of 'r' must be [batch_size, cell_size] + OP_REQUIRES( + ctx, TensorShapeUtils::IsMatrix(r_tensor->shape()), + errors::InvalidArgument("Rank of r must be 2, got ", r_tensor->dims())); OP_REQUIRES(ctx, r_tensor->dim_size(0) == batch_size, errors::InvalidArgument( "r.dims(0) != batch_size: ", r_tensor->dim_size(0), " vs. ", @@ -273,6 +286,9 @@ class GRUBlockCellGradOp : public OpKernel { cell_size)); // Shape of 'u' must be [batch_size, cell_size] + OP_REQUIRES( + ctx, TensorShapeUtils::IsMatrix(u_tensor->shape()), + errors::InvalidArgument("Rank of u must be 2, got ", u_tensor->dims())); OP_REQUIRES(ctx, u_tensor->dim_size(0) == batch_size, errors::InvalidArgument( "u.dims(0) != batch_size: ", u_tensor->dim_size(0), " vs. ", @@ -283,6 +299,9 @@ class GRUBlockCellGradOp : public OpKernel { cell_size)); // Shape of 'c' must be [batch_size, cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(c_tensor->shape()), + errors::InvalidArgument("Rank of w_c must be 2, got ", + c_tensor->dims())); OP_REQUIRES(ctx, c_tensor->dim_size(0) == batch_size, errors::InvalidArgument( "c.dims(0) != batch_size: ", c_tensor->dim_size(0), " vs. ", @@ -293,6 +312,9 @@ class GRUBlockCellGradOp : public OpKernel { cell_size)); // Shape of 'd_h' must be [batch_size, cell_size] + OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(d_h_tensor->shape()), + errors::InvalidArgument("Rank of d_h must be 2, got ", + d_h_tensor->dims())); OP_REQUIRES(ctx, d_h_tensor->dim_size(0) == batch_size, errors::InvalidArgument( "d_h.dims(0) != batch_size: ", d_h_tensor->dim_size(0), From cbda0becdb646fe95b793994219a91d19961e720 Mon Sep 17 00:00:00 2001 From: He Jiang Date: Fri, 27 Jan 2023 07:26:44 -0800 Subject: [PATCH 08/52] Support forceBackend option in GPU Delegate Options. PiperOrigin-RevId: 505112394 --- .../org/tensorflow/lite/gpu/GpuDelegate.java | 6 ++-- .../lite/gpu/GpuDelegateFactory.java | 33 +++++++++++++++++++ .../delegates/gpu/java/src/main/native/BUILD | 1 + .../java/src/main/native/gpu_delegate_jni.cc | 15 ++++++++- .../tensorflow/lite/gpu/GpuDelegateTest.java | 18 ++++++++++ 5 files changed, 70 insertions(+), 3 deletions(-) diff --git a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java index f92b47f3ffd63d..b69f387b922232 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java +++ b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java @@ -44,7 +44,8 @@ public GpuDelegate(GpuDelegateFactory.Options options) { options.areQuantizedModelsAllowed(), options.getInferencePreference(), options.getSerializationDir(), - options.getModelToken()); + options.getModelToken(), + options.getForceBackend().value()); } @UsedByReflection("TFLiteSupport/model/GpuDelegateProxy") @@ -83,7 +84,8 @@ private static native long createDelegate( boolean quantizedModelsAllowed, int preference, String serializationDir, - String modelToken); + String modelToken, + int forceBackend); private static native void deleteDelegate(long delegateHandle); } diff --git a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java index c4e3d104a59113..f673cbd7de4e65 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java +++ b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java @@ -42,6 +42,26 @@ public Options() {} */ public static final int INFERENCE_PREFERENCE_SUSTAINED_SPEED = 1; + /** Which GPU backend to select. */ + public enum GpuBackend { + /** Try OpenCL first. Fall back to OpenGL if it's not available. */ + UNSET(0), + /** Enforces execution with Open CL. */ + OPENCL(1), + /** Enforces execution with Open GL. */ + OPENGL(2); + private final int value; + + GpuBackend(int value) { + this.value = value; + } + + /** int value of this enum. */ + public int value() { + return value; + } + } + /** * Sets whether precision loss is allowed. * @@ -96,6 +116,14 @@ public Options setSerializationParams(String serializationDir, String modelToken return this; } + /** + * Sets the GPU Backend. + */ + public Options setForceBackend(GpuBackend forceBackend) { + this.forceBackend = forceBackend; + return this; + } + public boolean isPrecisionLossAllowed() { return precisionLossAllowed; } @@ -116,11 +144,16 @@ public String getModelToken() { return modelToken; } + public GpuBackend getForceBackend() { + return forceBackend; + } + private boolean precisionLossAllowed = true; boolean quantizedModelsAllowed = true; int inferencePreference = INFERENCE_PREFERENCE_FAST_SINGLE_ANSWER; String serializationDir = null; String modelToken = null; + GpuBackend forceBackend = GpuBackend.UNSET; } public GpuDelegateFactory() { diff --git a/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD b/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD index d8f87c782deb25..8dc04b326983c2 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD +++ b/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD @@ -50,6 +50,7 @@ cc_library_with_tflite( ], deps = [ "//tensorflow/lite/delegates/gpu:delegate", + "//tensorflow/lite/delegates/gpu:delegate_options", "//tensorflow/lite/delegates/gpu/common:gpu_info", "//tensorflow/lite/delegates/gpu/gl:egl_environment", "//tensorflow/lite/delegates/gpu/gl:request_gpu_info", diff --git a/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc b/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc index 0f32d254a2b492..16ad8fba81ecbd 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc +++ b/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc @@ -26,6 +26,7 @@ limitations under the License. #include "tensorflow/lite/experimental/acceleration/configuration/configuration_generated.h" #else #include "tensorflow/lite/delegates/gpu/delegate.h" +#include "tensorflow/lite/delegates/gpu/delegate_options.h" #endif #if TFLITE_DISABLE_SELECT_JAVA_APIS @@ -36,6 +37,9 @@ using tflite::GPUSettings; using tflite::GPUSettingsBuilder; using tflite::TFLiteSettings; using tflite::TFLiteSettingsBuilder; +#else +constexpr int kGpuBackendOpenCl = 1; +constexpr int kGpuBackendOpenGl = 2; #endif extern "C" { @@ -43,7 +47,7 @@ extern "C" { JNIEXPORT jlong JNICALL Java_org_tensorflow_lite_gpu_GpuDelegate_createDelegate( JNIEnv* env, jclass clazz, jboolean precision_loss_allowed, jboolean quantized_models_allowed, jint inference_preference, - jstring serialization_dir, jstring model_token) { + jstring serialization_dir, jstring model_token, jint force_backend) { if (!tflite::jni::CheckJniInitializedOrThrow(env)) return 0; #if TFLITE_DISABLE_SELECT_JAVA_APIS @@ -99,6 +103,10 @@ JNIEXPORT jlong JNICALL Java_org_tensorflow_lite_gpu_GpuDelegate_createDelegate( gpu_settings_builder.add_inference_preference( static_cast(inference_preference)); } + if (force_backend) { + gpu_settings_builder.add_force_backend( + static_cast(force_backend)); + } Offset gpu_settings = gpu_settings_builder.Finish(); TFLiteSettingsBuilder tflite_settings_builder(flatbuffer_builder); tflite_settings_builder.add_gpu_settings(gpu_settings); @@ -138,6 +146,11 @@ JNIEXPORT jlong JNICALL Java_org_tensorflow_lite_gpu_GpuDelegate_createDelegate( options.experimental_flags |= TFLITE_GPU_EXPERIMENTAL_FLAGS_ENABLE_SERIALIZATION; } + if (force_backend == kGpuBackendOpenCl) { + options.experimental_flags |= TFLITE_GPU_EXPERIMENTAL_FLAGS_CL_ONLY; + } else if (force_backend == kGpuBackendOpenGl) { + options.experimental_flags |= TFLITE_GPU_EXPERIMENTAL_FLAGS_GL_ONLY; + } return reinterpret_cast(TfLiteGpuDelegateV2Create(&options)); #endif // TFLITE_DISABLE_SELECT_JAVA_APIS } diff --git a/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java b/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java index 7994b67c061b5a..2dd9e00a182da9 100644 --- a/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java +++ b/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java @@ -17,6 +17,8 @@ import static com.google.common.truth.Truth.assertThat; import static java.util.concurrent.TimeUnit.MICROSECONDS; +import static org.junit.Assert.assertThrows; +import static org.tensorflow.lite.gpu.GpuDelegateFactory.Options.GpuBackend.OPENCL; import com.google.common.base.Stopwatch; import java.io.File; @@ -120,6 +122,22 @@ public void testInterpreterWithGpu_QuantModelRunOnCPU() throws Exception { } } + @Test + public void testInterpreterWithGpu_forceOpenCl_throwsException() { + Interpreter.Options options = new Interpreter.Options(); + try (GpuDelegate delegate = + new GpuDelegate(new GpuDelegateFactory.Options().setForceBackend(OPENCL))) { + IllegalArgumentException e = + assertThrows( + IllegalArgumentException.class, + // Create interpreter fails because OpenCL is not available on device. + () -> + new Interpreter(MOBILENET_QUANTIZED_MODEL_BUFFER, options.addDelegate(delegate))); + + assertThat(e).hasMessageThat().contains("Can not open OpenCL library"); + } + } + @Test public void testDelegateSerialization() throws Exception { ByteBuffer img = From 0b0dfcc72685b3c43dd3808f0d256a255eab9bd5 Mon Sep 17 00:00:00 2001 From: Vinila S <106367904+vinila21@users.noreply.github.com> Date: Fri, 27 Jan 2023 22:07:47 -0800 Subject: [PATCH 09/52] Revert "r2.11 cherry-pick: 7a3d9601d6a "Support forceBackend option in GPU Delegate Options."" --- .../org/tensorflow/lite/gpu/GpuDelegate.java | 6 ++-- .../lite/gpu/GpuDelegateFactory.java | 33 ------------------- .../delegates/gpu/java/src/main/native/BUILD | 1 - .../java/src/main/native/gpu_delegate_jni.cc | 15 +-------- .../tensorflow/lite/gpu/GpuDelegateTest.java | 18 ---------- 5 files changed, 3 insertions(+), 70 deletions(-) diff --git a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java index b69f387b922232..f92b47f3ffd63d 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java +++ b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegate.java @@ -44,8 +44,7 @@ public GpuDelegate(GpuDelegateFactory.Options options) { options.areQuantizedModelsAllowed(), options.getInferencePreference(), options.getSerializationDir(), - options.getModelToken(), - options.getForceBackend().value()); + options.getModelToken()); } @UsedByReflection("TFLiteSupport/model/GpuDelegateProxy") @@ -84,8 +83,7 @@ private static native long createDelegate( boolean quantizedModelsAllowed, int preference, String serializationDir, - String modelToken, - int forceBackend); + String modelToken); private static native void deleteDelegate(long delegateHandle); } diff --git a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java index f673cbd7de4e65..c4e3d104a59113 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java +++ b/tensorflow/lite/delegates/gpu/java/src/main/java/org/tensorflow/lite/gpu/GpuDelegateFactory.java @@ -42,26 +42,6 @@ public Options() {} */ public static final int INFERENCE_PREFERENCE_SUSTAINED_SPEED = 1; - /** Which GPU backend to select. */ - public enum GpuBackend { - /** Try OpenCL first. Fall back to OpenGL if it's not available. */ - UNSET(0), - /** Enforces execution with Open CL. */ - OPENCL(1), - /** Enforces execution with Open GL. */ - OPENGL(2); - private final int value; - - GpuBackend(int value) { - this.value = value; - } - - /** int value of this enum. */ - public int value() { - return value; - } - } - /** * Sets whether precision loss is allowed. * @@ -116,14 +96,6 @@ public Options setSerializationParams(String serializationDir, String modelToken return this; } - /** - * Sets the GPU Backend. - */ - public Options setForceBackend(GpuBackend forceBackend) { - this.forceBackend = forceBackend; - return this; - } - public boolean isPrecisionLossAllowed() { return precisionLossAllowed; } @@ -144,16 +116,11 @@ public String getModelToken() { return modelToken; } - public GpuBackend getForceBackend() { - return forceBackend; - } - private boolean precisionLossAllowed = true; boolean quantizedModelsAllowed = true; int inferencePreference = INFERENCE_PREFERENCE_FAST_SINGLE_ANSWER; String serializationDir = null; String modelToken = null; - GpuBackend forceBackend = GpuBackend.UNSET; } public GpuDelegateFactory() { diff --git a/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD b/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD index 8dc04b326983c2..d8f87c782deb25 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD +++ b/tensorflow/lite/delegates/gpu/java/src/main/native/BUILD @@ -50,7 +50,6 @@ cc_library_with_tflite( ], deps = [ "//tensorflow/lite/delegates/gpu:delegate", - "//tensorflow/lite/delegates/gpu:delegate_options", "//tensorflow/lite/delegates/gpu/common:gpu_info", "//tensorflow/lite/delegates/gpu/gl:egl_environment", "//tensorflow/lite/delegates/gpu/gl:request_gpu_info", diff --git a/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc b/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc index 16ad8fba81ecbd..0f32d254a2b492 100644 --- a/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc +++ b/tensorflow/lite/delegates/gpu/java/src/main/native/gpu_delegate_jni.cc @@ -26,7 +26,6 @@ limitations under the License. #include "tensorflow/lite/experimental/acceleration/configuration/configuration_generated.h" #else #include "tensorflow/lite/delegates/gpu/delegate.h" -#include "tensorflow/lite/delegates/gpu/delegate_options.h" #endif #if TFLITE_DISABLE_SELECT_JAVA_APIS @@ -37,9 +36,6 @@ using tflite::GPUSettings; using tflite::GPUSettingsBuilder; using tflite::TFLiteSettings; using tflite::TFLiteSettingsBuilder; -#else -constexpr int kGpuBackendOpenCl = 1; -constexpr int kGpuBackendOpenGl = 2; #endif extern "C" { @@ -47,7 +43,7 @@ extern "C" { JNIEXPORT jlong JNICALL Java_org_tensorflow_lite_gpu_GpuDelegate_createDelegate( JNIEnv* env, jclass clazz, jboolean precision_loss_allowed, jboolean quantized_models_allowed, jint inference_preference, - jstring serialization_dir, jstring model_token, jint force_backend) { + jstring serialization_dir, jstring model_token) { if (!tflite::jni::CheckJniInitializedOrThrow(env)) return 0; #if TFLITE_DISABLE_SELECT_JAVA_APIS @@ -103,10 +99,6 @@ JNIEXPORT jlong JNICALL Java_org_tensorflow_lite_gpu_GpuDelegate_createDelegate( gpu_settings_builder.add_inference_preference( static_cast(inference_preference)); } - if (force_backend) { - gpu_settings_builder.add_force_backend( - static_cast(force_backend)); - } Offset gpu_settings = gpu_settings_builder.Finish(); TFLiteSettingsBuilder tflite_settings_builder(flatbuffer_builder); tflite_settings_builder.add_gpu_settings(gpu_settings); @@ -146,11 +138,6 @@ JNIEXPORT jlong JNICALL Java_org_tensorflow_lite_gpu_GpuDelegate_createDelegate( options.experimental_flags |= TFLITE_GPU_EXPERIMENTAL_FLAGS_ENABLE_SERIALIZATION; } - if (force_backend == kGpuBackendOpenCl) { - options.experimental_flags |= TFLITE_GPU_EXPERIMENTAL_FLAGS_CL_ONLY; - } else if (force_backend == kGpuBackendOpenGl) { - options.experimental_flags |= TFLITE_GPU_EXPERIMENTAL_FLAGS_GL_ONLY; - } return reinterpret_cast(TfLiteGpuDelegateV2Create(&options)); #endif // TFLITE_DISABLE_SELECT_JAVA_APIS } diff --git a/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java b/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java index 2dd9e00a182da9..7994b67c061b5a 100644 --- a/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java +++ b/tensorflow/lite/java/src/test/java/org/tensorflow/lite/gpu/GpuDelegateTest.java @@ -17,8 +17,6 @@ import static com.google.common.truth.Truth.assertThat; import static java.util.concurrent.TimeUnit.MICROSECONDS; -import static org.junit.Assert.assertThrows; -import static org.tensorflow.lite.gpu.GpuDelegateFactory.Options.GpuBackend.OPENCL; import com.google.common.base.Stopwatch; import java.io.File; @@ -122,22 +120,6 @@ public void testInterpreterWithGpu_QuantModelRunOnCPU() throws Exception { } } - @Test - public void testInterpreterWithGpu_forceOpenCl_throwsException() { - Interpreter.Options options = new Interpreter.Options(); - try (GpuDelegate delegate = - new GpuDelegate(new GpuDelegateFactory.Options().setForceBackend(OPENCL))) { - IllegalArgumentException e = - assertThrows( - IllegalArgumentException.class, - // Create interpreter fails because OpenCL is not available on device. - () -> - new Interpreter(MOBILENET_QUANTIZED_MODEL_BUFFER, options.addDelegate(delegate))); - - assertThat(e).hasMessageThat().contains("Can not open OpenCL library"); - } - } - @Test public void testDelegateSerialization() throws Exception { ByteBuffer img = From 87bf986bb958da468a4365ed9a1f7bec4696fe37 Mon Sep 17 00:00:00 2001 From: Zhufeng Pan Date: Wed, 16 Nov 2022 10:24:38 -0800 Subject: [PATCH 10/52] Add inputs check for AvgPoolGrad PiperOrigin-RevId: 488975844 --- tensorflow/core/kernels/avgpooling_op.cc | 13 +++++++++++++ .../kernel_tests/nn_ops/pooling_ops_test.py | 15 +++++++++++++++ 2 files changed, 28 insertions(+) diff --git a/tensorflow/core/kernels/avgpooling_op.cc b/tensorflow/core/kernels/avgpooling_op.cc index bc536d187512af..4ee4b8ff324748 100644 --- a/tensorflow/core/kernels/avgpooling_op.cc +++ b/tensorflow/core/kernels/avgpooling_op.cc @@ -342,6 +342,19 @@ class AvgPoolingGradOp : public OpKernel { const T* out_backprop_ptr = out_backprop.flat().data(); T* input_backprop_ptr = output->flat().data(); + for (int64_t r = 0; r < out_backprop_rows; ++r) { + int rindex, rsize; + OP_REQUIRES_OK(context, + GetBroadcastSize(r, in_rows, window_rows, row_stride, + pad_rows, &rindex, &rsize)); + for (int64_t c = 0; c < out_backprop_cols; ++c) { + int cindex, csize; + OP_REQUIRES_OK(context, + GetBroadcastSize(c, in_cols, window_cols, col_stride, + pad_cols, &cindex, &csize)); + } + } + auto shard = [context, out_backprop_ptr, input_backprop_ptr, out_backprop_rows, out_backprop_cols, out_backprop_depth, in_rows, in_cols, window_rows, window_cols, row_stride, diff --git a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py index e279f6e0027da2..3c6f974d784513 100644 --- a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py @@ -2510,6 +2510,21 @@ def testAvgPoolGradInvalidInputShapeRaiseError(self): data_format="NHWC") self.evaluate(t) + def testAvgPoolGradInvalidStrideRaiseErrorProperly(self): + with self.assertRaises(errors_impl.InvalidArgumentError): + with self.cached_session(): + orig_input_shape = [11, 9, 78, 9] + grad = constant_op.constant( + 0.1, shape=[16, 16, 16, 16], dtype=dtypes.float64) + t = gen_nn_ops.AvgPoolGrad( + orig_input_shape=orig_input_shape, + grad=grad, + ksize=[1, 40, 128, 1], + strides=[1, 128, 128, 30], + padding="SAME", + data_format="NHWC") + self.evaluate(t) + def GetMaxPoolFwdTest(input_size, filter_size, strides, padding): From 2bae31327f77df0510542561d0a08843396dc0d4 Mon Sep 17 00:00:00 2001 From: Zhufeng Pan Date: Tue, 8 Nov 2022 11:12:10 -0800 Subject: [PATCH 11/52] Fix nn.conv3d_transpose security vulnerability with illegal input shape PiperOrigin-RevId: 487000656 --- tensorflow/core/kernels/conv_grad_ops_3d.cc | 5 +++++ .../kernel_tests/nn_ops/conv3d_transpose_test.py | 16 +++++++++++++++- 2 files changed, 20 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/kernels/conv_grad_ops_3d.cc b/tensorflow/core/kernels/conv_grad_ops_3d.cc index 73cb7b02c7c0d0..2728ae9f7dd252 100644 --- a/tensorflow/core/kernels/conv_grad_ops_3d.cc +++ b/tensorflow/core/kernels/conv_grad_ops_3d.cc @@ -1322,6 +1322,11 @@ class Conv3DBackpropInputOp : public OpKernel { Tensor* in_backprop; OP_REQUIRES_OK(context, context->allocate_output(0, input_shape, &in_backprop)); + for (std::size_t i = 0; i < input_shape.dims(); ++i) { + if (input_shape.dim_size(i) == 0) { + return; + } + } auto* stream = context->op_device_context()->stream(); OP_REQUIRES(context, stream, errors::Internal("No GPU stream available.")); diff --git a/tensorflow/python/kernel_tests/nn_ops/conv3d_transpose_test.py b/tensorflow/python/kernel_tests/nn_ops/conv3d_transpose_test.py index 1e683808296af9..6d185bb76463e0 100644 --- a/tensorflow/python/kernel_tests/nn_ops/conv3d_transpose_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/conv3d_transpose_test.py @@ -126,7 +126,7 @@ def testConv3DTransposeShapeMismatch(self): x_value = np.random.random_sample(x_shape).astype(np.float64) f_value = np.random.random_sample(f_shape).astype(np.float64) nn_ops.conv3d_transpose( - x_value, f_value, y_shape, strides, data_format='NCDHW') + x_value, f_value, y_shape, strides, data_format="NCDHW") def testConv3DTransposeOutputShapeType(self): # Test case for GitHub issue 18887 @@ -218,6 +218,20 @@ def testGradient(self): err_tolerance = 0.00055 self.assertLess(err, err_tolerance) + def testConv3DTransposeZeroShapeDoNotRaiseError(self): + with self.cached_session(): + x_value = np.zeros([10, 0, 2, 3, 3]) + f_value = np.ones((3, 3, 3, 3, 3)) + y_shape = np.stack([10, 0, 2, 3, 3]) + output = nn_ops.conv3d_transpose( + x_value, + f_value, + y_shape, + strides=(1, 1, 1), + data_format="NDHWC", + padding="SAME", + ) + _ = self.evaluate(output) if __name__ == "__main__": test.main() From fb5edf1558beaa0f553deccad904b44a0378f14b Mon Sep 17 00:00:00 2001 From: Sagun Bajra Date: Wed, 9 Nov 2022 15:40:12 -0800 Subject: [PATCH 12/52] Fix asan issue with QuantizeAndDequantizeV2/V3/V4/V4Grad shape inference functions. PiperOrigin-RevId: 487366553 --- tensorflow/core/ops/array_ops.cc | 16 +++++ .../kernel_tests/array_ops/array_ops_test.py | 66 +++++++++++++++++++ 2 files changed, 82 insertions(+) diff --git a/tensorflow/core/ops/array_ops.cc b/tensorflow/core/ops/array_ops.cc index b57a44f44fdd73..8e52dde528eaf6 100644 --- a/tensorflow/core/ops/array_ops.cc +++ b/tensorflow/core/ops/array_ops.cc @@ -2879,6 +2879,10 @@ REGISTER_OP("QuantizeAndDequantizeV2") axis); } else if (axis != -1) { ShapeHandle input; + if (axis >= kint32max) { + return errors::InvalidArgument( + "Axis cannot be >= kint32max value, got ", axis); + } TF_RETURN_IF_ERROR(c->WithRankAtLeast(c->input(0), axis + 1, &input)); DimensionHandle depth; TF_RETURN_IF_ERROR( @@ -2914,6 +2918,10 @@ REGISTER_OP("QuantizeAndDequantizeV4") axis); } else if (axis != -1) { ShapeHandle input; + if (axis >= kint32max) { + return errors::InvalidArgument( + "Axis cannot be >= kint32max value, got ", axis); + } TF_RETURN_IF_ERROR(c->WithRankAtLeast(c->input(0), axis + 1, &input)); DimensionHandle depth; TF_RETURN_IF_ERROR( @@ -2945,6 +2953,10 @@ REGISTER_OP("QuantizeAndDequantizeV4Grad") axis); } else if (axis != -1) { ShapeHandle input; + if (axis >= kint32max) { + return errors::InvalidArgument( + "Axis cannot be >= kint32max value, got ", axis); + } TF_RETURN_IF_ERROR(c->WithRankAtLeast(c->input(0), axis + 1, &input)); DimensionHandle depth; TF_RETURN_IF_ERROR( @@ -2981,6 +2993,10 @@ REGISTER_OP("QuantizeAndDequantizeV3") axis); } else if (axis != -1) { ShapeHandle input; + if (axis >= kint32max) { + return errors::InvalidArgument( + "Axis cannot be >= kint32max value, got ", axis); + } TF_RETURN_IF_ERROR(c->WithRankAtLeast(c->input(0), axis + 1, &input)); DimensionHandle depth; TF_RETURN_IF_ERROR( diff --git a/tensorflow/python/kernel_tests/array_ops/array_ops_test.py b/tensorflow/python/kernel_tests/array_ops/array_ops_test.py index c14703cde39348..06c7dc43c4d5a1 100644 --- a/tensorflow/python/kernel_tests/array_ops/array_ops_test.py +++ b/tensorflow/python/kernel_tests/array_ops/array_ops_test.py @@ -1855,6 +1855,72 @@ def testOutOfBoundAxis(self): max_range=input_max, axis=2**31 - 1)) + @test_util.run_v2_only + def testInvalidAxis(self): + + @def_function.function + def test_quantize_and_dequantize_v2(): + gen_array_ops.quantize_and_dequantize_v2( + input=[2.5], + input_min=[1.0], + input_max=[10.0], + signed_input=True, + num_bits=1, + range_given=True, + round_mode="HALF_TO_EVEN", + narrow_range=True, + axis=0x7fffffff) + + @def_function.function + def test_quantize_and_dequantize_v3(): + gen_array_ops.quantize_and_dequantize_v3( + input=[2.5], + input_min=[1.0], + input_max=[10.0], + num_bits=1, + signed_input=True, + range_given=True, + narrow_range=True, + axis=0x7fffffff) + + @def_function.function + def test_quantize_and_dequantize_v4(): + gen_array_ops.quantize_and_dequantize_v4( + input=[2.5], + input_min=[1.0], + input_max=[10.0], + signed_input=True, + num_bits=1, + range_given=True, + round_mode="HALF_TO_EVEN", + narrow_range=True, + axis=0x7fffffff) + + @def_function.function + def test_quantize_and_dequantize_v4_grad(): + gen_array_ops.quantize_and_dequantize_v4_grad( + gradients=[2.5], + input=[2.5], + input_min=[1.0], + input_max=[10.0], + axis=0x7fffffff) + + with self.assertRaisesRegex( + ValueError, "Axis cannot be >= kint32max value, got 2147483647"): + test_quantize_and_dequantize_v2() + + with self.assertRaisesRegex( + ValueError, "Axis cannot be >= kint32max value, got 2147483647"): + test_quantize_and_dequantize_v3() + + with self.assertRaisesRegex( + ValueError, "Axis cannot be >= kint32max value, got 2147483647"): + test_quantize_and_dequantize_v4() + + with self.assertRaisesRegex( + ValueError, "Axis cannot be >= kint32max value, got 2147483647"): + test_quantize_and_dequantize_v4_grad() + @test_util.run_all_in_graph_and_eager_modes class SortedSearchTest(test_util.TensorFlowTestCase): From a6853af17dd5d90265988432c7cd4f07c3707773 Mon Sep 17 00:00:00 2001 From: Edward Schwartz Date: Wed, 7 Dec 2022 13:25:19 -0800 Subject: [PATCH 13/52] Add overflow test for num_results*dim to sobol_sample op PiperOrigin-RevId: 493692010 --- tensorflow/core/kernels/sobol_op.cc | 4 ++++ tensorflow/python/ops/sobol_ops_test.py | 10 ++++++++++ 2 files changed, 14 insertions(+) diff --git a/tensorflow/core/kernels/sobol_op.cc b/tensorflow/core/kernels/sobol_op.cc index 484bac99463300..144257bda5496d 100644 --- a/tensorflow/core/kernels/sobol_op.cc +++ b/tensorflow/core/kernels/sobol_op.cc @@ -158,6 +158,10 @@ class SobolSampleOp : public OpKernel { num_results < std::numeric_limits::max() - skip, errors::InvalidArgument("num_results+skip must be less than ", std::numeric_limits::max())); + OP_REQUIRES(context, + num_results < std::numeric_limits::max() / dim, + errors::InvalidArgument("num_results*dim must be less than ", + std::numeric_limits::max())); Tensor* output = nullptr; OP_REQUIRES_OK(context, diff --git a/tensorflow/python/ops/sobol_ops_test.py b/tensorflow/python/ops/sobol_ops_test.py index 24abf790704ae0..a795cbdcf4f06e 100644 --- a/tensorflow/python/ops/sobol_ops_test.py +++ b/tensorflow/python/ops/sobol_ops_test.py @@ -139,5 +139,15 @@ def test_non_scalar_input(self): num_results=constant_op.constant([1, 0]), skip=constant_op.constant([1]))) + @test_util.run_in_graph_and_eager_modes + def testDimNumResultsOverflow(self): + with self.assertRaisesRegex( + (ValueError, errors.InvalidArgumentError), + r'num_results\*dim must be less than 2147483647'): + self.evaluate( + gen_math_ops.sobol_sample( + dim=2560, num_results=16384000, skip=0, dtype=dtypes.float32)) + + if __name__ == '__main__': googletest.main() From 4c49ff8d07cb0851e63d6ee4a1e344a0922cee50 Mon Sep 17 00:00:00 2001 From: Antonio Sanchez Date: Tue, 3 Jan 2023 14:22:31 -0800 Subject: [PATCH 14/52] Fix sparse tensor to CSR batch index OOB error. Added a check for the batch index. PiperOrigin-RevId: 499315915 --- tensorflow/core/kernels/sparse/BUILD | 1 + tensorflow/core/kernels/sparse/kernels.cc | 6 +++++ .../core/kernels/sparse/kernels_test.cc | 25 +++++++++++++++++++ 3 files changed, 32 insertions(+) diff --git a/tensorflow/core/kernels/sparse/BUILD b/tensorflow/core/kernels/sparse/BUILD index 8be732152fa060..96aba2faef9e78 100644 --- a/tensorflow/core/kernels/sparse/BUILD +++ b/tensorflow/core/kernels/sparse/BUILD @@ -106,6 +106,7 @@ tf_cc_test( "//tensorflow/core:lib", "//tensorflow/core:test", "//tensorflow/core:testlib", + "//tensorflow/core/platform:status_matchers", "//third_party/eigen3", ], ) diff --git a/tensorflow/core/kernels/sparse/kernels.cc b/tensorflow/core/kernels/sparse/kernels.cc index 4cf375cc3172d0..5be8a0dba3fe3d 100644 --- a/tensorflow/core/kernels/sparse/kernels.cc +++ b/tensorflow/core/kernels/sparse/kernels.cc @@ -75,6 +75,12 @@ Status SparseTensorToCSRSparseMatrixCPUFunctor::operator()( } else { // rank == 3 for (int64_t i = 0; i < total_nnz; ++i) { const int cur_batch = indices(i, 0); + if (cur_batch < 0 || cur_batch >= batch_size) { + return errors::InvalidArgument("Batch index ", cur_batch, + " is outside of valid batch range [", 0, + ", ", batch_size, ")"); + } + // For now, the rows pointers store the corresponding row counts. csr_row_ptr(cur_batch * (num_rows + 1) + indices(i, 1) + 1) += 1; csr_col_ind(i) = indices(i, 2); diff --git a/tensorflow/core/kernels/sparse/kernels_test.cc b/tensorflow/core/kernels/sparse/kernels_test.cc index 39ab6fa011b78a..d7aee43921575f 100644 --- a/tensorflow/core/kernels/sparse/kernels_test.cc +++ b/tensorflow/core/kernels/sparse/kernels_test.cc @@ -22,6 +22,8 @@ limitations under the License. #include "tensorflow/core/framework/tensor_types.h" #include "tensorflow/core/lib/core/status_test_util.h" #include "tensorflow/core/platform/test.h" +#include "tensorflow/tsl/platform/errors.h" +#include "tensorflow/tsl/platform/status_matchers.h" namespace tensorflow { namespace { @@ -73,6 +75,29 @@ TEST(SparseTensorToCSRSparseMatrix, BatchConversion) { test::ExpectTensorEqual(csr_col_ind, test::AsTensor({0, 3, 1})); } +TEST(SparseTensorToCSRSparseMatrix, InvalidBatchThrowsIllegalArgument) { + const auto indices = + test::AsTensor({0, 0, 0, // + 4, 2, 3, // Batch out of bounds. + 2, 0, 1}, + TensorShape({3, 3})); + Tensor batch_ptr(DT_INT32, {4}); + Tensor csr_col_ind(DT_INT32, {3}); + // row pointers have size = batch_size * (num_rows + 1) = 3 * 4 = 12 + Tensor csr_row_ptr(DT_INT32, {12}); + test::FillFn(&csr_row_ptr, [](int unused) { return 0; }); + + functor::SparseTensorToCSRSparseMatrixCPUFunctor coo_to_csr; + EXPECT_THAT( + coo_to_csr(3 /* batch_size */, 3 /* num_rows */, + indices.template matrix(), batch_ptr.vec(), + csr_row_ptr.vec(), csr_col_ind.vec()), + tsl::testing::StatusIs( + tsl::error::Code::INVALID_ARGUMENT, + ::testing::ContainsRegex( + "Batch index .* is outside of valid batch range"))); +} + } // namespace } // namespace tensorflow From 7ba601b3ffa63c89826815487b9fd036746723e0 Mon Sep 17 00:00:00 2001 From: Zhufeng Pan Date: Thu, 12 Jan 2023 13:26:50 -0800 Subject: [PATCH 15/52] Fix security vulnerability with FractionalMax(AVG)Pool with illegal pooling_ratio PiperOrigin-RevId: 501651261 --- tensorflow/core/kernels/fractional_avg_pool_op.cc | 2 +- tensorflow/core/kernels/fractional_max_pool_op.cc | 2 +- .../nn_ops/fractional_avg_pool_op_test.py | 12 +++++++++++- .../nn_ops/fractional_max_pool_op_test.py | 12 +++++++++++- 4 files changed, 24 insertions(+), 4 deletions(-) diff --git a/tensorflow/core/kernels/fractional_avg_pool_op.cc b/tensorflow/core/kernels/fractional_avg_pool_op.cc index 3bb206866089a0..82093255c2cb3a 100644 --- a/tensorflow/core/kernels/fractional_avg_pool_op.cc +++ b/tensorflow/core/kernels/fractional_avg_pool_op.cc @@ -51,7 +51,7 @@ class FractionalAvgPoolOp : public OpKernel { pooling_ratio_[i])); } OP_REQUIRES( - context, pooling_ratio_[0] == 1 || pooling_ratio_[3] == 1, + context, pooling_ratio_[0] == 1 && pooling_ratio_[3] == 1, errors::Unimplemented("Fractional average pooling is not yet " "supported on the batch nor channel dimension.")); OP_REQUIRES_OK(context, context->GetAttr("deterministic", &deterministic_)); diff --git a/tensorflow/core/kernels/fractional_max_pool_op.cc b/tensorflow/core/kernels/fractional_max_pool_op.cc index ec08b5c5028727..1e9238da3bd0f4 100644 --- a/tensorflow/core/kernels/fractional_max_pool_op.cc +++ b/tensorflow/core/kernels/fractional_max_pool_op.cc @@ -53,7 +53,7 @@ class FractionalMaxPoolOp : public OpKernel { } OP_REQUIRES( - context, pooling_ratio_[0] == 1 || pooling_ratio_[3] == 1, + context, pooling_ratio_[0] == 1 && pooling_ratio_[3] == 1, errors::Unimplemented("Fractional max pooling is not yet " "supported on the batch nor channel dimension.")); diff --git a/tensorflow/python/kernel_tests/nn_ops/fractional_avg_pool_op_test.py b/tensorflow/python/kernel_tests/nn_ops/fractional_avg_pool_op_test.py index 59b20de84b4c4e..21c48dd51b9e74 100644 --- a/tensorflow/python/kernel_tests/nn_ops/fractional_avg_pool_op_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/fractional_avg_pool_op_test.py @@ -351,7 +351,7 @@ def testPoolingRatioHasMoreDimThanInput(self): name=None) self.evaluate(result) - def testPoolingRatioValueOutOfRange(self): + def testPoolingRatioIllegalSmallValue(self): with self.cached_session() as _: # Whether turn on `TF2_BEHAVIOR` generates different error messages with self.assertRaisesRegex( @@ -368,6 +368,16 @@ def testPoolingRatioValueOutOfRange(self): ) self.evaluate(result) + def testPoolingIllegalRatioForBatch(self): + with self.cached_session() as _: + with self.assertRaises(errors.UnimplementedError): + result = nn_ops.gen_nn_ops.fractional_avg_pool( + np.zeros([3, 30, 50, 3]), + [2, 3, 1.5, 1], + True, + True) + self.evaluate(result) + class FractionalAvgPoolGradTest(test.TestCase): """Tests for FractionalAvgPoolGrad. diff --git a/tensorflow/python/kernel_tests/nn_ops/fractional_max_pool_op_test.py b/tensorflow/python/kernel_tests/nn_ops/fractional_max_pool_op_test.py index 9102973fa13b50..d1c7c3680e3cef 100644 --- a/tensorflow/python/kernel_tests/nn_ops/fractional_max_pool_op_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/fractional_max_pool_op_test.py @@ -338,7 +338,7 @@ def testPoolingRatioHasMoreDimThanInput(self): name=None) self.evaluate(result) - def testPoolingRatioValueOutOfRange(self): + def testPoolingRatioIllegalSmallValue(self): with self.cached_session() as _: # Whether turn on `TF2_BEHAVIOR` generates different error messages with self.assertRaisesRegex( @@ -355,6 +355,16 @@ def testPoolingRatioValueOutOfRange(self): ) self.evaluate(result) + def testPoolingIllegalRatioForBatch(self): + with self.cached_session() as _: + with self.assertRaises(errors.UnimplementedError): + result = nn_ops.fractional_max_pool( + np.zeros([3, 30, 50, 3]), + [2, 3, 1.5, 1], + True, + True) + self.evaluate(result) + class FractionalMaxPoolGradTest(test.TestCase): """Tests for FractionalMaxPoolGrad. From a9c1da2f4efc6d6871c19b4c0ba720eeded652a4 Mon Sep 17 00:00:00 2001 From: Mihai Maruseac Date: Thu, 19 Jan 2023 17:31:46 -0800 Subject: [PATCH 16/52] Quick fix for a vuln in printing empty tensors. PiperOrigin-RevId: 503308327 --- tensorflow/core/framework/tensor.cc | 3 +++ 1 file changed, 3 insertions(+) diff --git a/tensorflow/core/framework/tensor.cc b/tensorflow/core/framework/tensor.cc index 74f1a1bd03e8af..d41366e1f62dc2 100644 --- a/tensorflow/core/framework/tensor.cc +++ b/tensorflow/core/framework/tensor.cc @@ -1224,6 +1224,9 @@ template <> string SummarizeArray(int64_t limit, int64_t num_elts, const TensorShape& tensor_shape, const char* data, const bool print_v2) { + if (data == nullptr) { + return strings::StrCat(""); // we already print type and shape + } // We first convert all chars to be 0/1 to not get InvalidEnumValue sanitizer // error auto mutable_data = std::unique_ptr(new char[num_elts]); From 3d7a3e04394a77350a248a4498e3c8101105dcc9 Mon Sep 17 00:00:00 2001 From: Haoliang Zhang Date: Thu, 19 Jan 2023 14:30:16 -0800 Subject: [PATCH 17/52] Check filter_input_channel > 0 in conv kernel. PiperOrigin-RevId: 503266928 --- tensorflow/lite/kernels/conv.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/tensorflow/lite/kernels/conv.cc b/tensorflow/lite/kernels/conv.cc index 9e115bdcbc302c..52ea12832115c6 100644 --- a/tensorflow/lite/kernels/conv.cc +++ b/tensorflow/lite/kernels/conv.cc @@ -347,6 +347,7 @@ TfLiteStatus Prepare(KernelType kernel_type, TfLiteContext* context, // or equals (normal conv). auto input_channel = input->dims->data[3]; auto filter_input_channel = filter->dims->data[3]; + TF_LITE_ENSURE(context, filter_input_channel > 0); TF_LITE_ENSURE_EQ(context, input_channel % filter_input_channel, 0); data->groups = input_channel / filter_input_channel; From 5492d0394cdd0f7020f6b7054bc916617329b706 Mon Sep 17 00:00:00 2001 From: Andrei Pikas Date: Sat, 15 Oct 2022 22:26:47 +0300 Subject: [PATCH 18/52] Fix integer overflow for multiframe gifs. --- .../core/kernels/image/decode_image_op.cc | 4 ++-- tensorflow/core/lib/gif/gif_io.cc | 14 +++++++------- tensorflow/core/lib/gif/gif_io_test.cc | 5 +++-- .../core/lib/gif/testdata/3g_multiframe.gif | Bin 0 -> 22703 bytes tensorflow/core/lib/gif/testdata/BUILD | 1 + 5 files changed, 13 insertions(+), 11 deletions(-) create mode 100644 tensorflow/core/lib/gif/testdata/3g_multiframe.gif diff --git a/tensorflow/core/kernels/image/decode_image_op.cc b/tensorflow/core/kernels/image/decode_image_op.cc index 91137f8343c834..baa4407fd84dda 100644 --- a/tensorflow/core/kernels/image/decode_image_op.cc +++ b/tensorflow/core/kernels/image/decode_image_op.cc @@ -452,12 +452,12 @@ class DecodeImageV2Op : public OpKernel { // allocation til after dtype conversion is done. `gif`::Decode` supports // uint8 only. Tensor* output = nullptr; - int buffer_size = 0; + ptrdiff_t buffer_size = 0; string error_string; uint8* buffer = gif::Decode( input.data(), input.size(), [&](int num_frames, int width, int height, int channels) -> uint8* { - buffer_size = num_frames * height * width * channels; + buffer_size = ptrdiff_t(num_frames) * height * width * channels; Status status; // By the existing API, we support decoding GIF with `decode_jpeg` or diff --git a/tensorflow/core/lib/gif/gif_io.cc b/tensorflow/core/lib/gif/gif_io.cc index ba4aa1156db83c..dddb8708c5e02b 100644 --- a/tensorflow/core/lib/gif/gif_io.cc +++ b/tensorflow/core/lib/gif/gif_io.cc @@ -105,7 +105,7 @@ uint8* Decode(const void* srcdata, int datasize, uint8* const dstdata = allocate_output(target_num_frames, width, height, channel); if (!dstdata) return nullptr; - for (int k = 0; k < target_num_frames; k++) { + for (ptrdiff_t k = 0; k < target_num_frames; k++) { uint8* this_dst = dstdata + k * width * channel * height; SavedImage* this_image = &gif_file->SavedImages[k]; @@ -125,10 +125,10 @@ uint8* Decode(const void* srcdata, int datasize, if (k > 0) { uint8* last_dst = dstdata + (k - 1) * width * channel * height; - for (int i = 0; i < height; ++i) { + for (ptrdiff_t i = 0; i < height; ++i) { uint8* p_dst = this_dst + i * width * channel; uint8* l_dst = last_dst + i * width * channel; - for (int j = 0; j < width; ++j) { + for (ptrdiff_t j = 0; j < width; ++j) { p_dst[j * channel + 0] = l_dst[j * channel + 0]; p_dst[j * channel + 1] = l_dst[j * channel + 1]; p_dst[j * channel + 2] = l_dst[j * channel + 2]; @@ -141,9 +141,9 @@ uint8* Decode(const void* srcdata, int datasize, // If the first frame does not fill the entire canvas then fill the // unoccupied canvas with zeros (black). if (k == 0) { - for (int i = 0; i < height; ++i) { + for (ptrdiff_t i = 0; i < height; ++i) { uint8* p_dst = this_dst + i * width * channel; - for (int j = 0; j < width; ++j) { + for (ptrdiff_t j = 0; j < width; ++j) { p_dst[j * channel + 0] = 0; p_dst[j * channel + 1] = 0; p_dst[j * channel + 2] = 0; @@ -165,9 +165,9 @@ uint8* Decode(const void* srcdata, int datasize, return nullptr; } - for (int i = imgTop; i < imgBottom; ++i) { + for (ptrdiff_t i = imgTop; i < imgBottom; ++i) { uint8* p_dst = this_dst + i * width * channel; - for (int j = imgLeft; j < imgRight; ++j) { + for (ptrdiff_t j = imgLeft; j < imgRight; ++j) { GifByteType color_index = this_image->RasterBits[(i - img_desc->Top) * (img_desc->Width) + (j - img_desc->Left)]; diff --git a/tensorflow/core/lib/gif/gif_io_test.cc b/tensorflow/core/lib/gif/gif_io_test.cc index 38c181911697c2..32e63480720663 100644 --- a/tensorflow/core/lib/gif/gif_io_test.cc +++ b/tensorflow/core/lib/gif/gif_io_test.cc @@ -52,7 +52,7 @@ void TestDecodeGif(Env* env, DecodeGifTestCase testcase) { w = width; h = height; c = channels; - return new uint8[frame_cnt * height * width * channels]; + return new uint8[ptrdiff_t(frame_cnt) * height * width * channels]; }, &error_string)); ASSERT_NE(imgdata, nullptr); @@ -72,7 +72,8 @@ TEST(GifTest, Gif) { {testdata_path + "optimized.gif", 12, 20, 40, 3}, {testdata_path + "red_black.gif", 1, 16, 16, 3}, {testdata_path + "scan.gif", 12, 20, 40, 3}, - {testdata_path + "squares.gif", 2, 16, 16, 3}}); + {testdata_path + "squares.gif", 2, 16, 16, 3}, + {testdata_path + "3g_multiframe.gif", 519, 1920, 1080, 3}}); for (const auto& tc : testcases) { TestDecodeGif(env, tc); diff --git a/tensorflow/core/lib/gif/testdata/3g_multiframe.gif b/tensorflow/core/lib/gif/testdata/3g_multiframe.gif new file mode 100644 index 0000000000000000000000000000000000000000..f9286dc85c0084b3125e4ff7d605a695dbb8a5cd GIT binary patch literal 22703 zcmeI1Sxj4J7RO180!j%%Z4@j>Q@SV+LI_Hu;DM@Qx_|*MD2;+e(v(S~KnP7yTLp_$ zmDy|qwxLN1E)YV1U}HA3*%zBF8#Z3ah8-ZptTtxz?(2PmI+bRm>HA3Y-Ix3F={_8t z`}_UR`I3?ok$?LK(GTfM_VMw0y*}Up75IOx!28~}uwS0~>hr&+eR(W8=FJ_i>+rta z*Js!3chq)Gwto{Jf3MFD|NSUVRI6TNpOH8J0c>y!IhWp&6Ob+Z$FOvUYl;$82i^s^o}qi#fgE$(l7 z%1mocxs@`&e%_An%XujtXjgV;J}7@TmD};6FZ(I`YYDe=hL^V}zLm=BdMO~7bT1{m z?pLG4H{UvdhB)&NDDNCTCxvHlPQ4mnDpkFV9H_`!tHi&npgK`2DL-SzooE zlko@El$>uR4HyA;V+Q(hZB8cOZM_;Re(OE#5M)>s-8lS>d}z?5nQTtIjUDD&)>M6$ zU&)6Bwk`F83h#Tkks-TD`?UQV#mKPJreBoY#*L1+ov>-`m11=CHAHqssENM&vekwD z2U+T(fD>AEaZqB!e97L^+4H4{i>!Il{;S$~au_LMf$~vx_Ci@iD{G-VimP3yh#8Mq zq#mEjUZf#cSc~*etlGuOgq??$s=nBFY3Wzg!M3IA=o70;H_{RhE#EwK`qJ`m_=|1J z)n~7+F4ts|4y`aQR9{;8J-fASr8bwlx>A=nen?aQ)6^x+Ey7Bhrr~Gns^&IvXQZ~V zWM7W<4&@+Q+eAH~(>7NnMy|HpJe{+8mvNE3+FE~Ax5{iJMe0~B)j7I0Rx4Y_?%?Wl z?VNFMbg+HBS9)S?y>Bw{@J9c$)0Z~} zR2SPfxU*N+Hh5~%;Z10%`ts(WwzYkezs_CT6zInfZw;BJE^iIPEA3k&F6-LXsE=Qi ze#|#GSAWkxyhDFKAa-5uHg|PtHUt9KY!gI3?oMwMIY7V8Yd!{9Y%2! zZ`~+~5k{G$$CbGz8B)_>l7C`bHz^YQj+iIE2>#Ce7!}@Weu9qOFh5O8I%0WtD*Zdl zb9`2(W$JAHhDDi4K4Mi}sQJ$NBAeN1ozCTLSZDHtM{Kh{DZjJ5BxpKqb3fa>8@5+O zzi3!p5_|=or-XOG3)I+6c(E!e+P-u%{fd2=k=12isn6fEYZ}SX4sA=#6~`)z+2zo6 z@HQQ59AUI`yEF98 zs;q9ebvA#?ZBvtvdf=s+?>%-cv)kiX=WThMdf`#8%dGs~>xMPmUXRPR^|}KC8h!jw zV{e4$iT>!uH-bfDJLB|)LDWXyaMZnB8TulGv~g#w=w4ufzBnx4j(-yB{@bk$7&7KY;<#oQZ#;oZ=mH-oA+x_!nkPz zogi&KU=s;Hbr>p%0WD#E$)e9fj8znL%Ll;|qQp4kuT|8RkHV8DQZkI!8Pb-6u@e*6 z0^^OwfV&Y%$>Ni>#+xkk-N^I_@oB#CHxBh~R93PiW7=2^N$*DIPe{%=j5T8ctuf?e z>G=>7LxgTURx=^J7-#x@lG=KlnJml6Fx9H0ttWUBvMU9qI&}aODNL6CP;08!qM32Z z3HepN>6V_#{6v$iD3~@iz*6R?wh6D|7l-M#&u&(NU&>@rh`G@p!}=^(JV}Z(-wC3z zz6eiwT%KWWLdaN&vEs+{0&{cN?lx3X%9HE0=9UOdTS~h4Nj2YmH-^@R&PsV&J8f=7 z%G$8`;-?J`Gc#c~JB^(3tSQ98LSfh^YsAl5<1B4yH1;WG%5!#xg^ic7PxHjjy9zAr znY-KZ!j!4rT1!VZraeO`p5pQ?oq4qOvzinof7;STkhPz)iIpP`OE+VAlo zL&0>M50?~DgOjwbpO~pH zbs09kO4fCiCwaM9U=yf!cN2uEbH-ZRkQUQjpp?v6`LMO^zZ3LEe|6-H8 z@;GdxK7kyfA6orJC_Ltm`|3SwY?qaQ~5S*T-@g6(e#mT42?9gnN%{Vm~`)s!=iCk%Oi zYpiS)d(H8*F>rvDgwdU>b39{V2iWN{-D!d2Ifp*bk%d{ym~l)&@`0{=+1fd$Lpc`6 z<&ZJ!=R=(;5tiFqBU`^1?|d;y=k_x(8#!m3(<(Wa%ad(fx#pZv2l5~xX7h(S=d2dX z<11yGR|U?OdOB}NgV`#Wan8YV-iS@M^^4Q_%4ZKW=7-f6g}T)KIOu+`Tu+L3%?DLN z55uvB@-walgaR6ml^f{ST#I3Q21Q9&cd<61!~ z2Icv3Q-jl`N!Y`mBxB7@p>8b-$A3~IH@C*SSJNu_&zM*X`;1$MSMaBJa!c1W_gdy2 zfl7$A_SU)AvvGoHrQFIDxHs}D1+yBgjX&ewBq#)PHo0xY>E0sl8B+V<;QOH-Jq0(k z5Uha5<2{C|%AuujoLzjzV`M0XR$>))`8AKJanGs0R<^}<4&$5`ieX>p?trNUzsc-{KSQI`hibQ+Pej*Azw&-#`+5 z@x+}3qCbPU3nB)nh=DLM2vPVpvT!fHa37%%!6*!Y3iqoD55R?Ch@uaWMIYgd4ibtY z7)6m#QIx7E8ZL@K6dywtAIBG;AQU4R#c@#aC#vF4;o<~D$!ExtFYqOagc1~^Bn2u# zt4gqNNgAT`B(n4rzVtMq6wfHlfJ)D*O3%TinF!K(B!$wUOD2uUfyQ%D2~g+VEYC{z`NUI$aE5M|eqWjFC< z)r2wzqpTJxt5=mZz-5hy@+M??3%bS;CfgXrrj`X)@*BPxx^%3eg!f9BO4K0AC0w}b1SkH4?a4qxB@a_HYD zzrQ^6=U2Ahe|ZRe2~^yt5Oe|P0?-A3zS{xw}BSOD_)>mC3`*e(H>r9dtK zT>w}BSOARhhZKPkfDwQZfDwQZfDwQZeEui_*d2mN5$q1ZzF=F5fCYdBpbG#C01Lo9 z0!WAd3dmq<1abk`8i91UEk(crzyi<(fCYdBV1xWeKn6wtMgT?tMgT?tMgT?tcZ#p` z0N5S=Q37zA1oj2sHVJHvwxbbX0bl{>0>A>m0+0?tIs`MqpJNip1t1rITmW(b$OZo+ Hxgh?Z>XK=- literal 0 HcmV?d00001 diff --git a/tensorflow/core/lib/gif/testdata/BUILD b/tensorflow/core/lib/gif/testdata/BUILD index 540482ab5accbf..888b5f3adb5210 100644 --- a/tensorflow/core/lib/gif/testdata/BUILD +++ b/tensorflow/core/lib/gif/testdata/BUILD @@ -15,6 +15,7 @@ filegroup( "scan.gif", "red_black.gif", "squares.gif", + "3g_multiframe.gif", "pendulum_sm.gif", # Add groundtruth frames for `pendulum_sm.gif`. # PNG format because it's lossless. From 6b48e16c5311a4be90a0820286c39bd1a5874efe Mon Sep 17 00:00:00 2001 From: Mason Chang Date: Wed, 1 Feb 2023 14:09:32 -0800 Subject: [PATCH 19/52] Add out of bounds array check to dynamic_stitch_op. PiperOrigin-RevId: 506418249 --- .../tf2xla/kernels/dynamic_stitch_op.cc | 4 ++++ tensorflow/core/kernels/dynamic_stitch_op.cc | 17 +++++++++++------ .../data_structures/dynamic_stitch_op_test.py | 13 +++++++++++++ 3 files changed, 28 insertions(+), 6 deletions(-) diff --git a/tensorflow/compiler/tf2xla/kernels/dynamic_stitch_op.cc b/tensorflow/compiler/tf2xla/kernels/dynamic_stitch_op.cc index 7566cc79742d08..cd03b617158e7d 100644 --- a/tensorflow/compiler/tf2xla/kernels/dynamic_stitch_op.cc +++ b/tensorflow/compiler/tf2xla/kernels/dynamic_stitch_op.cc @@ -146,6 +146,10 @@ class DynamicStitchOp : public XlaOpKernel { for (int input_num = 0; input_num < indices.size(); input_num++) { for (int i = 0; i < indices[input_num].shape().dimensions(0); ++i) { int index = indices[input_num].Get({i}); + OP_REQUIRES( + ctx, index >= 0, + errors::InvalidArgument("indices[", index, "] is out of range")); + src_input_vector[index] = input_num; src_slice_vector[index] = i; if (!src_index_used[index]) { diff --git a/tensorflow/core/kernels/dynamic_stitch_op.cc b/tensorflow/core/kernels/dynamic_stitch_op.cc index 50cb3acbb031f9..9b405a80667a8c 100644 --- a/tensorflow/core/kernels/dynamic_stitch_op.cc +++ b/tensorflow/core/kernels/dynamic_stitch_op.cc @@ -97,6 +97,17 @@ class DynamicStitchOpImplBase : public OpKernel { *first_dim_size = max_index + 1; + for (const Tensor& indices : *indices_inputs) { + auto indices_vec = indices.flat(); + + for (int i = 0; i < indices_vec.size(); i++) { + int32_t index = internal::SubtleMustCopy(indices_vec(i)); + OP_REQUIRES( + c, FastBoundsCheck(index, *first_dim_size), + errors::InvalidArgument("indices[", i, "] is out of range")); + } + } + // Validate that data[i].shape = indices[i].shape + constant OP_REQUIRES_OK(c, c->input_list("data", data_inputs)); const Tensor& data0 = (*data_inputs)[0]; @@ -265,9 +276,6 @@ class DynamicStitchOpImplCPU : public DynamicStitchOpImplBase { const T* data_base = data_flat.data(); for (int i = 0; i < indices_vec.size(); i++) { int32_t index = internal::SubtleMustCopy(indices_vec(i)); - OP_REQUIRES( - c, FastBoundsCheck(index, first_dim_size), - errors::InvalidArgument("indices[", i, "] is out of range")); memcpy(merged_base + index * slice_size, data_base + i * slice_size, slice_bytes); } @@ -277,9 +285,6 @@ class DynamicStitchOpImplCPU : public DynamicStitchOpImplBase { // Copy slice data[i] to merged[indices[i]] Eigen::DSizes data_indices(i, 0); int32_t index = internal::SubtleMustCopy(indices_vec(i)); - OP_REQUIRES( - c, FastBoundsCheck(index, first_dim_size), - errors::InvalidArgument("indices[", i, "] is out of range")); Eigen::DSizes merged_indices(index, 0); merged_flat.slice(merged_indices, sizes) = data_flat.slice(data_indices, sizes); diff --git a/tensorflow/python/kernel_tests/data_structures/dynamic_stitch_op_test.py b/tensorflow/python/kernel_tests/data_structures/dynamic_stitch_op_test.py index 05e0cb6be2925a..8a90267f882f8e 100644 --- a/tensorflow/python/kernel_tests/data_structures/dynamic_stitch_op_test.py +++ b/tensorflow/python/kernel_tests/data_structures/dynamic_stitch_op_test.py @@ -221,6 +221,19 @@ def testErrorDataAndIndicesSizeMismatch(self): with self.assertRaises(ValueError): self.stitch_op(indices, data) + def testOutOfBoundsIndexRaisesInvalidArgument(self): + with self.assertRaisesRegex(errors.InvalidArgumentError, "out of range"): + indices = [[-1000], [405], [519], [758], [1015]] + data = [ + [110.27793884277344], + [120.29475402832031], + [157.2418212890625], + [157.2626953125], + [188.45382690429688], + ] + + self.evaluate(self.stitch_op(indices, data)) + class DynamicStitchTest(DynamicStitchTestBase, test.TestCase): From 372ebce2b88c768f146c2b93ba07b06ef7496f15 Mon Sep 17 00:00:00 2001 From: Antonio Sanchez Date: Tue, 17 Jan 2023 11:24:45 -0800 Subject: [PATCH 20/52] Fix histogram overflow vulnerability. Changed order of cast and subtraction to avoid overflows. PiperOrigin-RevId: 502640700 --- tensorflow/core/kernels/histogram_op.cc | 20 +++++++++++--------- tensorflow/python/ops/histogram_ops_test.py | 13 +++++++++++++ 2 files changed, 24 insertions(+), 9 deletions(-) diff --git a/tensorflow/core/kernels/histogram_op.cc b/tensorflow/core/kernels/histogram_op.cc index 59d23c60ffb59c..1e8bf7e04fc007 100644 --- a/tensorflow/core/kernels/histogram_op.cc +++ b/tensorflow/core/kernels/histogram_op.cc @@ -64,15 +64,17 @@ struct HistogramFixedWidthFunctor { // step = (b - a) / nbins // (x - a) / step // , then the entries are mapped to output. - - // Bug fix: Switch the order of cwiseMin and int32-casting to avoid - // producing a negative index when casting an big int64 number to int32 - index_to_bin.device(d) = - ((values.cwiseMax(value_range(0)) - values.constant(value_range(0))) - .template cast() / - step) - .cwiseMin(nbins_minus_1) - .template cast(); + // + // Bound range and cast to double _before_ subtracting the lower-bound to + // avoid overflow. Otherwise the difference may not fit within the type + // (e.g. int32). + index_to_bin.device(d) = ((values.cwiseMax(value_range(0)) + .cwiseMin(value_range(1)) + .template cast() - + static_cast(value_range(0))) / + step) + .cwiseMin(nbins_minus_1) + .template cast(); out.setZero(); for (int32_t i = 0; i < index_to_bin.size(); i++) { diff --git a/tensorflow/python/ops/histogram_ops_test.py b/tensorflow/python/ops/histogram_ops_test.py index cc17ed0a032c88..629507052dd206 100644 --- a/tensorflow/python/ops/histogram_ops_test.py +++ b/tensorflow/python/ops/histogram_ops_test.py @@ -176,5 +176,18 @@ def test_shape_inference(self): self.assertAllClose(expected_bin_counts, hist.eval({placeholder: 5})) + def test_large_range(self): + hist = histogram_ops.histogram_fixed_width( + values=constant_op.constant( + [-(2**31), 2**31 - 1], dtype=dtypes.int32 + ), + value_range=constant_op.constant( + [-(2**31), 2**31 - 1], dtype=dtypes.int32 + ), + nbins=2, + ) + self.assertAllEqual(hist, [1, 1]) + + if __name__ == '__main__': test.main() From a392b291b603b09d10feb0e6f8f2153bbdc9ab58 Mon Sep 17 00:00:00 2001 From: Roman Dzhabarov Date: Thu, 26 Jan 2023 12:34:51 -0800 Subject: [PATCH 21/52] [TFG] Fix index out of bound issue in TFG shape inference pass. Ensure that there is at least a single non control input to the Identity/IdentityN. PiperOrigin-RevId: 504913650 --- tensorflow/core/transforms/shape_inference/pass.cc | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/tensorflow/core/transforms/shape_inference/pass.cc b/tensorflow/core/transforms/shape_inference/pass.cc index 783eb60d9b8c29..a61edd7faa5d08 100644 --- a/tensorflow/core/transforms/shape_inference/pass.cc +++ b/tensorflow/core/transforms/shape_inference/pass.cc @@ -101,9 +101,8 @@ void ShapeInference::TryToCacheResultsTensorValue(Operation *op) { if (op_name == "Const") { cached_tensor_values_[op->getResult(0)] = op->getAttrOfType("value"); - } else if (op_name == "Identity" || - (op_name == "IdentityN" && - TFOp(op).getNonControlOperands().size() == 1)) { + } else if ((op_name == "Identity" || op_name == "IdentityN") && + TFOp(op).getNonControlOperands().size() == 1) { DenseElementsAttr operand_tensor_value = GetTensorValue(op->getOperand(0)); if (!operand_tensor_value) return; cached_tensor_values_[op->getResult(0)] = operand_tensor_value; From b33ee6dfbee399c976994604bca6528a55c87b2f Mon Sep 17 00:00:00 2001 From: Jeremy Meredith Date: Tue, 24 Jan 2023 11:35:18 -0800 Subject: [PATCH 22/52] Fixing null pointer read in TensorArrayConcat when the step container is missing. PiperOrigin-RevId: 504332194 --- tensorflow/core/kernels/tensor_array_ops.cc | 5 +++-- .../data_structures/tensor_array_ops_test.py | 16 ++++++++++++++++ 2 files changed, 19 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/kernels/tensor_array_ops.cc b/tensorflow/core/kernels/tensor_array_ops.cc index 72accb86aa2419..ef3c78dd6471fc 100644 --- a/tensorflow/core/kernels/tensor_array_ops.cc +++ b/tensorflow/core/kernels/tensor_array_ops.cc @@ -80,8 +80,9 @@ Status GetTensorArray(OpKernelContext* ctx, TensorArray** tensor_array) { TF_RETURN_IF_ERROR(GetHandle(ctx, &container, &ta_handle)); ResourceMgr* rm = ctx->resource_manager(); if (rm == nullptr) return errors::Internal("No resource manager."); - TF_RETURN_IF_ERROR( - ctx->step_container()->Lookup(rm, container + ta_handle, tensor_array)); + ScopedStepContainer* sc = ctx->step_container(); + if (sc == nullptr) return errors::Internal("No step container."); + TF_RETURN_IF_ERROR(sc->Lookup(rm, container + ta_handle, tensor_array)); return OkStatus(); } else { return LookupResource(ctx, HandleFromInput(ctx, 0), tensor_array); diff --git a/tensorflow/python/kernel_tests/data_structures/tensor_array_ops_test.py b/tensorflow/python/kernel_tests/data_structures/tensor_array_ops_test.py index 509ebdad6ddaf1..d89e78f3a9524a 100644 --- a/tensorflow/python/kernel_tests/data_structures/tensor_array_ops_test.py +++ b/tensorflow/python/kernel_tests/data_structures/tensor_array_ops_test.py @@ -1833,6 +1833,22 @@ def testStackShapeOnStaticSize(self): ta = ta.write(0, [0]) self.assertEqual([42, 1], ta.stack().shape.as_list()) + def testTensorArrayConcatFailsWhenMissingStepContainer(self): + @def_function.function + def func(): + y = data_flow_ops.TensorArrayConcatV2( + handle=["a", "b"], + flow_in=0.1, + dtype=dtypes.int32, + element_shape_except0=1, + ) + return y + + with self.assertRaisesRegex( + errors.NotFoundError, "Container .* does not exist" + ): + self.evaluate(func()) + class TensorArrayBenchmark(test.Benchmark): From 2a6ee549013198a6e3096a4a19a2b3f3e6b0a7f0 Mon Sep 17 00:00:00 2001 From: Jian Cai Date: Wed, 1 Feb 2023 21:48:51 -0800 Subject: [PATCH 23/52] [Tensorflow] Fix security vulnerability with DenseBincountOp PiperOrigin-RevId: 506514542 --- tensorflow/compiler/tests/BUILD | 16 ++++++++ tensorflow/compiler/tests/bincount_op_test.py | 40 +++++++++++++++++++ .../compiler/tf2xla/kernels/bincount_op.cc | 17 ++++---- 3 files changed, 66 insertions(+), 7 deletions(-) create mode 100644 tensorflow/compiler/tests/bincount_op_test.py diff --git a/tensorflow/compiler/tests/BUILD b/tensorflow/compiler/tests/BUILD index 38b45f89f3a07a..5bdf389672cf3f 100644 --- a/tensorflow/compiler/tests/BUILD +++ b/tensorflow/compiler/tests/BUILD @@ -2344,3 +2344,19 @@ tf_xla_py_test( "//tensorflow/python:training", ], ) + +tf_xla_py_test( + name = "bincount_op_test", + size = "small", + srcs = ["bincount_op_test.py"], + enable_mlir_bridge = False, + python_version = "PY3", + shard_count = 10, + tags = [ + "no_pip", # TODO(b/149738646): fix pip install so these tests run on kokoro pip + ], + deps = [ + ":xla_test", + "//tensorflow/python:platform_test", + ], +) diff --git a/tensorflow/compiler/tests/bincount_op_test.py b/tensorflow/compiler/tests/bincount_op_test.py new file mode 100644 index 00000000000000..79e8a7e91b8a08 --- /dev/null +++ b/tensorflow/compiler/tests/bincount_op_test.py @@ -0,0 +1,40 @@ +# Copyright 2023 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. +# ============================================================================== +"""Tests for bincount using the XLA JIT.""" +from tensorflow.compiler.tests import xla_test +from tensorflow.python.framework import errors +from tensorflow.python.ops import gen_math_ops +from tensorflow.python.platform import googletest + + +class BincountTest(xla_test.XLATestCase): + + def testInputRank0(self): + with self.session(): + with self.test_scope(): + bincount = gen_math_ops.bincount(arr=6, size=804, weights=[52, 351]) + + with self.assertRaisesRegex( + errors.InvalidArgumentError, + ( + "`weights` must be the same shape as `arr` or a length-0" + " `Tensor`, in which case it acts as all weights equal to 1." + ), + ): + self.evaluate(bincount) + + +if __name__ == "__main__": + googletest.main() diff --git a/tensorflow/compiler/tf2xla/kernels/bincount_op.cc b/tensorflow/compiler/tf2xla/kernels/bincount_op.cc index ba95c77f61250a..5cac4d98c50a9e 100644 --- a/tensorflow/compiler/tf2xla/kernels/bincount_op.cc +++ b/tensorflow/compiler/tf2xla/kernels/bincount_op.cc @@ -62,21 +62,15 @@ class DenseBincountOp : public XlaOpKernel { StatusOr input_shape_or = ctx->builder()->GetShape(input); OP_REQUIRES_OK(ctx, input_shape_or.status()); auto input_shape = input_shape_or.value(); - auto size = input_shape.dimensions(0); - if (!size) { - output = xla::Broadcast(zero, {output_size}); - ctx->SetOutput(0, output); - return; - } auto rank = input_shape.rank(); OP_REQUIRES(ctx, rank <= 2, errors::InvalidArgument( "Shape must be at most rank 2 but is rank ", rank)); - xla::XlaOp weights = ctx->Input(2); StatusOr weights_shape_or = ctx->builder()->GetShape(weights); + OP_REQUIRES_OK(ctx, weights_shape_or.status()); auto weights_shape = weights_shape_or.value(); @@ -91,11 +85,20 @@ class DenseBincountOp : public XlaOpKernel { "1. Received ", weights_shape.DebugString())); + auto size = input_shape.dimensions(0); + + if (!size) { + output = xla::Broadcast(zero, {output_size}); + ctx->SetOutput(0, output); + return; + } + auto weights_size = weights_shape.dimensions(0); bool has_weights = false; if (weights_size) { has_weights = true; } + xla::Shape output_shape = xla::ShapeUtil::MakeShape(dtype, {output_size}); xla::ScatterDimensionNumbers scatter_dnums; scatter_dnums.set_index_vector_dim(1); From dc791a716d935ef6f2fbac0abc1d688b47f9acbb Mon Sep 17 00:00:00 2001 From: Isha Arkatkar Date: Tue, 24 Jan 2023 13:53:58 -0800 Subject: [PATCH 24/52] Fix security vulnerability in EditDistance op shape function. PiperOrigin-RevId: 504367470 --- tensorflow/core/ops/array_ops.cc | 14 ++++- .../array_ops/edit_distance_op_test.py | 63 ++++++++++++++++++- 2 files changed, 75 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/ops/array_ops.cc b/tensorflow/core/ops/array_ops.cc index b57a44f44fdd73..7a58ce26e04f43 100644 --- a/tensorflow/core/ops/array_ops.cc +++ b/tensorflow/core/ops/array_ops.cc @@ -24,6 +24,7 @@ limitations under the License. #include "tensorflow/core/framework/types.h" #include "tensorflow/core/framework/types.pb.h" #include "tensorflow/core/lib/core/errors.h" +#include "tensorflow/core/platform/status.h" #include "tensorflow/core/platform/types.h" #include "tensorflow/core/util/mirror_pad_mode.h" #include "tensorflow/core/util/padding.h" @@ -1071,13 +1072,24 @@ REGISTER_OP("EditDistance") // or else the output shape is unknown. return shape_inference::UnknownShape(c); } - if (hypothesis_shape_t->NumElements() != truth_shape_t->NumElements()) { return errors::InvalidArgument( "Num elements of hypothesis_shape does not match truth_shape: ", hypothesis_shape_t->NumElements(), " vs. ", truth_shape_t->NumElements()); } + if (hypothesis_shape_t->NumElements() < 2) { + return errors::InvalidArgument( + "Input Hypothesis SparseTensors must have rank at least 2, but " + "hypothesis_shape rank is: ", + hypothesis_shape_t->NumElements()); + } + if (truth_shape_t->NumElements() < 2) { + return errors::InvalidArgument( + "Input Truth SparseTensors must have rank at least 2, but " + "truth_shape rank is: ", + truth_shape_t->NumElements()); + } auto h_values = hypothesis_shape_t->flat(); auto t_values = truth_shape_t->flat(); diff --git a/tensorflow/python/kernel_tests/array_ops/edit_distance_op_test.py b/tensorflow/python/kernel_tests/array_ops/edit_distance_op_test.py index 8aa22cb812a0e6..3dc2f4565b7e0f 100644 --- a/tensorflow/python/kernel_tests/array_ops/edit_distance_op_test.py +++ b/tensorflow/python/kernel_tests/array_ops/edit_distance_op_test.py @@ -15,8 +15,9 @@ """Tests for tensorflow.kernels.edit_distance_op.""" import numpy as np - +from tensorflow.python.eager import def_function from tensorflow.python.framework import constant_op +from tensorflow.python.framework import errors from tensorflow.python.framework import ops from tensorflow.python.framework import sparse_tensor from tensorflow.python.ops import array_ops @@ -225,6 +226,66 @@ def testEditDistanceBadIndices(self): "to outside of the buffer for the output tensor|" r"Dimension -\d+ must be >= 0")) + def testEmptyShapeWithEditDistanceRaisesError(self): + para = { + "hypothesis_indices": [[]], + "hypothesis_values": ["tmp/"], + "hypothesis_shape": [], + "truth_indices": [[]], + "truth_values": [""], + "truth_shape": [], + "normalize": False, + } + + # Check edit distance raw op with empty shape in eager mode. + with self.assertRaisesRegex( + (errors.InvalidArgumentError, ValueError), + ( + r"Input Hypothesis SparseTensors must have rank at least 2, but" + " hypothesis_shape rank is: 0|Input SparseTensors must have rank " + "at least 2, but truth_shape rank is: 0" + ), + ): + array_ops.gen_array_ops.EditDistance(**para) + + # Check raw op with tf.function + @def_function.function + def TestFunction(): + """Wrapper function for edit distance call.""" + array_ops.gen_array_ops.EditDistance(**para) + + with self.assertRaisesRegex( + ValueError, + ( + "Input Hypothesis SparseTensors must have rank at least 2, but" + " hypothesis_shape rank is: 0" + ), + ): + TestFunction() + + # Check with python wrapper API + hypothesis_indices = [[]] + hypothesis_values = [0] + hypothesis_shape = [] + truth_indices = [[]] + truth_values = [1] + truth_shape = [] + expected_output = [] # dummy ignored + + with self.assertRaisesRegex( + ValueError, + ( + "Input Hypothesis SparseTensors must have rank at least 2, but" + " hypothesis_shape rank is: 0" + ), + ): + self._testEditDistance( + hypothesis=(hypothesis_indices, hypothesis_values, hypothesis_shape), + truth=(truth_indices, truth_values, truth_shape), + normalize=False, + expected_output=expected_output, + ) + if __name__ == "__main__": test.main() From 6096ce39bb008ee365b2b8efc205329be20f6ebd Mon Sep 17 00:00:00 2001 From: Jeremy Meredith Date: Thu, 26 Jan 2023 11:46:58 -0800 Subject: [PATCH 25/52] Check for unexpected scalars in the shape argument to ParallelConcat. PiperOrigin-RevId: 504901518 --- tensorflow/core/kernels/inplace_ops.cc | 2 +- tensorflow/core/ops/array_ops.cc | 7 +++++++ .../kernel_tests/array_ops/stack_op_test.py | 5 +++-- tensorflow/python/ops/array_ops_test.py | 15 +++++++++++++++ 4 files changed, 26 insertions(+), 3 deletions(-) diff --git a/tensorflow/core/kernels/inplace_ops.cc b/tensorflow/core/kernels/inplace_ops.cc index abda03164fd2d4..9b144a14f7e23f 100644 --- a/tensorflow/core/kernels/inplace_ops.cc +++ b/tensorflow/core/kernels/inplace_ops.cc @@ -78,7 +78,7 @@ class ParallelConcatUpdate : public OpKernel { OP_REQUIRES( ctx, value.dim_size(0) > loc_, errors::InvalidArgument("0th dimension of value = ", value.dim_size(0), - " is less than loc_=", loc_)); + " must be greater than loc_ = ", loc_)); auto update = ctx->input(1); diff --git a/tensorflow/core/ops/array_ops.cc b/tensorflow/core/ops/array_ops.cc index b57a44f44fdd73..debfff9452f242 100644 --- a/tensorflow/core/ops/array_ops.cc +++ b/tensorflow/core/ops/array_ops.cc @@ -15,6 +15,7 @@ limitations under the License. #include #include +#include #include "tensorflow/core/framework/common_shape_fns.h" #include "tensorflow/core/framework/kernel_shape_util.h" @@ -307,6 +308,12 @@ REGISTER_OP("ParallelConcat") return errors::InvalidArgument( "All input shapes must be fully defined."); } + if (c->Rank(c->input(i)) < 1) { + return errors::InvalidArgument( + "The rank of all input shapes must be greater than 0, " + "but input ", + i, " had rank ", c->Rank(c->input(i)), "."); + } DimensionHandle unused; if (!c->WithValue(c->Dim(c->input(i), 0), 1, &unused).ok()) { return errors::InvalidArgument("Size of first dimension must be 1."); diff --git a/tensorflow/python/kernel_tests/array_ops/stack_op_test.py b/tensorflow/python/kernel_tests/array_ops/stack_op_test.py index 378f5174bc7edc..28c1d65c782dd6 100644 --- a/tensorflow/python/kernel_tests/array_ops/stack_op_test.py +++ b/tensorflow/python/kernel_tests/array_ops/stack_op_test.py @@ -82,8 +82,9 @@ def f(): y = gen_array_ops.parallel_concat(values=[["tf"]], shape=0) return y - with self.assertRaisesRegex(errors.InvalidArgumentError, - r"0th dimension of value .* is less than"): + with self.assertRaisesRegex( + errors.InvalidArgumentError, r"0th dimension .* must be greater than" + ): f() def testSimpleParallelGPU(self): diff --git a/tensorflow/python/ops/array_ops_test.py b/tensorflow/python/ops/array_ops_test.py index 3f063a24a27cd2..0c82f5ac0987df 100644 --- a/tensorflow/python/ops/array_ops_test.py +++ b/tensorflow/python/ops/array_ops_test.py @@ -18,6 +18,7 @@ from tensorflow.python.eager import def_function from tensorflow.python.framework import dtypes from tensorflow.python.framework import tensor_spec +from tensorflow.python.framework import test_util from tensorflow.python.ops import array_ops from tensorflow.python.ops import math_ops from tensorflow.python.ops import random_ops @@ -91,6 +92,20 @@ def g(x): conc = g.get_concrete_function(tensor_spec.TensorSpec([10, None])) self.assertAllEqual(conc.output_shapes.as_list(), [10]) + @test_util.run_in_graph_and_eager_modes + def testParallelConcatFailsWithRankZeroShape(self): + op = array_ops.ParallelConcat + para = {"shape": 0, "values": [1]} + + def func(): + y = op(**para) + return y + + with self.assertRaisesRegex( + Exception, "(rank|dimension) of .* must be greater than .* 0" + ): + func() + if __name__ == "__main__": test.main() From d6f0416e1cc042c4f23cac112da43c5573db6c20 Mon Sep 17 00:00:00 2001 From: Peng Wang Date: Tue, 31 Jan 2023 15:02:20 -0800 Subject: [PATCH 26/52] Fixes shape inference of LookupTableImportV2 to handle scalar values. PiperOrigin-RevId: 506126405 --- tensorflow/core/ops/lookup_ops.cc | 5 +++-- .../data_structures/lookup_ops_test.py | 15 +++++++++++++++ 2 files changed, 18 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/ops/lookup_ops.cc b/tensorflow/core/ops/lookup_ops.cc index 2c279776dca938..4e074e0ef746e6 100644 --- a/tensorflow/core/ops/lookup_ops.cc +++ b/tensorflow/core/ops/lookup_ops.cc @@ -309,9 +309,10 @@ REGISTER_OP("LookupTableImportV2") ShapeHandle keys; TF_RETURN_IF_ERROR(c->WithRank(c->input(1), 1, &keys)); + ShapeHandle values; + TF_RETURN_IF_ERROR(c->WithRankAtLeast(c->input(2), 1, &values)); DimensionHandle unused; - TF_RETURN_IF_ERROR( - c->Merge(c->Dim(keys, 0), c->Dim(c->input(2), 0), &unused)); + TF_RETURN_IF_ERROR(c->Merge(c->Dim(keys, 0), c->Dim(values, 0), &unused)); return OkStatus(); }); diff --git a/tensorflow/python/kernel_tests/data_structures/lookup_ops_test.py b/tensorflow/python/kernel_tests/data_structures/lookup_ops_test.py index 1ce760afd091ec..816918ecb95fb7 100644 --- a/tensorflow/python/kernel_tests/data_structures/lookup_ops_test.py +++ b/tensorflow/python/kernel_tests/data_structures/lookup_ops_test.py @@ -42,6 +42,7 @@ from tensorflow.python.framework import test_util from tensorflow.python.ops import array_ops from tensorflow.python.ops import control_flow_ops +from tensorflow.python.ops import gen_lookup_ops from tensorflow.python.ops import lookup_ops from tensorflow.python.ops import map_fn from tensorflow.python.ops import variables @@ -574,6 +575,20 @@ def false_fn(): self.evaluate(lookup_ops.tables_initializer()) self.assertAllEqual(grad, -10.) + def testImportShapeInference(self, is_anonymous): + v = variables.Variable(1) + + @def_function.function(jit_compile=True) + def foo(): + return gen_lookup_ops.lookup_table_import_v2( + table_handle=v.handle, keys=[1.1, 2.2], values=1 + ) + + with self.assertRaisesRegex( + ValueError, r"Shape must be at least rank 1 but is rank 0" + ): + foo() + def testExportShapeInference(self, is_anonymous): table = self.getHashTable()( lookup_ops.KeyValueTensorInitializer( From 08e79ae2d71f83ce1f3bdbcb6ff0df9dc9a28594 Mon Sep 17 00:00:00 2001 From: Felix Chern Date: Fri, 3 Feb 2023 07:30:19 -0800 Subject: [PATCH 27/52] Fixes crashes due to buzz input for ApproxTopK PiperOrigin-RevId: 506898015 --- tensorflow/core/ops/nn_ops.cc | 11 ++++++++++ .../kernel_tests/math_ops/approx_topk_test.py | 21 +++++++++++++++++++ 2 files changed, 32 insertions(+) diff --git a/tensorflow/core/ops/nn_ops.cc b/tensorflow/core/ops/nn_ops.cc index c16563001e48ca..d75c173b658812 100644 --- a/tensorflow/core/ops/nn_ops.cc +++ b/tensorflow/core/ops/nn_ops.cc @@ -1395,10 +1395,17 @@ Status ApproxTopKShape(shape_inference::InferenceContext* c) { TF_RETURN_IF_ERROR(c->GetAttr("aggregate_to_topk", &aggregate_to_topk)); ShapeHandle input_shape; TF_RETURN_IF_ERROR(c->WithRankAtLeast(c->input(0), 1, &input_shape)); + int64_t r_dim_copy = reduction_dimension; + int64_t rank = c->Rank(input_shape); if (reduction_dimension < 0) { // Reverse index reduction_dimension += c->Rank(input_shape); } + if (reduction_dimension >= c->Rank(input_shape)) { + return errors::InvalidArgument("Invalid reduction dimension: ", r_dim_copy, + ". Must be within the range of [", -rank, + ", ", rank - 1, "]"); + } int64_t reduction_dim_value = c->Value(c->Dim(input_shape, reduction_dimension)); @@ -1406,6 +1413,10 @@ Status ApproxTopKShape(shape_inference::InferenceContext* c) { return errors::InvalidArgument("input must have last dimension >= k = ", k, " but was ", reduction_dim_value); } + if (recall_target > 1.0 || recall_target <= 0.) { + return errors::InvalidArgument("Invalid recall target: ", recall_target, + ". Valid value range in : [0, 1.0]."); + } int64_t output_dim_value = [&] { if (aggregate_to_topk) { diff --git a/tensorflow/python/kernel_tests/math_ops/approx_topk_test.py b/tensorflow/python/kernel_tests/math_ops/approx_topk_test.py index d3b04f64b5943b..c3c8126f325ea1 100644 --- a/tensorflow/python/kernel_tests/math_ops/approx_topk_test.py +++ b/tensorflow/python/kernel_tests/math_ops/approx_topk_test.py @@ -23,6 +23,7 @@ from tensorflow.python.eager.def_function import function from tensorflow.python.framework import constant_op from tensorflow.python.framework import dtypes +from tensorflow.python.framework import errors from tensorflow.python.framework import test_util from tensorflow.python.ops import array_ops from tensorflow.python.ops import math_ops @@ -215,6 +216,26 @@ def ann_with_grads(db, out_grads): ann_with_grads(db_op, out_grads_op)) self.assertAllClose(expected_in_grads, result_in_grads) + def test_invalid_input(self): + @function(jit_compile=True) + def fuzz_jit(): + return nn_ops.approx_max_k( + [ + 183.39395141601562, + 62.6842041015625, + 83.8385238647461, + 204.36642456054688, + ], + 4774, + reduction_dimension=0x8282828, + recall_target=135.9822179933652, + reduction_input_size_override=6154, + aggregate_to_topk=True, + ) + + with self.assertRaises((errors.InvalidArgumentError, ValueError)): + fuzz_jit() + if __name__ == '__main__': test.main() From a27a30d3bc119f39199fcae65f261a0eb1bef023 Mon Sep 17 00:00:00 2001 From: Matthias Kramm Date: Wed, 1 Feb 2023 15:02:59 -0800 Subject: [PATCH 28/52] Fix memory corruption vulnerability in reverse_sequence_op. PiperOrigin-RevId: 506433062 --- tensorflow/compiler/tests/BUILD | 24 +++++++++ .../tests/reverse_sequence_op_args_test.py | 52 +++++++++++++++++++ .../tf2xla/kernels/reverse_sequence_op.cc | 5 ++ 3 files changed, 81 insertions(+) create mode 100644 tensorflow/compiler/tests/reverse_sequence_op_args_test.py diff --git a/tensorflow/compiler/tests/BUILD b/tensorflow/compiler/tests/BUILD index 38b45f89f3a07a..539901e6423d5d 100644 --- a/tensorflow/compiler/tests/BUILD +++ b/tensorflow/compiler/tests/BUILD @@ -1203,6 +1203,30 @@ tf_xla_py_test( ], ) +# copybara:uncomment_begin(google-only) +# tf_xla_py_test( +# name = "reverse_sequence_op_args_test", +# size = "medium", +# srcs = ["reverse_sequence_op_args_test.py"], +# enable_mlir_bridge = False, +# main = "reverse_sequence_op_args_test.py", +# python_version = "PY3", +# tags = [ +# "no_pip", +# "optonly", +# ], +# deps = [ +# ":xla_test", +# "//tensorflow/compiler/jit:xla_cpu_jit", # DisableOnExport +# "//tensorflow/python:array_ops", +# "//tensorflow/python:framework", +# "//tensorflow/python:platform_test", +# "//tensorflow/python/compat:v2_compat", +# "//tensorflow/python/eager:function", +# ], +# ) +# copybara:uncomment_end + tf_xla_py_test( name = "rmsprop_test", size = "small", diff --git a/tensorflow/compiler/tests/reverse_sequence_op_args_test.py b/tensorflow/compiler/tests/reverse_sequence_op_args_test.py new file mode 100644 index 00000000000000..3ccb9b1df27880 --- /dev/null +++ b/tensorflow/compiler/tests/reverse_sequence_op_args_test.py @@ -0,0 +1,52 @@ +# Copyright 2023 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. +# ============================================================================== +"""Tests for tensorflow.ops.reverse_sequence_op.""" + +from tensorflow.compiler.tests import xla_test +from tensorflow.python.compat import v2_compat +from tensorflow.python.eager import def_function +from tensorflow.python.framework import errors +from tensorflow.python.ops import array_ops +from tensorflow.python.platform import test + + +class ReverseSequenceArgsTest(xla_test.XLATestCase): + """Tests argument verification of array_ops.reverse_sequence.""" + + def testInvalidArguments(self): + # seq_axis negative + with self.assertRaisesRegex( + (errors.InvalidArgumentError, ValueError), "seq_dim must be >=0" + ): + + @def_function.function(jit_compile=True) + def f(x): + return array_ops.reverse_sequence(x, [2, 2], seq_axis=-1) + + f([[1, 2], [3, 4]]) + + # batch_axis negative + with self.assertRaisesRegex(ValueError, "batch_dim must be >=0"): + + @def_function.function(jit_compile=True) + def g(x): + return array_ops.reverse_sequence(x, [2, 2], seq_axis=1, batch_axis=-1) + + g([[1, 2], [3, 4]]) + + +if __name__ == "__main__": + v2_compat.enable_v2_behavior() + test.main() diff --git a/tensorflow/compiler/tf2xla/kernels/reverse_sequence_op.cc b/tensorflow/compiler/tf2xla/kernels/reverse_sequence_op.cc index a9a67bc3b170e0..08a8545ec689ca 100644 --- a/tensorflow/compiler/tf2xla/kernels/reverse_sequence_op.cc +++ b/tensorflow/compiler/tf2xla/kernels/reverse_sequence_op.cc @@ -42,10 +42,15 @@ class ReverseSequenceOp : public XlaOpKernel { seq_lens_shape.dims())); OP_REQUIRES(context, batch_dim_ != seq_dim_, errors::InvalidArgument("batch_dim == seq_dim == ", seq_dim_)); + OP_REQUIRES(context, seq_dim_ >= 0, + errors::InvalidArgument("seq_dim must be >=0, got ", seq_dim_)); OP_REQUIRES( context, seq_dim_ < input_shape.dims(), errors::InvalidArgument("seq_dim must be < input rank", " ( ", seq_dim_, " vs. ", input_shape.dims(), ")")); + OP_REQUIRES( + context, batch_dim_ >= 0, + errors::InvalidArgument("batch_dim must be >=0, got ", batch_dim_)); OP_REQUIRES( context, batch_dim_ < input_shape.dims(), errors::InvalidArgument("batch_dim must be < input rank", " ( ", From c7f0b8e334061ab86c3d45d46ea397b17c34b3e4 Mon Sep 17 00:00:00 2001 From: Justin Szaday Date: Wed, 18 Jan 2023 11:44:11 -0800 Subject: [PATCH 29/52] Add bounds-checking to `TPUPartitionedInput` and `TPUPartitionedOutput` ops. PiperOrigin-RevId: 502936861 --- .../core/tpu/ops/tpu_partitioned_input_op.cc | 104 ++++++++++++++++++ .../core/tpu/ops/tpu_partitioned_output_op.cc | 79 ++++++++++++- 2 files changed, 177 insertions(+), 6 deletions(-) diff --git a/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc b/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc index 1a185ad210733a..c905664cac95bc 100644 --- a/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc +++ b/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc @@ -45,7 +45,30 @@ REGISTER_OP("TPUPartitionedInput") int partition_dim; TF_RETURN_IF_ERROR(c->GetAttr("partition_dim", &partition_dim)); + if (c->num_inputs() == 0) { + return errors::InvalidArgument( + "Expected at least one input to TPUPartitionedInput."); + } + ShapeHandle cur = c->input(c->num_inputs() - 1); + int rank = InferenceContext::kUnknownRank; + if (dtype == DT_RESOURCE) { + auto* shapes_and_types = + c->input_handle_shapes_and_types(c->num_inputs() - 1); + if (shapes_and_types) { + ShapeHandle shape_handle = shapes_and_types->at(0).shape; + rank = InferenceContext::Rank(shape_handle); + } + } else { + rank = InferenceContext::Rank(cur); + } + + // limitation: can only validate rank when it is known + if ((rank != InferenceContext::kUnknownRank && partition_dim >= rank) || + (partition_dim < -1)) + return errors::InvalidArgument("Cannot partition dim ", partition_dim, + " of rank ", rank, " tensor."); + for (int i = c->num_inputs() - 2; i >= 0; --i) { TF_RETURN_WITH_CONTEXT_IF_ERROR(c->Merge(c->input(i), cur, &cur), "From merging shape ", i, @@ -101,4 +124,85 @@ REGISTER_OP("TPUPartitionedInput") return OkStatus(); }); +REGISTER_OP("TPUPartitionedInputV2") + .Input("inputs: N * T") + .Output("output: T") + .Attr("N: int >= 1") + .Attr("T: type") + .Attr("partition_dims: list(int)") + .Attr("is_packed: bool = false") + .SetShapeFn([](InferenceContext* c) { + DataType dtype; + TF_RETURN_IF_ERROR(c->GetAttr("T", &dtype)); + std::vector partition_dims; + TF_RETURN_IF_ERROR(c->GetAttr("partition_dims", &partition_dims)); + bool is_packed; + TF_RETURN_IF_ERROR(c->GetAttr("is_packed", &is_packed)); + + int num_partitions = 1; + for (const int& partition_dim : partition_dims) { + num_partitions *= partition_dim; + } + + bool replicated = partition_dims.empty(); + int num_inputs_expected = is_packed ? 1 : num_partitions; + if (!((replicated && !is_packed) || + (c->num_inputs() == num_inputs_expected))) { + // we cannot validate the number of inputs for replicated, unpacked ops + // since we cannot infer the number of partitions from partition_dims + return errors::InvalidArgument("Expected ", num_inputs_expected, + " inputs, got ", c->num_inputs(), "."); + } else if (c->num_inputs() == 0) { + return errors::InvalidArgument( + "Expected at least one input to TPUPartitionedInputV2."); + } + + ShapeHandle output_shape; + if (dtype == DT_RESOURCE) { + ShapeHandle previous_shape_handle; + const std::vector* shapes_and_types = + nullptr; + for (int i = c->num_inputs() - 1; i >= 0; --i) { + shapes_and_types = c->input_handle_shapes_and_types(i); + if (shapes_and_types) { + ShapeHandle shape_handle = shapes_and_types->at(0).shape; + if (!c->FullyDefined(shape_handle)) { + return errors::InvalidArgument("Inputs must have static shape,", + "input[", i, + "] has unknown dimension."); + } + + if (i != c->num_inputs() - 1) { + ShapeHandle tmp; + if (!c->Merge(shape_handle, previous_shape_handle, &tmp).ok()) { + return errors::InvalidArgument( + "Inputs must have the same shape."); + } + } else { + previous_shape_handle = shape_handle; + } + } + } + + if (shapes_and_types) { + TF_ASSIGN_OR_RETURN( + output_shape, + _ComputeOutputShape(c, previous_shape_handle, partition_dims)); + std::vector output_shapes_and_types; + output_shapes_and_types.push_back(shape_inference::ShapeAndType( + output_shape, shapes_and_types->at(0).dtype)); + c->set_output_handle_shapes_and_types(0, output_shapes_and_types); + } + } + + if (!c->FullyDefined(output_shape)) { + TF_ASSIGN_OR_RETURN( + output_shape, _ComputeOutputShape(c, c->input(0), partition_dims)); + } + + c->set_output(0, output_shape); + + return OkStatus(); + }); + } // namespace tensorflow diff --git a/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc b/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc index 9ab183ad6d3dcc..da9fd60923cd25 100644 --- a/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc +++ b/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc @@ -38,19 +38,86 @@ REGISTER_OP("TPUPartitionedOutput") TF_RETURN_IF_ERROR(c->GetAttr("num_splits", &num_splits)); if (dtype == DT_RESOURCE) { return errors::Unimplemented("Not implemented."); + } else if (c->num_inputs() == 0) { + return errors::InvalidArgument( + "Expected at least one input to TPUPartitionedOutput."); } ShapeHandle input = c->input(0); + int rank = InferenceContext::Rank(input); + // limitation: can only validate rank when it is known + if ((rank != InferenceContext::kUnknownRank && partition_dim >= rank) || + (partition_dim < -1)) + return errors::InvalidArgument("Cannot partition dim ", partition_dim, + " of rank ", rank, " tensor."); + ShapeHandle newoutput0; - shape_inference::DimensionHandle new_dim; - TF_RETURN_WITH_CONTEXT_IF_ERROR( - c->Divide(c->Dim(input, partition_dim), num_splits, - true /* evenly_divisible */, &new_dim), - "Number of ways to split should evenly divide the split dimension"); - TF_CHECK_OK(c->ReplaceDim(input, partition_dim, new_dim, &newoutput0)); + if (partition_dim == -1) { + newoutput0 = input; // replicated input/output share shapes + } else { + shape_inference::DimensionHandle new_dim; + TF_RETURN_WITH_CONTEXT_IF_ERROR( + c->Divide(c->Dim(input, partition_dim), num_splits, + true /* evenly_divisible */, &new_dim), + "Number of ways to split should evenly divide the split dimension"); + TF_CHECK_OK(c->ReplaceDim(input, partition_dim, new_dim, &newoutput0)); + } + for (int i = num_splits - 1; i >= 0; --i) { c->set_output(i, newoutput0); } + + return OkStatus(); + }); + +REGISTER_OP("TPUPartitionedOutputV2") + .Input("inputs: T") + .Output("output: num_splits * T") + .Attr("T: type") + .Attr("num_splits: int >= 1") + .Attr("partition_dims: list(int)") + .SetShapeFn([](InferenceContext* c) { + DataType dtype; + TF_RETURN_IF_ERROR(c->GetAttr("T", &dtype)); + std::vector partition_dims; + TF_RETURN_IF_ERROR(c->GetAttr("partition_dims", &partition_dims)); + int num_splits; + TF_RETURN_IF_ERROR(c->GetAttr("num_splits", &num_splits)); + if (dtype == DT_RESOURCE) { + return errors::Unimplemented("Not implemented."); + } else if (c->num_inputs() == 0) { + return errors::InvalidArgument( + "Expected at least one input to TPUPartitionedOutputV2."); + } + + ShapeHandle handle = c->input(0); + int rank = InferenceContext::Rank(handle); + int num_cores_per_replica = 1; + for (const int& partition_dim : partition_dims) { + num_cores_per_replica *= partition_dim; + } + + if (num_splits != num_cores_per_replica) { + return errors::InvalidArgument("Expected ", num_cores_per_replica, + " splits."); + } else if (rank > (int)partition_dims.size()) { + return errors::InvalidArgument("Expected at least ", rank, + " partition dimensions."); + } + + for (int i = 0; i < rank; ++i) { + shape_inference::DimensionHandle dim; + TF_RETURN_WITH_CONTEXT_IF_ERROR( + c->Divide(c->Dim(handle, i), partition_dims[i], + true /* evenly_divisible */, &dim), + "Number of ways to split should evenly divide the split dimension"); + TF_CHECK_OK(c->ReplaceDim(handle, i, dim, &handle)); + } + + for (int i = num_splits - 1; i >= 0; --i) { + c->set_output(i, handle); + } + return OkStatus(); }); From bf27d68697235d85256dfec09f83526bb8781e2d Mon Sep 17 00:00:00 2001 From: Zhufeng Pan Date: Thu, 26 Jan 2023 13:43:37 -0800 Subject: [PATCH 30/52] Add check for "ksize" argument in MaxPoolWithArgmax PiperOrigin-RevId: 504930976 --- tensorflow/core/kernels/maxpooling_op.cc | 5 +++ tensorflow/core/ops/nn_ops.cc | 8 +++++ .../kernel_tests/nn_ops/pooling_ops_test.py | 31 +++++++++++++------ 3 files changed, 35 insertions(+), 9 deletions(-) diff --git a/tensorflow/core/kernels/maxpooling_op.cc b/tensorflow/core/kernels/maxpooling_op.cc index 4fb198690d5ad2..927f6331d78265 100644 --- a/tensorflow/core/kernels/maxpooling_op.cc +++ b/tensorflow/core/kernels/maxpooling_op.cc @@ -966,6 +966,11 @@ class MaxPoolingWithArgmaxOp : public OpKernel { OP_REQUIRES(context, ksize_.size() == 4, errors::InvalidArgument("Sliding window ksize field must " "specify 4 dimensions")); + for (int i = 0; i < ksize_.size(); ++i) { + OP_REQUIRES(context, ksize_[i] > 0, + errors::InvalidArgument( + "ksize must be a postive int32 value, got:", ksize_[i])); + } OP_REQUIRES_OK(context, context->GetAttr("strides", &stride_)); OP_REQUIRES(context, stride_.size() == 4, errors::InvalidArgument("Sliding window stride field must " diff --git a/tensorflow/core/ops/nn_ops.cc b/tensorflow/core/ops/nn_ops.cc index c16563001e48ca..e3ad108420700a 100644 --- a/tensorflow/core/ops/nn_ops.cc +++ b/tensorflow/core/ops/nn_ops.cc @@ -972,6 +972,14 @@ REGISTER_OP("MaxPoolWithArgmax") .Output("argmax: Targmax") .Attr("T: realnumbertype") .SetShapeFn([](InferenceContext* c) { + std::vector ksize; + TF_RETURN_IF_ERROR(c->GetAttr("ksize", &ksize)); + for (int i = 0; i < ksize.size(); ++i) { + if (ksize[i] <= 0) { + return errors::InvalidArgument( + "ksize must be a postive int32 value, got:", ksize[i]); + } + } TF_RETURN_IF_ERROR(shape_inference::MaxPoolShape(c)); c->set_output(1, c->output(0)); return OkStatus(); diff --git a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py index e279f6e0027da2..0888953abe4699 100644 --- a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py @@ -784,6 +784,19 @@ def testMaxPoolInvalidFilterSize(self, **kwargs): t = self.evaluate( nn_ops.max_pool(t, ksize=[1, 1, 2, 1], strides=1, padding="VALID")) + @test_util.run_in_graph_and_eager_modes + def testMaxPoolWithArgmaxKsizeOverflow(self): + with self.assertRaisesRegex( + (ValueError, errors_impl.InvalidArgumentError), + "ksize must be a postive int32 value"): + with self.cached_session(): + t = gen_nn_ops.max_pool_with_argmax( + input=[[[[1, 1, 1]]]], + ksize=[1, -2**31, 4, 1], + strides=[1, 1000, 1, 7], + padding="SAME") + self.evaluate(t) + # Tests for DepthwiseMaxPooling on CPU only. @parameterized.parameters( GetTestConfigsDicts( @@ -1069,9 +1082,9 @@ def testMaxPoolingGradThrowDeterminismError(self): t = constant_op.constant(tensor_input, shape=[2, 2, 2, 1]) argmax_t = constant_op.constant( [0, 1, 3, 5, 0, 2, 6, 8], shape=[2, 2, 2, 1], dtype=dtypes.int64) - with self.assertRaisesRegexp( - errors_impl.UnimplementedError, "Determinism is not yet supported " - "for MaxPoolGradWithArgmax."): + with self.assertRaisesRegex( + errors_impl.UnimplementedError, "Determinism is not yet supported" + " for MaxPoolGradWithArgmax."): out_op = gen_nn_ops.max_pool_grad_with_argmax( orig_in, t, @@ -2335,7 +2348,7 @@ def testOpEdgeCases(self): @test_util.run_deprecated_v1 def testEdgeCasesRaiseErrors(self): - with self.assertRaisesRegexp( + with self.assertRaisesRegex( ValueError, "NCHW_VECT_C.*is not supported with " "explicit padding|XLA does not support pooling ops with explicit " "padding"): @@ -2345,7 +2358,7 @@ def testEdgeCasesRaiseErrors(self): strides=[1, 2, 2, 1], padding=[[0, 0], [0, 1], [0, 1], [0, 0]], data_format="NCHW_VECT_C") - with self.assertRaisesRegexp( + with self.assertRaisesRegex( ValueError, "Explicit padding is not supported with an input " "tensor of rank 5"): nn_ops.max_pool_v2( @@ -2354,7 +2367,7 @@ def testEdgeCasesRaiseErrors(self): strides=[1, 2, 2, 1, 1], padding=[[0, 0], [0, 1], [0, 1], [0, 0]], data_format="NCHW") - with self.assertRaisesRegexp( + with self.assertRaisesRegex( ValueError, "Attr 'padding' of 'MaxPoolV2' Op passed " "string 'EXPLICIT'"): gen_nn_ops.max_pool_v2( @@ -2367,7 +2380,7 @@ def testEdgeCasesRaiseErrors(self): @test_util.run_deprecated_v1 def testEdgeCasesExcessPadding(self): with self.session(use_gpu=test.is_gpu_available()) as sess: - with self.assertRaisesRegexp( + with self.assertRaisesRegex( (errors_impl.UnimplementedError, errors_impl.InvalidArgumentError), "Right padding 2 needs to be smaller than the window size 2|" "XLA does not support pooling ops with explicit padding"): @@ -2385,7 +2398,7 @@ def testEdgeCasesExcessPadding(self): @test_util.run_deprecated_v1 def testNegativePadding(self): with self.session(use_gpu=test.is_gpu_available()) as sess: - with self.assertRaisesRegexp( + with self.assertRaisesRegex( ValueError, "All elements of explicit_paddings must be " "nonnegative for"): input_sizes = [1, 3, 3, 1] @@ -2402,7 +2415,7 @@ def testNegativePadding(self): @test_util.run_deprecated_v1 def testExplicitPaddingBatch(self): with self.session(use_gpu=test.is_gpu_available()) as sess: - with self.assertRaisesRegexp( + with self.assertRaisesRegex( ValueError, "Nonzero explicit padding in the batch or depth " "dimensions is not supported"): input_sizes = [1, 3, 3, 1] From 539df36e6da271467c5308829a68162fd709b728 Mon Sep 17 00:00:00 2001 From: unda <4316856+fcoUnda@users.noreply.github.com> Date: Fri, 3 Feb 2023 21:52:21 +0000 Subject: [PATCH 31/52] fixing pylint complaints --- tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py index 0888953abe4699..1b8aac0c377619 100644 --- a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py @@ -1072,8 +1072,8 @@ def testMaxPoolingGradThrowDeterminismError(self): try: config_exec.enable_op_determinism() orig_input = [ - 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 0.0, 0.0, - 0.0, 1.0, 0.0, 1.0 + 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 0.0, + 0.0, 0.0, 1.0, 0.0, 1.0 ] tensor_input = [11.0, 12.0, 13.0, 14.0, 21.0, 22.0, 23.0, 24.0] @@ -1100,8 +1100,8 @@ def testMaxPoolingGradThrowDeterminismError(self): try: config_exec.enable_op_determinism() orig_input = [ - 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 0.0, 0.0, - 0.0, 1.0, 0.0, 1.0 + 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 0.0, + 0.0, 0.0, 1.0, 0.0, 1.0 ] tensor_input = [11.0, 12.0, 13.0, 14.0, 21.0, 22.0, 23.0, 24.0] From 53fccb429a53b52ed03de4efcda45afe1cee924e Mon Sep 17 00:00:00 2001 From: Jian Cai Date: Wed, 1 Feb 2023 15:35:07 -0800 Subject: [PATCH 32/52] [Tensorflow] Fix security vulnerability with TensorListSplitOp PiperOrigin-RevId: 506441188 --- tensorflow/compiler/tests/tensor_list_ops_test.py | 11 +++++++++++ tensorflow/compiler/tf2xla/kernels/tensor_list_ops.cc | 2 ++ 2 files changed, 13 insertions(+) diff --git a/tensorflow/compiler/tests/tensor_list_ops_test.py b/tensorflow/compiler/tests/tensor_list_ops_test.py index 659d9f41e8d5eb..3c9b29b1835c79 100644 --- a/tensorflow/compiler/tests/tensor_list_ops_test.py +++ b/tensorflow/compiler/tests/tensor_list_ops_test.py @@ -236,6 +236,17 @@ def testZerosLikeForTensorList(self): self.assertAllEqual(z.shape.as_list(), [None]) self.assertAllEqual(z, [0.0, 0.0]) + def testInvalidSplitLength(self): + with self.session(), self.test_scope(): + tensor_list_split = list_ops.tensor_list_split( + tensor=[1], element_shape=[-1], lengths=[0] + ) + with self.assertRaisesRegex( + errors.UnimplementedError, "All lengths must be positive" + ): + self.evaluate(tensor_list_split) + + if __name__ == "__main__": os.environ["TF_XLA_FLAGS"] = ("--tf_xla_min_cluster_size=2 " + os.environ.get("TF_XLA_FLAGS", "")) diff --git a/tensorflow/compiler/tf2xla/kernels/tensor_list_ops.cc b/tensorflow/compiler/tf2xla/kernels/tensor_list_ops.cc index 980ca07e11719d..5d299fde60071e 100644 --- a/tensorflow/compiler/tf2xla/kernels/tensor_list_ops.cc +++ b/tensorflow/compiler/tf2xla/kernels/tensor_list_ops.cc @@ -553,6 +553,8 @@ class TensorListSplitOp : public XlaOpKernel { OP_REQUIRES(ctx, len == length, errors::Unimplemented("All lengths have to be the same")); } + OP_REQUIRES(ctx, length, + errors::Unimplemented("All lengths must be positive")); OP_REQUIRES( ctx, element_dims[0] % length == 0, errors::Unimplemented("Buffer size has to be a multiple of length")); From 038f2eaeabb5e651274b2439e90f918cb06b31e0 Mon Sep 17 00:00:00 2001 From: Mihai Maruseac Date: Wed, 28 Dec 2022 08:33:13 -0800 Subject: [PATCH 33/52] Update tensorflow/core/kernels/image/decode_image_op.cc --- tensorflow/core/kernels/image/decode_image_op.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/core/kernels/image/decode_image_op.cc b/tensorflow/core/kernels/image/decode_image_op.cc index baa4407fd84dda..ed6f09ff7fb9ca 100644 --- a/tensorflow/core/kernels/image/decode_image_op.cc +++ b/tensorflow/core/kernels/image/decode_image_op.cc @@ -457,7 +457,7 @@ class DecodeImageV2Op : public OpKernel { uint8* buffer = gif::Decode( input.data(), input.size(), [&](int num_frames, int width, int height, int channels) -> uint8* { - buffer_size = ptrdiff_t(num_frames) * height * width * channels; + buffer_size = static_cast(num_frames) * height * width * channels; Status status; // By the existing API, we support decoding GIF with `decode_jpeg` or From 11eb850b33138ba1bbab1ca4972be5bca661ba4e Mon Sep 17 00:00:00 2001 From: Andrei Pikas Date: Wed, 28 Dec 2022 14:59:28 +0300 Subject: [PATCH 34/52] Replace ptrdiff_t with int64_t. --- tensorflow/core/kernels/image/decode_image_op.cc | 4 ++-- tensorflow/core/lib/gif/gif_io.cc | 14 +++++++------- tensorflow/core/lib/gif/gif_io_test.cc | 2 +- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/tensorflow/core/kernels/image/decode_image_op.cc b/tensorflow/core/kernels/image/decode_image_op.cc index baa4407fd84dda..f09b4624c91cac 100644 --- a/tensorflow/core/kernels/image/decode_image_op.cc +++ b/tensorflow/core/kernels/image/decode_image_op.cc @@ -452,12 +452,12 @@ class DecodeImageV2Op : public OpKernel { // allocation til after dtype conversion is done. `gif`::Decode` supports // uint8 only. Tensor* output = nullptr; - ptrdiff_t buffer_size = 0; + int64_t buffer_size = 0; string error_string; uint8* buffer = gif::Decode( input.data(), input.size(), [&](int num_frames, int width, int height, int channels) -> uint8* { - buffer_size = ptrdiff_t(num_frames) * height * width * channels; + buffer_size = int64_t(num_frames) * height * width * channels; Status status; // By the existing API, we support decoding GIF with `decode_jpeg` or diff --git a/tensorflow/core/lib/gif/gif_io.cc b/tensorflow/core/lib/gif/gif_io.cc index dddb8708c5e02b..dfe1ee91f76b59 100644 --- a/tensorflow/core/lib/gif/gif_io.cc +++ b/tensorflow/core/lib/gif/gif_io.cc @@ -105,7 +105,7 @@ uint8* Decode(const void* srcdata, int datasize, uint8* const dstdata = allocate_output(target_num_frames, width, height, channel); if (!dstdata) return nullptr; - for (ptrdiff_t k = 0; k < target_num_frames; k++) { + for (int64_t k = 0; k < target_num_frames; k++) { uint8* this_dst = dstdata + k * width * channel * height; SavedImage* this_image = &gif_file->SavedImages[k]; @@ -125,10 +125,10 @@ uint8* Decode(const void* srcdata, int datasize, if (k > 0) { uint8* last_dst = dstdata + (k - 1) * width * channel * height; - for (ptrdiff_t i = 0; i < height; ++i) { + for (int64_t i = 0; i < height; ++i) { uint8* p_dst = this_dst + i * width * channel; uint8* l_dst = last_dst + i * width * channel; - for (ptrdiff_t j = 0; j < width; ++j) { + for (int64_t j = 0; j < width; ++j) { p_dst[j * channel + 0] = l_dst[j * channel + 0]; p_dst[j * channel + 1] = l_dst[j * channel + 1]; p_dst[j * channel + 2] = l_dst[j * channel + 2]; @@ -141,9 +141,9 @@ uint8* Decode(const void* srcdata, int datasize, // If the first frame does not fill the entire canvas then fill the // unoccupied canvas with zeros (black). if (k == 0) { - for (ptrdiff_t i = 0; i < height; ++i) { + for (int64_t i = 0; i < height; ++i) { uint8* p_dst = this_dst + i * width * channel; - for (ptrdiff_t j = 0; j < width; ++j) { + for (int64_t j = 0; j < width; ++j) { p_dst[j * channel + 0] = 0; p_dst[j * channel + 1] = 0; p_dst[j * channel + 2] = 0; @@ -165,9 +165,9 @@ uint8* Decode(const void* srcdata, int datasize, return nullptr; } - for (ptrdiff_t i = imgTop; i < imgBottom; ++i) { + for (int64_t i = imgTop; i < imgBottom; ++i) { uint8* p_dst = this_dst + i * width * channel; - for (ptrdiff_t j = imgLeft; j < imgRight; ++j) { + for (int64_t j = imgLeft; j < imgRight; ++j) { GifByteType color_index = this_image->RasterBits[(i - img_desc->Top) * (img_desc->Width) + (j - img_desc->Left)]; diff --git a/tensorflow/core/lib/gif/gif_io_test.cc b/tensorflow/core/lib/gif/gif_io_test.cc index 32e63480720663..5cd4999c799563 100644 --- a/tensorflow/core/lib/gif/gif_io_test.cc +++ b/tensorflow/core/lib/gif/gif_io_test.cc @@ -52,7 +52,7 @@ void TestDecodeGif(Env* env, DecodeGifTestCase testcase) { w = width; h = height; c = channels; - return new uint8[ptrdiff_t(frame_cnt) * height * width * channels]; + return new uint8[int64_t(frame_cnt) * height * width * channels]; }, &error_string)); ASSERT_NE(imgdata, nullptr); From 579775ba7d9a4f77ee3b42d0a76e7523af7c3ae3 Mon Sep 17 00:00:00 2001 From: shuw Date: Thu, 13 Oct 2022 08:02:27 -0700 Subject: [PATCH 35/52] Fix conjugate transpose for 0/1D complex tensor --- .../core/kernels/transpose_functor_gpu.cu.cc | 23 +++++++++++++++---- 1 file changed, 18 insertions(+), 5 deletions(-) diff --git a/tensorflow/core/kernels/transpose_functor_gpu.cu.cc b/tensorflow/core/kernels/transpose_functor_gpu.cu.cc index 0747685853ee07..35101d8ce3bb3f 100644 --- a/tensorflow/core/kernels/transpose_functor_gpu.cu.cc +++ b/tensorflow/core/kernels/transpose_functor_gpu.cu.cc @@ -31,6 +31,14 @@ typedef Eigen::GpuDevice GPUDevice; namespace tensorflow { namespace internal { +template +__global__ void ConjugateKernel(int nthreads, const T* __restrict__ src, + T* __restrict__ dst) { + GPU_1D_KERNEL_LOOP(idx, nthreads) { + dst[idx] = Eigen::numext::conj(ldg(src + idx)); + } +} + template __global__ void TransposeKernel(int nthreads, const T* __restrict__ src, const int32* __restrict__ buf, @@ -62,6 +70,15 @@ void TransposeSimple(const GPUDevice& d, const Tensor& in, CHECK_LT(nelem, kint32max) << "Tensor too large to transpose on GPU"; // Pack strides and permutation into one buffer. const int32 ndims = in.dims(); + GpuLaunchConfig cfg = GetGpuLaunchConfig(nelem, d); + const T* p = reinterpret_cast(in.tensor_data().data()); + T* q = reinterpret_cast(const_cast((out->tensor_data().data()))); + if (conjugate && ndims < 2) { + TF_CHECK_OK(GpuLaunchKernel(ConjugateKernel, cfg.block_count, + cfg.thread_per_block, 0, d.stream(), + cfg.virtual_thread_count, p, q)); + return; + } gtl::InlinedVector host_buf(ndims * 3); gtl::InlinedVector in_strides = ComputeStride(in.shape()); gtl::InlinedVector out_strides = ComputeStride(out->shape()); @@ -78,9 +95,6 @@ void TransposeSimple(const GPUDevice& d, const Tensor& in, // therefore we are doing a sync copy effectively. d.memcpyHostToDevice(dev_buf, host_buf.data(), num_bytes); // Launch kernel to q[...] = p[...]. - const T* p = reinterpret_cast(in.tensor_data().data()); - T* q = reinterpret_cast(const_cast((out->tensor_data().data()))); - GpuLaunchConfig cfg = GetGpuLaunchConfig(nelem, d); TF_CHECK_OK(GpuLaunchKernel( TransposeKernel, cfg.block_count, cfg.thread_per_block, 0, d.stream(), cfg.virtual_thread_count, p, @@ -179,8 +193,7 @@ template struct Transpose { static void run(const GPUDevice& d, const Tensor& in, const gtl::ArraySlice perm, Tensor* out) { - if (in.dims() < 2) return; - if (internal::TransposeUsingTile::run(d, in, perm, out)) { + if (in.dims() > 1 && internal::TransposeUsingTile::run(d, in, perm, out)) { return; } From 20e119c1b7d06aac9d6d455c8e32b4e998713f47 Mon Sep 17 00:00:00 2001 From: Jian Cai Date: Thu, 26 Jan 2023 17:27:04 -0800 Subject: [PATCH 36/52] Copy input tensor in a RandomShuffleOp to output directly when its number of elements is less than or equal to 1. This fixes the crash in TFXLA MLIR bridge when the input tensor has a rank of 0. This also makes its behavoir more consistent with TF core kernel. PiperOrigin-RevId: 504983211 --- .../compiler/mlir/xla/tests/legalize-tf.mlir | 11 +++++++++++ .../mlir/xla/transforms/legalize_tf.cc | 19 ++++++++++++------- tensorflow/compiler/tests/BUILD | 1 + tensorflow/compiler/tests/random_ops_test.py | 6 ++++++ 4 files changed, 30 insertions(+), 7 deletions(-) diff --git a/tensorflow/compiler/mlir/xla/tests/legalize-tf.mlir b/tensorflow/compiler/mlir/xla/tests/legalize-tf.mlir index 39d6198bfbebf6..b8405b9cb70192 100644 --- a/tensorflow/compiler/mlir/xla/tests/legalize-tf.mlir +++ b/tensorflow/compiler/mlir/xla/tests/legalize-tf.mlir @@ -5175,6 +5175,17 @@ func.func @tensor_scatter_max(%tensor: tensor, %indices: tensor tensor { + // CHECK: [[INPUT:%.*]] = mhlo.constant dense<1.000000e+20> : tensor + // CHECK-NEXT: return [[INPUT]] + %cst = "tf.Const"() {value = dense<1.000000e+20> : tensor} : () -> tensor + %0 = "tf.RandomShuffle"(%cst) {device = "", seed = -4294967297 : i64, seed2 = -2147483649 : i64} : (tensor) -> tensor + return %0 : tensor +} + +// ----- + // CHECK-LABEL: @random_shuffle_first_dim_1 // CHECK-SAME: [[INPUT:%.*]]: tensor<1x?xf32> func.func @random_shuffle_first_dim_1(%input: tensor<1x?xf32>) -> tensor<1x?xf32> { diff --git a/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc b/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc index 04ec56efb4861c..1b08efb9ee8726 100644 --- a/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc +++ b/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc @@ -5818,19 +5818,24 @@ class ConvertRandomShuffleOp : public OpRewritePattern { LogicalResult matchAndRewrite(TF::RandomShuffleOp op, PatternRewriter &rewriter) const override { - auto input_type = op.value().getType().dyn_cast(); + auto no_op = [&]() { + rewriter.replaceOp(op, op.getValue()); + return success(); + }; + + auto input_type = op.getValue().getType().dyn_cast(); if (!input_type) return failure(); + if (input_type.hasStaticShape() && input_type.getNumElements() <= 1) + // No shuffling is required, so copy input directly to output. + return no_op(); int64_t input_rank = input_type.getRank(); int64_t first_dim_size = input_type.getDimSize(0); if (ShapedType::isDynamic(first_dim_size)) return failure(); - // We are shuffling along the first dimension. If its size is <= 1, then - // shuffling is a no-op. - if (first_dim_size <= 1) { - rewriter.replaceOp(op, op.value()); - return success(); - } + if (first_dim_size <= 1) + // No shuffling is required, so copy input directly to output. + return no_op(); // For vectors, shuffle values by sorting instead of the obvious // Fisher-Yates algorithm. Fisher-Yates is simple to implement and correct, diff --git a/tensorflow/compiler/tests/BUILD b/tensorflow/compiler/tests/BUILD index 9fb09816fa8187..3ba67c8acbf458 100644 --- a/tensorflow/compiler/tests/BUILD +++ b/tensorflow/compiler/tests/BUILD @@ -1119,6 +1119,7 @@ tf_xla_py_test( ], deps = [ ":xla_test", + "//tensorflow:tensorflow_py", "//tensorflow/python:array_ops", "//tensorflow/python:framework", "//tensorflow/python:math_ops", diff --git a/tensorflow/compiler/tests/random_ops_test.py b/tensorflow/compiler/tests/random_ops_test.py index 03e919dee9107d..80be5b5e836daf 100644 --- a/tensorflow/compiler/tests/random_ops_test.py +++ b/tensorflow/compiler/tests/random_ops_test.py @@ -279,6 +279,12 @@ def testShuffle2d(self): self.assertAllEqual(len(result.flatten()), len(expected)) self.assertAllEqual(set(result.flatten()), set(expected)) + def testRandomShuffleInputRank0(self): + with self.session(): + with self.test_scope(): + shuffle = random_ops.random_shuffle(value=1e20) + self.evaluate(shuffle) + if __name__ == '__main__': googletest.main() From 49b7db4dbacf684bab48245df2d4286fb4ad6afb Mon Sep 17 00:00:00 2001 From: mdfaijul Date: Mon, 23 Jan 2023 15:45:13 -0800 Subject: [PATCH 37/52] Check min-max values to be scalars. --- tensorflow/core/kernels/mkl/mkl_qmatmul_op.cc | 60 ++++++++++++-- .../core/kernels/mkl/mkl_qmatmul_op_test.cc | 80 +++++++++---------- 2 files changed, 92 insertions(+), 48 deletions(-) diff --git a/tensorflow/core/kernels/mkl/mkl_qmatmul_op.cc b/tensorflow/core/kernels/mkl/mkl_qmatmul_op.cc index ec722bb66a6583..eb069a9aa275d3 100644 --- a/tensorflow/core/kernels/mkl/mkl_qmatmul_op.cc +++ b/tensorflow/core/kernels/mkl/mkl_qmatmul_op.cc @@ -317,8 +317,20 @@ class MklDnnQuantizedMatMulOp : public MklDnnMatMulOpBase { // This is the case the inner-product and requantization are fused. // "min_freezed_output" and "max_freezed_output" are the requested range // for the output. - min_output_value = context->input(7).flat()(0); - max_output_value = context->input(8).flat()(0); + const Tensor& min_freezed_tensor = context->input(7); + const Tensor& max_freezed_tensor = context->input(8); + OP_REQUIRES(context, + TensorShapeUtils::IsScalar(min_freezed_tensor.shape()), + errors::InvalidArgument( + "`min_freezed_output` must be rank 0 but is rank ", + min_freezed_tensor.dims())); + OP_REQUIRES(context, + TensorShapeUtils::IsScalar(max_freezed_tensor.shape()), + errors::InvalidArgument( + "`max_freezed_output` must be rank 0 but is rank ", + max_freezed_tensor.dims())); + min_output_value = min_freezed_tensor.scalar()(); + max_output_value = max_freezed_tensor.scalar()(); } else { ComputeOutputRangeForInt32(context, &min_output_value, &max_output_value); } @@ -344,10 +356,10 @@ class MklDnnQuantizedMatMulOp : public MklDnnMatMulOpBase { void ComputeOutputRangeForInt32(OpKernelContext* context, float* min_output_value, float* max_output_value) { - const float min_input = context->input(3).flat()(0); - const float max_input = context->input(4).flat()(0); - const float min_weight = context->input(5).flat()(0); - const float max_weight = context->input(6).flat()(0); + const float min_input = context->input(3).scalar()(); + const float max_input = context->input(4).scalar()(); + const float min_weight = context->input(5).scalar()(); + const float max_weight = context->input(6).scalar()(); MklQuantizationRangeForMultiplication( min_input, max_input, min_weight, max_weight, min_output_value, max_output_value); @@ -361,6 +373,25 @@ class MklDnnQuantizedMatMulOp : public MklDnnMatMulOpBase { params.dtypes.append(typeid(Tbias).name()); params.dtypes.append(typeid(Toutput).name()); + // min-max values for input and weight should be scalar. + const Tensor& min_input_tensor = context->input(3); + const Tensor& max_input_tensor = context->input(4); + const Tensor& min_weight_tensor = context->input(5); + const Tensor& max_weight_tensor = context->input(6); + + OP_REQUIRES(context, TensorShapeUtils::IsScalar(min_input_tensor.shape()), + errors::InvalidArgument("`min_a` must be rank 0 but is rank ", + min_input_tensor.dims())); + OP_REQUIRES(context, TensorShapeUtils::IsScalar(max_input_tensor.shape()), + errors::InvalidArgument("`max_a` must be rank 0 but is rank ", + max_input_tensor.dims())); + OP_REQUIRES(context, TensorShapeUtils::IsScalar(min_weight_tensor.shape()), + errors::InvalidArgument("`min_b` must be rank 0 but is rank ", + min_weight_tensor.dims())); + OP_REQUIRES(context, TensorShapeUtils::IsScalar(max_weight_tensor.shape()), + errors::InvalidArgument("`max_b` must be rank 0 but is rank ", + max_weight_tensor.dims())); + // When the output type is quint8, the output data is requantized into // quint8. A post_op "output_scale" is added to do the conversion. if (std::is_same::value || @@ -371,8 +402,21 @@ class MklDnnQuantizedMatMulOp : public MklDnnMatMulOpBase { ComputeOutputRangeForInt32(context, &min_output_value, &max_output_value); float scale_int32 = std::max(std::abs(min_output_value), std::abs(max_output_value)); - const float min_freezed_output = context->input(7).flat()(0); - const float max_freezed_output = context->input(8).flat()(0); + const Tensor& min_freezed_tensor = context->input(7); + const Tensor& max_freezed_tensor = context->input(8); + // min-max values of freezed output range should be scalar. + OP_REQUIRES(context, + TensorShapeUtils::IsScalar(min_freezed_tensor.shape()), + errors::InvalidArgument( + "`min_freezed_output` must be rank 0 but is rank ", + min_freezed_tensor.dims())); + OP_REQUIRES(context, + TensorShapeUtils::IsScalar(max_freezed_tensor.shape()), + errors::InvalidArgument( + "`max_freezed_output` must be rank 0 but is rank ", + max_freezed_tensor.dims())); + const float min_freezed_output = min_freezed_tensor.scalar()(); + const float max_freezed_output = max_freezed_tensor.scalar()(); float scale_eightbit = std::max(std::abs(min_freezed_output), std::abs(max_freezed_output)); float scale = 1.0; diff --git a/tensorflow/core/kernels/mkl/mkl_qmatmul_op_test.cc b/tensorflow/core/kernels/mkl/mkl_qmatmul_op_test.cc index de41b3d1264129..724f1758d91580 100644 --- a/tensorflow/core/kernels/mkl/mkl_qmatmul_op_test.cc +++ b/tensorflow/core/kernels/mkl/mkl_qmatmul_op_test.cc @@ -12,7 +12,7 @@ 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. ==============================================================================*/ -#if defined(INTEL_MKL) && defined(ENABLE_MKL) +#if defined(INTEL_MKL) #define EIGEN_USE_THREADS #include @@ -64,10 +64,10 @@ TEST_F(QuantizedMatMulTest, Small_withBias) { AddInputFromArray(TensorShape({3, 4}), {7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}); AddInputFromArray(TensorShape({4}), {1, 2, 3, 4}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); TF_ASSERT_OK(RunOpKernel()); // Here are the results we expect, from hand calculations: @@ -116,10 +116,10 @@ TEST_F(QuantizedMatMulTest, Small_withNegBias) { AddInputFromArray(TensorShape({3, 4}), {7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}); AddInputFromArray(TensorShape({4}), {100, -200, 300, -400}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); TF_ASSERT_OK(RunOpKernel()); // Here are the results we expect, from hand calculations: @@ -178,10 +178,10 @@ TEST_F(QuantizedMatMulTest, Small_WithNegInp) { AddInputFromArray(TensorShape({3, 2}), {1, 4, 2, 5, 3, 6}); // Bias AddInputFromArray(TensorShape({2}), {10.0f, 20.0f}); - AddInputFromArray(TensorShape({1}), {-12.0f}); - AddInputFromArray(TensorShape({1}), {243.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); + AddInputFromArray(TensorShape({}), {-12.0f}); + AddInputFromArray(TensorShape({}), {243.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); TF_ASSERT_OK(RunOpKernel()); // First calculate C = A * B, @@ -240,12 +240,12 @@ TEST_F(QuantizedMatMulTest, Small_withBiasAndReq) { AddInputFromArray(TensorShape({3, 4}), {7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}); AddInputFromArray(TensorShape({4}), {10, -20, 30, -40}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); TF_ASSERT_OK(RunOpKernel()); // Here are the results we expect, from hand calculations: @@ -308,12 +308,12 @@ TEST_F(QuantizedMatMulTest, Small_withBiasAndDeq) { AddInputFromArray(TensorShape({3, 4}), {7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}); AddInputFromArray(TensorShape({4}), {10, -20, 30, -40}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); TF_ASSERT_OK(RunOpKernel()); // Here are the results we expect, from hand calculations: @@ -375,10 +375,10 @@ TEST_F(QuantizedMatMulTest, Small_withBiasAndRelu) { {7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}); AddInputFromArray(TensorShape({4}), {100.0f, -200.0f, 300.0f, -400.0f}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); TF_ASSERT_OK(RunOpKernel()); // Here are the results we expect, from hand calculations: @@ -431,12 +431,12 @@ TEST_F(QuantizedMatMulTest, Small_withBiasAndReluAndReq) { AddInputFromArray(TensorShape({3, 4}), {7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}); AddInputFromArray(TensorShape({4}), {10, -20, 30, -40}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); TF_ASSERT_OK(RunOpKernel()); // Here are the results we expect, from hand calculations: @@ -502,10 +502,10 @@ TEST_F(QuantizedMatMulTest, Small_withWeightCached) { AddInputFromArray(TensorShape({3, 4}), {7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}); AddInputFromArray(TensorShape({4}), {1, 2, 3, 4}); - AddInputFromArray(TensorShape({1}), {0}); - AddInputFromArray(TensorShape({1}), {255.0f}); - AddInputFromArray(TensorShape({1}), {-127.0f}); - AddInputFromArray(TensorShape({1}), {127.0f}); + AddInputFromArray(TensorShape({}), {0}); + AddInputFromArray(TensorShape({}), {255.0f}); + AddInputFromArray(TensorShape({}), {-127.0f}); + AddInputFromArray(TensorShape({}), {127.0f}); int64 start_time = Env::Default()->NowMicros(); TF_ASSERT_OK(RunOpKernel()); @@ -543,4 +543,4 @@ TEST_F(QuantizedMatMulTest, Small_withWeightCached) { } // namespace tensorflow -#endif // INTEL_MKL && ENABLE_MKL +#endif // INTEL_MKL From 8b586523b42a9bebf7bcb83eae432b2b08121d80 Mon Sep 17 00:00:00 2001 From: learning-to-play <66660475+learning-to-play@users.noreply.github.com> Date: Fri, 3 Mar 2023 12:46:28 -0800 Subject: [PATCH 38/52] Adds Windows WSL2 GPU support to RELEASE.md --- RELEASE.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/RELEASE.md b/RELEASE.md index d55e4604985c6c..6321f26defbe97 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -18,6 +18,8 @@ * `tensorflow/python/keras` code is a legacy copy of Keras since the TensorFlow v2.7 release, and will be deleted in the v2.12 release. Please remove any import of `tensorflow.python.keras` and use the public API with `from tensorflow import keras` or `import tensorflow as tf; tf.keras`. +* For using TensorFlow GPU on Windows, you will need to install TensorFlow in Windows Subsystem for Linux (WSL2) (comment added during the TensorFlow 2.11.1 release). + ## Major Features and Improvements * `tf.lite`: From a24c0bd840c5ee9d8f165e3e70e0906a133fc1db Mon Sep 17 00:00:00 2001 From: learning-to-play <66660475+learning-to-play@users.noreply.github.com> Date: Fri, 3 Mar 2023 13:08:39 -0800 Subject: [PATCH 39/52] Update RELEASE.md --- RELEASE.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/RELEASE.md b/RELEASE.md index 6321f26defbe97..d2975a352bb7fd 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -18,7 +18,7 @@ * `tensorflow/python/keras` code is a legacy copy of Keras since the TensorFlow v2.7 release, and will be deleted in the v2.12 release. Please remove any import of `tensorflow.python.keras` and use the public API with `from tensorflow import keras` or `import tensorflow as tf; tf.keras`. -* For using TensorFlow GPU on Windows, you will need to install TensorFlow in Windows Subsystem for Linux (WSL2) (comment added during the TensorFlow 2.11.1 release). +* Starting TensorFlow 2.11, for using TensorFlow GPU on Windows, you will need to install TensorFlow in Windows Subsystem for Linux (WSL2) (comment added during the TensorFlow 2.11.1 release). ## Major Features and Improvements From f07e5aa7b962208dd30b1a13b48445aade05e6f2 Mon Sep 17 00:00:00 2001 From: learning-to-play <66660475+learning-to-play@users.noreply.github.com> Date: Fri, 3 Mar 2023 13:23:25 -0800 Subject: [PATCH 40/52] Updates RELEASE.md --- RELEASE.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/RELEASE.md b/RELEASE.md index d2975a352bb7fd..031918693ac9f3 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -18,7 +18,7 @@ * `tensorflow/python/keras` code is a legacy copy of Keras since the TensorFlow v2.7 release, and will be deleted in the v2.12 release. Please remove any import of `tensorflow.python.keras` and use the public API with `from tensorflow import keras` or `import tensorflow as tf; tf.keras`. -* Starting TensorFlow 2.11, for using TensorFlow GPU on Windows, you will need to install TensorFlow in Windows Subsystem for Linux (WSL2) (comment added during the TensorFlow 2.11.1 release). +* TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. ## Major Features and Improvements From 5655611b8a47fc27e717b9695673b7b624147516 Mon Sep 17 00:00:00 2001 From: Mihai Maruseac Date: Fri, 3 Mar 2023 15:20:23 -0800 Subject: [PATCH 41/52] Keep published release notes intact, update for patch release Signed-off-by: Mihai Maruseac --- RELEASE.md | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/RELEASE.md b/RELEASE.md index 031918693ac9f3..9e39dd4d3ca9a3 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -1,3 +1,11 @@ +# Release 2.11.1 + +**Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu (developed by Intel, see 2.11.0 release notes) and, optionally, try the TensorFlow-DirectML-Plugin. + +This release also introduces several vulnerability fixes: + +* TODO (in another PR) + # Release 2.11.0 ## Breaking Changes @@ -18,8 +26,6 @@ * `tensorflow/python/keras` code is a legacy copy of Keras since the TensorFlow v2.7 release, and will be deleted in the v2.12 release. Please remove any import of `tensorflow.python.keras` and use the public API with `from tensorflow import keras` or `import tensorflow as tf; tf.keras`. -* TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. - ## Major Features and Improvements * `tf.lite`: From 3d0ddec49174b017fb146f068d85c5d6754e39ba Mon Sep 17 00:00:00 2001 From: learning-to-play <66660475+learning-to-play@users.noreply.github.com> Date: Fri, 3 Mar 2023 15:29:23 -0800 Subject: [PATCH 42/52] Updates RELEASE.md --- RELEASE.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/RELEASE.md b/RELEASE.md index 9e39dd4d3ca9a3..c8c00080ea3f6e 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -1,6 +1,6 @@ # Release 2.11.1 -**Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu (developed by Intel, see 2.11.0 release notes) and, optionally, try the TensorFlow-DirectML-Plugin. +**Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. This release also introduces several vulnerability fixes: From 9db4c5af5622e7861601e7ab93630df39ea704e9 Mon Sep 17 00:00:00 2001 From: TensorFlow Release Automation Date: Mon, 13 Mar 2023 23:45:08 +0000 Subject: [PATCH 43/52] Insert release notes place-fill --- RELEASE.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/RELEASE.md b/RELEASE.md index c8c00080ea3f6e..2d3a0f53e5e301 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -1,5 +1,9 @@ # Release 2.11.1 + + +# Release 2.11.1 + **Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. This release also introduces several vulnerability fixes: From 39e6332029d34340a327ada41926c073d32bc40d Mon Sep 17 00:00:00 2001 From: venkat2469 <104534215+venkat2469@users.noreply.github.com> Date: Mon, 13 Mar 2023 17:11:13 -0700 Subject: [PATCH 44/52] Update RELEASE.md --- RELEASE.md | 30 ++++++++++++++++++++---------- 1 file changed, 20 insertions(+), 10 deletions(-) diff --git a/RELEASE.md b/RELEASE.md index 2d3a0f53e5e301..a97d21f37a7f52 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -1,14 +1,24 @@ # Release 2.11.1 - - - -# Release 2.11.1 - -**Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. - -This release also introduces several vulnerability fixes: - -* TODO (in another PR) + +* Fixes an FPE in TFLite in conv kernel [CVE-2023-27579] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-27579) +* Fixes a double free in Fractional(Max/Avg)Pool [CVE-2023-25801] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25801) +* Fixes a null dereference on ParallelConcat with XLA [CVE-2023-25676] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25676) +* Fixes a segfault in Bincount with XLA [CVE-2023-25675] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25675) +* Fixes an NPE in RandomShuffle with XLA enable [CVE-2023-25674] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25674) +* Fixes an FPE in TensorListSplit with XLA [CVE-2023-25673] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25673) +* Fixes segmentation fault in tfg-translate [CVE-2023-25671] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25671) +* Fixes an NPE in QuantizedMatMulWithBiasAndDequantize [CVE-2023-25670] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25670) +* Fixes an FPE in AvgPoolGrad with XLA [CVE-2023-25669] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25669) +* Fixes a heap out-of-buffer read vulnerability in the QuantizeAndDequantize operation [CVE-2023-25668] (https://cve.mitre.org/cgi-bin/cvename.cgi? name=CVE-2023-25668) +* Fixes a segfault when opening multiframe gif [CVE-2023-25667] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25667) +* Fixes an NPE in SparseSparseMaximum [CVE-2023-25665] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25665) +* Fixes an FPE in AudioSpectrogram [CVE-2023-25666] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25666) +* Fixes a heap-buffer-overflow in AvgPoolGrad [CVE-2023-25664] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25664) +* Fixes a NPE in TensorArrayConcatV2 [CVE-2023-25663] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25663) +* Fixes a Integer overflow in EditDistance [CVE-2023-25662] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25662) +* Fixes a Seg fault in `tf.raw_ops.Print` [CVE-2023-25660] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25660) +* Fixes a OOB read in DynamicStitch [CVE-2023-25659] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25659) +* Fixes a OOB Read in GRUBlockCellGrad [CVE-2023-25658] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25658) # Release 2.11.0 From 8a9d48690b6251f391521698a1b9c8e849c99ca4 Mon Sep 17 00:00:00 2001 From: learning-to-play <66660475+learning-to-play@users.noreply.github.com> Date: Mon, 13 Mar 2023 17:21:29 -0700 Subject: [PATCH 45/52] Update RELEASE.md --- RELEASE.md | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/RELEASE.md b/RELEASE.md index a97d21f37a7f52..8ff21eba50b13f 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -1,5 +1,9 @@ # Release 2.11.1 - + +**Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. + +This release also introduces several vulnerability fixes: + * Fixes an FPE in TFLite in conv kernel [CVE-2023-27579] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-27579) * Fixes a double free in Fractional(Max/Avg)Pool [CVE-2023-25801] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25801) * Fixes a null dereference on ParallelConcat with XLA [CVE-2023-25676] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25676) From 15a371d6fdae31317db6ca25b1a70955b7bc3451 Mon Sep 17 00:00:00 2001 From: TensorFlow Release Automation Date: Tue, 14 Mar 2023 00:34:33 +0000 Subject: [PATCH 46/52] Update version numbers to 2.11.1 --- tensorflow/core/public/version.h | 2 +- tensorflow/tensorflow.bzl | 2 +- tensorflow/tools/pip_package/setup.py | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/tensorflow/core/public/version.h b/tensorflow/core/public/version.h index 2a94515d20d141..ff912017b11ba8 100644 --- a/tensorflow/core/public/version.h +++ b/tensorflow/core/public/version.h @@ -22,7 +22,7 @@ limitations under the License. // tensorflow/tools/pip_package/setup.py #define TF_MAJOR_VERSION 2 #define TF_MINOR_VERSION 11 -#define TF_PATCH_VERSION 0 +#define TF_PATCH_VERSION 1 // TF_VERSION_SUFFIX is non-empty for pre-releases (e.g. "-alpha", "-alpha.1", // "-beta", "-rc", "-rc.1") diff --git a/tensorflow/tensorflow.bzl b/tensorflow/tensorflow.bzl index 72bd0bcb94aa7f..fa3da78ddcd398 100644 --- a/tensorflow/tensorflow.bzl +++ b/tensorflow/tensorflow.bzl @@ -62,7 +62,7 @@ def register_extension_info(**kwargs): # not contain rc or alpha, only numbers. # Also update tensorflow/core/public/version.h # and tensorflow/tools/pip_package/setup.py -VERSION = "2.11.0" +VERSION = "2.11.1" VERSION_MAJOR = VERSION.split(".")[0] two_gpu_tags = ["requires-gpu-nvidia:2", "notap", "manual", "no_pip"] diff --git a/tensorflow/tools/pip_package/setup.py b/tensorflow/tools/pip_package/setup.py index 74a1962c3344d7..0e9a4623daaef1 100644 --- a/tensorflow/tools/pip_package/setup.py +++ b/tensorflow/tools/pip_package/setup.py @@ -46,7 +46,7 @@ # result for pip. # Also update tensorflow/tensorflow.bzl and # tensorflow/core/public/version.h -_VERSION = '2.11.0' +_VERSION = '2.11.1' # We use the same setup.py for all tensorflow_* packages and for the nightly From 609e5ad545433d4a425fa5e5a07bb482cc6077b3 Mon Sep 17 00:00:00 2001 From: Mihai Maruseac Date: Tue, 14 Mar 2023 16:00:13 -0700 Subject: [PATCH 47/52] Fix 2.11 build --- tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc b/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc index 1b08efb9ee8726..1be22841bfa9e2 100644 --- a/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc +++ b/tensorflow/compiler/mlir/xla/transforms/legalize_tf.cc @@ -5819,11 +5819,11 @@ class ConvertRandomShuffleOp : public OpRewritePattern { LogicalResult matchAndRewrite(TF::RandomShuffleOp op, PatternRewriter &rewriter) const override { auto no_op = [&]() { - rewriter.replaceOp(op, op.getValue()); + rewriter.replaceOp(op, op.value()); return success(); }; - auto input_type = op.getValue().getType().dyn_cast(); + auto input_type = op.value().getType().dyn_cast(); if (!input_type) return failure(); if (input_type.hasStaticShape() && input_type.getNumElements() <= 1) // No shuffling is required, so copy input directly to output. From 2757416dcd4a2d00ea36512c2ffd347030c1196b Mon Sep 17 00:00:00 2001 From: Mihai Maruseac Date: Wed, 15 Mar 2023 08:20:04 -0700 Subject: [PATCH 48/52] Fix bad cherrypick Signed-off-by: Mihai Maruseac --- .../core/tpu/ops/tpu_partitioned_input_op.cc | 81 ------------------- .../core/tpu/ops/tpu_partitioned_output_op.cc | 51 ------------ 2 files changed, 132 deletions(-) diff --git a/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc b/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc index c905664cac95bc..fa43ef51d80d70 100644 --- a/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc +++ b/tensorflow/core/tpu/ops/tpu_partitioned_input_op.cc @@ -124,85 +124,4 @@ REGISTER_OP("TPUPartitionedInput") return OkStatus(); }); -REGISTER_OP("TPUPartitionedInputV2") - .Input("inputs: N * T") - .Output("output: T") - .Attr("N: int >= 1") - .Attr("T: type") - .Attr("partition_dims: list(int)") - .Attr("is_packed: bool = false") - .SetShapeFn([](InferenceContext* c) { - DataType dtype; - TF_RETURN_IF_ERROR(c->GetAttr("T", &dtype)); - std::vector partition_dims; - TF_RETURN_IF_ERROR(c->GetAttr("partition_dims", &partition_dims)); - bool is_packed; - TF_RETURN_IF_ERROR(c->GetAttr("is_packed", &is_packed)); - - int num_partitions = 1; - for (const int& partition_dim : partition_dims) { - num_partitions *= partition_dim; - } - - bool replicated = partition_dims.empty(); - int num_inputs_expected = is_packed ? 1 : num_partitions; - if (!((replicated && !is_packed) || - (c->num_inputs() == num_inputs_expected))) { - // we cannot validate the number of inputs for replicated, unpacked ops - // since we cannot infer the number of partitions from partition_dims - return errors::InvalidArgument("Expected ", num_inputs_expected, - " inputs, got ", c->num_inputs(), "."); - } else if (c->num_inputs() == 0) { - return errors::InvalidArgument( - "Expected at least one input to TPUPartitionedInputV2."); - } - - ShapeHandle output_shape; - if (dtype == DT_RESOURCE) { - ShapeHandle previous_shape_handle; - const std::vector* shapes_and_types = - nullptr; - for (int i = c->num_inputs() - 1; i >= 0; --i) { - shapes_and_types = c->input_handle_shapes_and_types(i); - if (shapes_and_types) { - ShapeHandle shape_handle = shapes_and_types->at(0).shape; - if (!c->FullyDefined(shape_handle)) { - return errors::InvalidArgument("Inputs must have static shape,", - "input[", i, - "] has unknown dimension."); - } - - if (i != c->num_inputs() - 1) { - ShapeHandle tmp; - if (!c->Merge(shape_handle, previous_shape_handle, &tmp).ok()) { - return errors::InvalidArgument( - "Inputs must have the same shape."); - } - } else { - previous_shape_handle = shape_handle; - } - } - } - - if (shapes_and_types) { - TF_ASSIGN_OR_RETURN( - output_shape, - _ComputeOutputShape(c, previous_shape_handle, partition_dims)); - std::vector output_shapes_and_types; - output_shapes_and_types.push_back(shape_inference::ShapeAndType( - output_shape, shapes_and_types->at(0).dtype)); - c->set_output_handle_shapes_and_types(0, output_shapes_and_types); - } - } - - if (!c->FullyDefined(output_shape)) { - TF_ASSIGN_OR_RETURN( - output_shape, _ComputeOutputShape(c, c->input(0), partition_dims)); - } - - c->set_output(0, output_shape); - - return OkStatus(); - }); - } // namespace tensorflow diff --git a/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc b/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc index da9fd60923cd25..aaf08ca1ca0536 100644 --- a/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc +++ b/tensorflow/core/tpu/ops/tpu_partitioned_output_op.cc @@ -70,55 +70,4 @@ REGISTER_OP("TPUPartitionedOutput") return OkStatus(); }); -REGISTER_OP("TPUPartitionedOutputV2") - .Input("inputs: T") - .Output("output: num_splits * T") - .Attr("T: type") - .Attr("num_splits: int >= 1") - .Attr("partition_dims: list(int)") - .SetShapeFn([](InferenceContext* c) { - DataType dtype; - TF_RETURN_IF_ERROR(c->GetAttr("T", &dtype)); - std::vector partition_dims; - TF_RETURN_IF_ERROR(c->GetAttr("partition_dims", &partition_dims)); - int num_splits; - TF_RETURN_IF_ERROR(c->GetAttr("num_splits", &num_splits)); - if (dtype == DT_RESOURCE) { - return errors::Unimplemented("Not implemented."); - } else if (c->num_inputs() == 0) { - return errors::InvalidArgument( - "Expected at least one input to TPUPartitionedOutputV2."); - } - - ShapeHandle handle = c->input(0); - int rank = InferenceContext::Rank(handle); - int num_cores_per_replica = 1; - for (const int& partition_dim : partition_dims) { - num_cores_per_replica *= partition_dim; - } - - if (num_splits != num_cores_per_replica) { - return errors::InvalidArgument("Expected ", num_cores_per_replica, - " splits."); - } else if (rank > (int)partition_dims.size()) { - return errors::InvalidArgument("Expected at least ", rank, - " partition dimensions."); - } - - for (int i = 0; i < rank; ++i) { - shape_inference::DimensionHandle dim; - TF_RETURN_WITH_CONTEXT_IF_ERROR( - c->Divide(c->Dim(handle, i), partition_dims[i], - true /* evenly_divisible */, &dim), - "Number of ways to split should evenly divide the split dimension"); - TF_CHECK_OK(c->ReplaceDim(handle, i, dim, &handle)); - } - - for (int i = num_splits - 1; i >= 0; --i) { - c->set_output(i, handle); - } - - return OkStatus(); - }); - } // namespace tensorflow From a2ba9f16f0154bf93f21132878b154238d89fad6 Mon Sep 17 00:00:00 2001 From: venkat2469 <104534215+venkat2469@users.noreply.github.com> Date: Wed, 15 Mar 2023 17:54:33 +0000 Subject: [PATCH 49/52] Updating Release.md with Legal Language for Release Notes --- RELEASE.md | 1 + 1 file changed, 1 insertion(+) diff --git a/RELEASE.md b/RELEASE.md index 8ff21eba50b13f..109f873d3494b9 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -1,6 +1,7 @@ # Release 2.11.1 **Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. +* Security vulnerability fixes will no longer be patched to this Tensorflow version. The latest Tensorflow version includes the security vulnerability fixes. You can update to the latest version (recommended) or patch security vulnerabilities yourself (steps). You can refer to the release notes of the latest Tensorflow version for a list of newly fixed vulnerabilities. If you have any questions, please create a GitHub issue to let us know. This release also introduces several vulnerability fixes: From 93dea7a67df44bde557e580dfdcde5ba0a7a344d Mon Sep 17 00:00:00 2001 From: venkat2469 <104534215+venkat2469@users.noreply.github.com> Date: Wed, 15 Mar 2023 11:16:49 -0700 Subject: [PATCH 50/52] Update RELEASE.md --- RELEASE.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/RELEASE.md b/RELEASE.md index 109f873d3494b9..08ea106f941dcc 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -1,7 +1,7 @@ # Release 2.11.1 **Note**: TensorFlow 2.10 was the last TensorFlow release that supported GPU on native-Windows. Starting with TensorFlow 2.11, you will need to install TensorFlow in WSL2, or install tensorflow-cpu and, optionally, try the TensorFlow-DirectML-Plugin. -* Security vulnerability fixes will no longer be patched to this Tensorflow version. The latest Tensorflow version includes the security vulnerability fixes. You can update to the latest version (recommended) or patch security vulnerabilities yourself (steps). You can refer to the release notes of the latest Tensorflow version for a list of newly fixed vulnerabilities. If you have any questions, please create a GitHub issue to let us know. +* Security vulnerability fixes will no longer be patched to this Tensorflow version. The latest Tensorflow version includes the security vulnerability fixes. You can update to the latest version (recommended) or patch security vulnerabilities yourself [steps](https://github.com/tensorflow/tensorflow#patching-guidelines). You can refer to the [release notes](https://github.com/tensorflow/tensorflow/releases) of the latest Tensorflow version for a list of newly fixed vulnerabilities. If you have any questions, please create a GitHub issue to let us know. This release also introduces several vulnerability fixes: From eea48f50d6982879909bf8e0d0151bbce3f9bf4a Mon Sep 17 00:00:00 2001 From: Mihai Maruseac Date: Thu, 16 Mar 2023 09:04:00 -0700 Subject: [PATCH 51/52] Disable a test that results in OOM+segfault Signed-off-by: Mihai Maruseac --- tensorflow/python/ops/histogram_ops_test.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tensorflow/python/ops/histogram_ops_test.py b/tensorflow/python/ops/histogram_ops_test.py index 629507052dd206..bb2800f79c1d65 100644 --- a/tensorflow/python/ops/histogram_ops_test.py +++ b/tensorflow/python/ops/histogram_ops_test.py @@ -176,7 +176,8 @@ def test_shape_inference(self): self.assertAllClose(expected_bin_counts, hist.eval({placeholder: 5})) - def test_large_range(self): + # This test is disabled on 2.11 as it OOMs + segfaults afterwards + def DISABLED_test_large_range(self): hist = histogram_ops.histogram_fixed_width( values=constant_op.constant( [-(2**31), 2**31 - 1], dtype=dtypes.int32 From 13b85dcf966d0c94b2e5c21291be039db2dec7b9 Mon Sep 17 00:00:00 2001 From: Mihai Maruseac Date: Thu, 16 Mar 2023 09:51:51 -0700 Subject: [PATCH 52/52] Fix release notes Signed-off-by: Mihai Maruseac --- RELEASE.md | 38 +++++++++++++++++++------------------- 1 file changed, 19 insertions(+), 19 deletions(-) diff --git a/RELEASE.md b/RELEASE.md index 08ea106f941dcc..80aad26d6dff87 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -5,25 +5,25 @@ This release also introduces several vulnerability fixes: -* Fixes an FPE in TFLite in conv kernel [CVE-2023-27579] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-27579) -* Fixes a double free in Fractional(Max/Avg)Pool [CVE-2023-25801] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25801) -* Fixes a null dereference on ParallelConcat with XLA [CVE-2023-25676] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25676) -* Fixes a segfault in Bincount with XLA [CVE-2023-25675] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25675) -* Fixes an NPE in RandomShuffle with XLA enable [CVE-2023-25674] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25674) -* Fixes an FPE in TensorListSplit with XLA [CVE-2023-25673] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25673) -* Fixes segmentation fault in tfg-translate [CVE-2023-25671] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25671) -* Fixes an NPE in QuantizedMatMulWithBiasAndDequantize [CVE-2023-25670] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25670) -* Fixes an FPE in AvgPoolGrad with XLA [CVE-2023-25669] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25669) -* Fixes a heap out-of-buffer read vulnerability in the QuantizeAndDequantize operation [CVE-2023-25668] (https://cve.mitre.org/cgi-bin/cvename.cgi? name=CVE-2023-25668) -* Fixes a segfault when opening multiframe gif [CVE-2023-25667] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25667) -* Fixes an NPE in SparseSparseMaximum [CVE-2023-25665] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25665) -* Fixes an FPE in AudioSpectrogram [CVE-2023-25666] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25666) -* Fixes a heap-buffer-overflow in AvgPoolGrad [CVE-2023-25664] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25664) -* Fixes a NPE in TensorArrayConcatV2 [CVE-2023-25663] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25663) -* Fixes a Integer overflow in EditDistance [CVE-2023-25662] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25662) -* Fixes a Seg fault in `tf.raw_ops.Print` [CVE-2023-25660] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25660) -* Fixes a OOB read in DynamicStitch [CVE-2023-25659] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25659) -* Fixes a OOB Read in GRUBlockCellGrad [CVE-2023-25658] (https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25658) +* Fixes an FPE in TFLite in conv kernel [CVE-2023-27579](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-27579) +* Fixes a double free in Fractional(Max/Avg)Pool [CVE-2023-25801](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25801) +* Fixes a null dereference on ParallelConcat with XLA [CVE-2023-25676](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25676) +* Fixes a segfault in Bincount with XLA [CVE-2023-25675](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25675) +* Fixes an NPE in RandomShuffle with XLA enable [CVE-2023-25674](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25674) +* Fixes an FPE in TensorListSplit with XLA [CVE-2023-25673](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25673) +* Fixes segmentation fault in tfg-translate [CVE-2023-25671](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25671) +* Fixes an NPE in QuantizedMatMulWithBiasAndDequantize [CVE-2023-25670](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25670) +* Fixes an FPE in AvgPoolGrad with XLA [CVE-2023-25669](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25669) +* Fixes a heap out-of-buffer read vulnerability in the QuantizeAndDequantize operation [CVE-2023-25668](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25668) +* Fixes a segfault when opening multiframe gif [CVE-2023-25667](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25667) +* Fixes an NPE in SparseSparseMaximum [CVE-2023-25665](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25665) +* Fixes an FPE in AudioSpectrogram [CVE-2023-25666](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25666) +* Fixes a heap-buffer-overflow in AvgPoolGrad [CVE-2023-25664](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25664) +* Fixes a NPE in TensorArrayConcatV2 [CVE-2023-25663](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25663) +* Fixes a Integer overflow in EditDistance [CVE-2023-25662](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25662) +* Fixes a Seg fault in `tf.raw_ops.Print` [CVE-2023-25660](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25660) +* Fixes a OOB read in DynamicStitch [CVE-2023-25659](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25659) +* Fixes a OOB Read in GRUBlockCellGrad [CVE-2023-25658](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2023-25658) # Release 2.11.0