CVE-2023-25658: Out-of-bounds Read (CWE-125)
value: TensorFlow is an open source platform for machine learning. Prior to versions 2.12.0 and 2.11.1, an out of bounds read is in GRUBlockCellGrad. A fix is included in TensorFlow 2.12.0 and 2.11.1.

CWE description: The product reads data past the end, or before the beginning, of the intended buffer.
CVSS v3 base score: 7.5
Published: 2023-03-25T00:15Z
Repository: https://github.com/tensorflow/tensorflow
Language: C++
Files changed: .bazelrc, tensorflow/compiler/jit/BUILD, tensorflow/compiler/xla/backends/interpreter/BUILD, tensorflow/compiler/xla/backends/interpreter/platform.cc, tensorflow/compiler/xla/hlo/evaluator/hlo_evaluator.cc, tensorflow/compiler/xla/mlir/tools/mlir_replay/BUILD, tensorflow/compiler/xla/service/hlo_graph_dumper.cc, tensorflow/compiler/xla/stream_executor/BUILD, tensorflow/compiler/xla/stream_executor/cuda/cuda_blas.cc, tensorflow/compiler/xla/stream_executor/cuda/cuda_dnn.cc, tensorflow/compiler/xla/stream_executor/cuda/cuda_driver.cc, tensorflow/compiler/xla/stream_executor/cuda/cuda_fft.cc, tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.cc, tensorflow/compiler/xla/stream_executor/cuda/cuda_platform.cc, tensorflow/compiler/xla/stream_executor/cuda/cuda_rng.cc, tensorflow/compiler/xla/stream_executor/host/BUILD, tensorflow/compiler/xla/stream_executor/host/host_gpu_executor.h, tensorflow/compiler/xla/stream_executor/host/host_platform.cc, tensorflow/compiler/xla/stream_executor/lib/BUILD, tensorflow/compiler/xla/stream_executor/multi_platform_manager.cc, tensorflow/compiler/xla/stream_executor/multi_platform_manager.h, tensorflow/compiler/xla/stream_executor/platform.cc, tensorflow/compiler/xla/stream_executor/plugin_registry.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_blas.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_diagnostics.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_dnn.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_driver.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_fft.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_gpu_executor.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_platform.cc, tensorflow/compiler/xla/stream_executor/rocm/rocm_rng.cc, tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.cc, tensorflow/compiler/xla/stream_executor/tf_allocator_adapter.cc, tensorflow/core/common_runtime/function.cc, tensorflow/core/data/BUILD, tensorflow/core/data/service/server_lib.cc, tensorflow/core/data/service/server_lib.h, tensorflow/core/data/service/snapshot/BUILD, tensorflow/core/data/service/snapshot/path_utils.cc, tensorflow/core/data/service/snapshot/path_utils.h, tensorflow/core/data/service/snapshot/path_utils_test.cc, tensorflow/core/data/service/snapshot/snapshot_manager.cc, tensorflow/core/data/tfdataz_metrics.cc, tensorflow/core/data/tfdataz_metrics.h, tensorflow/core/data/tfdataz_metrics_test.cc, tensorflow/core/distributed_runtime/eager/remote_mgr.cc, tensorflow/core/distributed_runtime/integration_test/coordination_test_opkernel_registration.cc, tensorflow/core/framework/dataset.h, tensorflow/core/framework/op_requires.h, tensorflow/core/kernels/data/iterator_ops.cc, tensorflow/core/kernels/function_ops.cc, tensorflow/core/kernels/functional_ops.cc, tensorflow/core/kernels/rnn/gru_ops.cc, tensorflow/core/lib/core/status_test.cc, tensorflow/core/platform/error_payloads.cc, tensorflow/core/tpu/kernels/tpu_compile_op_common.cc, tensorflow/core/tpu/kernels/tpu_functional_ops.cc, tensorflow/core/tpu/tpu_embedding_errors.cc, tensorflow/core/tpu/tpu_embedding_errors.h, tensorflow/core/util/zen_util.h, tensorflow/lite/delegates/xnnpack/README.md, tensorflow/python/data/experimental/kernel_tests/service/snapshot_ft_test.py, tensorflow/python/framework/BUILD, tensorflow/python/framework/errors_test_helper.cc, tensorflow/python/tpu/tpu_strategy_util.py, tensorflow/tools/ci_build/release/requirements_common.txt, tensorflow/tools/ci_build/release/requirements_mac.txt, tensorflow/tools/tf_sig_build_dockerfiles/devel.requirements.txt, tensorflow/tsl/c/tsl_status.cc, tensorflow/tsl/c/tsl_status_helper_test.cc, tensorflow/tsl/distributed_runtime/coordination/coordination_service_error_util.h, tensorflow/tsl/distributed_runtime/rpc/grpc_util.h, tensorflow/tsl/distributed_runtime/rpc/grpc_util_test.cc, tensorflow/tsl/platform/errors.h, tensorflow/tsl/platform/test.h, tensorflow/tsl/profiler/lib/BUILD, tensorflow/tsl/profiler/lib/traceme_encode.h, tensorflow/tsl/profiler/lib/traceme_encode_test.cc, third_party/tf_runtime/workspace.bzl
Commit date: 2023-01-25 16:05:35 -0800
Commit message: Merged commit includes the following changes:
504684855  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Remove legacy references from `ops.py`.

    This is done to eventually remove the lazy loads in `indexed_slices.py`.

--
504682660  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Add missing reference symbol.

--
504682467  by A. Unique TensorFlower<gardener@tensorflow.org>:

    #tf-data-service Add a function to parse split file names.

--
504680829  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Update XNNPACK README to document feature parity for FP16 operators

--
504677404  by A. Unique TensorFlower<gardener@tensorflow.org>:
    Automated rollback of changelist 504103068.

504676634  by A. Unique TensorFlower<gardener@tensorflow.org>:
    Automated rollback of changelist 504656381.

504672100  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Delete compiler/xla/stream_executor/lib/initialize.h, update users to depend on stream_executor/platform/intialize.h directly

--
504669543  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Refactor transforms to make buildozer object injectable.

--
504667726  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Update all CI dependencies to non vulnerable versions.

    This MUST land before branch cut.

--
504661218  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Update TFRT dependency to use revision
    http://github.com/tensorflow/runtime/commit/c653281a1a23c0c3d41536a983c7d10fcc5b1fbf.

--
504660212  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Change TraceMeArg to keep a reference to an AlphaNum constructed on the stack rather than constructing it by value.

    The eliminates a lifetime safety bug where AlphaNum uses a default constructed argument as a string buffer that was a dangling reference before.

    Also add a LIFETIME_BOUND attribute to ensure this is used correctly.

--
504660069  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Add kIsOpenSource to tsl/platform/test.h

--
504659330  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Update usages of tensorflow::Status::SetPayload.

--
504656381  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Remove the scrubbing of constants in tpu_cluster_formation.

    Instead, just ignore tf.Const attributes, in CollectAndGroupClusterOps. This
    is functionally equivalent (for this pass), but a bit more clean since it
    doesn't leave modified tf.Const ops behind.

--
504654874  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Add missing dependency for the Mac build.

--
504653641  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Internal visibility change

--
504648821  by A. Unique TensorFlower<gardener@tensorflow.org>:
    Automated rollback of changelist 504006190.
BEGIN_PUBLIC
Rollback of PR #59418: [aarch64] enable acl linking for cpu xla for mkl_aarch64_threadpool config
Reason: The configuration is still not very well-tested.
END_PUBLIC

504645490  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Fix input validation for GRU ops.

--
504633571  by A. Unique TensorFlower<gardener@tensorflow.org>:

    [SE] Replace uses of stream_executor::port::error with tsl::error.

    This lets us remove stream_executor/lib/error.h!

    No functional change; these are aliases of one another.

--
504633188  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Update tf.tpu.experimental.initialize_tpu_system to disable running on eager mode.

--
504618714  by A. Unique TensorFlower<gardener@tensorflow.org>:

    [XLA:GPU] Update hlo graph dumper to additionally dump the gemm epilogue from the GemmBackendConfig

--
504618661  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Fix ODR violation from zen_util.h

--
504616623  by A. Unique TensorFlower<gardener@tensorflow.org>:

    [tf.data] Collect the iterator's memory usage metrics for /tfdataz page.

--
504612989  by A. Unique TensorFlower<gardener@tensorflow.org>:

    Internal change

PiperOrigin-RevId: 504684855

Diff: diff --git a/.bazelrc b/.bazelrc
index 8fc7a73f224fd..930313e8c2589 100644
--- a/.bazelrc
+++ b/.bazelrc
@@ -240,11 +240,10 @@ build:mkl_aarch64 --define=build_with_acl=true
 build:mkl_aarch64 -c opt
 
 # Config setting to build oneDNN with Compute Library for the Arm Architecture (ACL).
 # with Eigen threadpool support
 build:mkl_aarch64_threadpool --define=build_with_mkl_aarch64=true
-build:mkl_aarch64_threadpool --define=build_with_acl=true
 build:mkl_aarch64_threadpool -c opt
 
 # This config refers to building CUDA op kernels with nvcc.
 build:cuda --repo_env TF_NEED_CUDA=1
 build:cuda --crosstool_top=@local_config_cuda//crosstool:toolchain
diff --git a/tensorflow/compiler/jit/BUILD b/tensorflow/compiler/jit/BUILD
index 2c90c32bf77b0..70b07f90dbd74 100644
--- a/tensorflow/compiler/jit/BUILD
+++ b/tensorflow/compiler/jit/BUILD
@@ -409,10 +409,11 @@ cc_library(
     hdrs = ["xla_launch_util.h"],
     visibility = [
         ":internal",
         # We reuse VariableInfo in TFRT's implementation of TpuExecuteOp.
         "//learning/brain/tfrt/tf_tpu:__pkg__",
+        "//learning/brain/tfrt/tpu_plugin:__pkg__",
         "//learning/brain/tfrt/tpu_common:__pkg__",
         "//tensorflow/core/common_runtime/next_pluggable_device:__pkg__",
     ],
     deps = [
         ":common",
diff --git a/tensorflow/compiler/xla/backends/interpreter/BUILD b/tensorflow/compiler/xla/backends/interpreter/BUILD
index 137e83b9d7efa..f431901a4e865 100644
--- a/tensorflow/compiler/xla/backends/interpreter/BUILD
+++ b/tensorflow/compiler/xla/backends/interpreter/BUILD
@@ -136,11 +136,11 @@ cc_library(
     hdrs = ["platform.h"],
     deps = [
         ":executor",
         ":platform_id",
         "//tensorflow/compiler/xla/stream_executor",
-        "//tensorflow/compiler/xla/stream_executor/lib",
+        "//tensorflow/compiler/xla/stream_executor/platform",
         "//tensorflow/tsl/platform:status",
         "@com_google_absl//absl/strings:str_format",
     ],
     alwayslink = True,  # Registers itself with the MultiPlatformManager.
 )
diff --git a/tensorflow/compiler/xla/backends/interpreter/platform.cc b/tensorflow/compiler/xla/backends/interpreter/platform.cc
index 2d60804f7f29c..0259baf822148 100644
--- a/tensorflow/compiler/xla/backends/interpreter/platform.cc
+++ b/tensorflow/compiler/xla/backends/interpreter/platform.cc
@@ -19,13 +19,13 @@ limitations under the License.
 #include <utility>
 
 #include "absl/strings/str_format.h"
 #include "tensorflow/compiler/xla/backends/interpreter/executor.h"
 #include "tensorflow/compiler/xla/stream_executor/device_options.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/multi_platform_manager.h"
 #include "tensorflow/compiler/xla/stream_executor/platform.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/tsl/platform/status.h"
 
 namespace stream_executor {
 namespace interpreter {
 
diff --git a/tensorflow/compiler/xla/hlo/evaluator/hlo_evaluator.cc b/tensorflow/compiler/xla/hlo/evaluator/hlo_evaluator.cc
index fdabe23557c5c..3dea63111adcb 100644
--- a/tensorflow/compiler/xla/hlo/evaluator/hlo_evaluator.cc
+++ b/tensorflow/compiler/xla/hlo/evaluator/hlo_evaluator.cc
@@ -220,11 +220,11 @@ Status MakeEvalErrorDueToParamOrInfeed(const HloInstruction& eval_instruction) {
   std::string error_payload;
   error_payload.resize(sizeof(EvalErrorDetail));
   absl::little_endian::Store32(
       const_cast<char*>(error_payload.data()),
       static_cast<uint32_t>(EvalErrorDetail::kDynamicValueDependence));
-  error.SetPayload(kEvalErrorDetailUrl, error_payload);
+  error.SetPayload(kEvalErrorDetailUrl, absl::Cord(error_payload));
   return error;
 }
 
 // Repesents a value that might or might not be determined statically.
 struct DynamicOrStaticInteger {
diff --git a/tensorflow/compiler/xla/mlir/tools/mlir_replay/BUILD b/tensorflow/compiler/xla/mlir/tools/mlir_replay/BUILD
index bcd3c5d4cff6d..e7ac4534382f9 100644
--- a/tensorflow/compiler/xla/mlir/tools/mlir_replay/BUILD
+++ b/tensorflow/compiler/xla/mlir/tools/mlir_replay/BUILD
@@ -7,10 +7,11 @@ xla_cc_binary(
     deps = [
         ":mlir_replay_lib",
         "//tensorflow/compiler/xla:debug_options_flags",
         "//tensorflow/compiler/xla/mlir/runtime/ir:rt",
         "//tensorflow/compiler/xla/mlir/tools/mlir_replay/public:compiler_trace_proto_cc",
+        "//tensorflow/compiler/xla/mlir/tools/mlir_replay/public:compiler_trace_proto_cc_impl",
         "//tensorflow/compiler/xla/mlir/tools/mlir_replay/public:execution_trace_proto_cc",
         "//tensorflow/compiler/xla/mlir_hlo:gml_st",
         "//tensorflow/compiler/xla/mlir_hlo:hlo_dialect_registration",
         "//tensorflow/compiler/xla/mlir_hlo:lhlo",
         "//tensorflow/compiler/xla/mlir_hlo:lhlo_gpu",
diff --git a/tensorflow/compiler/xla/service/hlo_graph_dumper.cc b/tensorflow/compiler/xla/service/hlo_graph_dumper.cc
index 9e72f7b87e29c..40d9513bb8165 100644
--- a/tensorflow/compiler/xla/service/hlo_graph_dumper.cc
+++ b/tensorflow/compiler/xla/service/hlo_graph_dumper.cc
@@ -1249,10 +1249,14 @@ ExtractGemmBackendConfigProps(const gpu::GemmBackendConfig& config,
               DotDimensionNumbersToString(config.dot_dimension_numbers()),
               {{", ", "<br/>"}}));
   if (config.algorithm_case() == gpu::GemmBackendConfig::kSelectedAlgorithm) {
     props.emplace_back("algorithm", StrCat(config.selected_algorithm()));
   }
+  if (config.epilogue() != gpu::GemmBackendConfig::DEFAULT) {
+    props.emplace_back(
+        "epilogue", gpu::GemmBackendConfig::Epilogue_Name(config.epilogue()));
+  }
   return props;
 }
 
 std::string HloDotDumper::GetInstructionNodeBackendConfig(
     const HloInstruction* instr) {
diff --git a/tensorflow/compiler/xla/stream_executor/BUILD b/tensorflow/compiler/xla/stream_executor/BUILD
index c0b1ea16b72c7..7c5d91212269f 100644
--- a/tensorflow/compiler/xla/stream_executor/BUILD
+++ b/tensorflow/compiler/xla/stream_executor/BUILD
@@ -237,10 +237,11 @@ cc_library(
     deps = [
         ":plugin",
         ":stream_executor_headers",
         "//tensorflow/compiler/xla/stream_executor/lib",
         "//tensorflow/compiler/xla/stream_executor/platform",
+        "//tensorflow/tsl/platform:errors",
         "//tensorflow/tsl/platform:status",
         "//tensorflow/tsl/platform:statusor",
         "@com_google_absl//absl/strings",
     ],
 )
@@ -448,10 +449,11 @@ tsl_gpu_library(
         ":temporary_memory_manager",
         ":timer",
         "//tensorflow/compiler/xla/stream_executor/lib",
         "//tensorflow/compiler/xla/stream_executor/platform",
         "//tensorflow/tsl/platform:env",
+        "//tensorflow/tsl/platform:errors",
         "//tensorflow/tsl/platform:logging",
         "//tensorflow/tsl/platform:stacktrace",
         "//tensorflow/tsl/platform:status",
         "//tensorflow/tsl/platform:statusor",
         "//tensorflow/tsl/protobuf:dnn_proto_cc",
@@ -659,11 +661,11 @@ cc_library(
         ":fft",
         ":multi_platform_manager",
         ":platform",
         ":plugin",
         ":stream_executor_headers",
-        "//tensorflow/compiler/xla/stream_executor/lib",
+        "//tensorflow/tsl/platform:errors",
         "//tensorflow/tsl/platform:status",
         "//tensorflow/tsl/platform:statusor",
         "@com_google_absl//absl/base:core_headers",
         "@com_google_absl//absl/strings",
         "@com_google_absl//absl/strings:str_format",
diff --git a/tensorflow/compiler/xla/stream_executor/cuda/cuda_blas.cc b/tensorflow/compiler/xla/stream_executor/cuda/cuda_blas.cc
index fd5e483b45db0..6b84c0444c1f8 100644
--- a/tensorflow/compiler/xla/stream_executor/cuda/cuda_blas.cc
+++ b/tensorflow/compiler/xla/stream_executor/cuda/cuda_blas.cc
@@ -61,11 +61,11 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_helpers.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_timer.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_types.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/scratch_allocator.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor.h"
diff --git a/tensorflow/compiler/xla/stream_executor/cuda/cuda_dnn.cc b/tensorflow/compiler/xla/stream_executor/cuda/cuda_dnn.cc
index fb0ca7fcdf7c0..c331220602bc0 100644
--- a/tensorflow/compiler/xla/stream_executor/cuda/cuda_dnn.cc
+++ b/tensorflow/compiler/xla/stream_executor/cuda/cuda_dnn.cc
@@ -34,12 +34,11 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_timer.h"
 #include "tensorflow/compiler/xla/stream_executor/dnn.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/scratch_allocator.h"
 #include "tensorflow/compiler/xla/stream_executor/stream.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_internal.h"
@@ -83,22 +82,22 @@ static_assert(CUDNN_VERSION >= 7300, "cuDNN needs to be version 7.3 or higher");
     cudnnStatus_t _status = (expr);                                     \
     if (!ABSL_PREDICT_TRUE(_status == CUDNN_STATUS_SUCCESS)) {          \
       std::ostringstream oss;                                           \
       oss << CudnnStatusToString(_status) << "\nin " << __FILE__ << "(" \
           << __LINE__ << "): '" << #expr << "'";                        \
-      return tsl::Status(port::error::UNKNOWN, oss.str());              \
+      return tsl::Status(tsl::error::UNKNOWN, oss.str());               \
     }                                                                   \
   } while (false)
 
 #define RETURN_MSG_IF_CUDNN_ERROR(expr)                                 \
   do {                                                                  \
     cudnnStatus_t _status = (expr).get_status();                        \
     if (!ABSL_PREDICT_TRUE(_status == CUDNN_STATUS_SUCCESS)) {          \
       std::ostringstream oss;                                           \
       oss << CudnnStatusToString(_status) << "\nin " << __FILE__ << "(" \
           << __LINE__ << "): '" << #expr << "' " << (expr).get_error(); \
-      return tsl::Status(port::error::UNKNOWN, oss.str());              \
+      return tsl::Status(tsl::error::UNKNOWN, oss.str());               \
     }                                                                   \
   } while (false)
 
 #define RETURN_FALSE_IF_CUDNN_ERROR(expr)                                  \
   do {                                                                     \
@@ -415,11 +414,11 @@ tsl::Status CudnnSupport::Init() {
           "library.  If building from sources, make sure the library loaded at "
           "runtime is compatible with the version specified during compile "
           "configuration.");
       LOG(ERROR) << error;
       cudnnDestroy(cudnn_handle);
-      return tsl::Status(port::error::INTERNAL, error);
+      return tsl::Status(tsl::error::INTERNAL, error);
     }
 
     cudnn_.reset(new CudnnAccess(cudnn_handle));
 
     LOG(INFO) << "Loaded cuDNN version " << cudnnGetVersion();
@@ -439,11 +438,11 @@ tsl::Status CudnnSupport::Init() {
       LOG(ERROR) << "Possibly insufficient driver version: "
                  << cuda::DriverVersionToString(version);
     }
   }
 
-  return tsl::Status(port::error::INTERNAL,
+  return tsl::Status(tsl::error::INTERNAL,
                      absl::StrCat("cudnn library could not create a handle: ",
                                   CudnnStatusToString(status)));
 }
 
 void CudnnSupport::NotifyStreamDestroyed(Stream* stream) /* override */ {
@@ -1297,11 +1296,11 @@ class CudnnRnnDescriptor : public dnn::RnnDescriptor {
     bool use_tensor_ops =
         algorithm_config.algorithm().has_value()
             ? algorithm_config.algorithm()->tensor_ops_enabled()
             : allow_tensor_ops;
     if (use_tensor_ops && !allow_tensor_ops) {
-      return tsl::Status(port::error::INVALID_ARGUMENT,
+      return tsl::Status(tsl::error::INVALID_ARGUMENT,
                          "Algo requests disallowed tensor op evaluation.");
     }
 
 #if CUDNN_VERSION >= 8000
     cudnnMathType_t math_type =
@@ -1656,11 +1655,11 @@ class CudnnRnnSequenceTensorDescriptor
 
   static tsl::StatusOr<CudnnRnnSequenceTensorDescriptor> Create(
       GpuExecutor* parent, int max_seq_length, int batch_size, int data_size,
       cudnnDataType_t data_type) {
     if (max_seq_length <= 0) {
-      return tsl::Status(port::error::INVALID_ARGUMENT, "max_seq_length <= 0");
+      return tsl::Status(tsl::error::INVALID_ARGUMENT, "max_seq_length <= 0");
     }
     int dims[] = {batch_size, data_size, 1};
     int strides[] = {dims[1] * dims[2], dims[2], 1};
     TensorDescriptor tensor_desc = CreateTensorDescriptor();
     RETURN_IF_CUDNN_ERROR(cudnnSetTensorNdDescriptor(
@@ -1675,11 +1674,11 @@ class CudnnRnnSequenceTensorDescriptor
   static tsl::StatusOr<CudnnRnnSequenceTensorDescriptor> Create(
       GpuExecutor* parent, int max_seq_length, int batch_size, int data_size,
       const absl::Span<const int>& seq_lengths, bool time_major,
       cudnnDataType_t data_type) {
     if (max_seq_length <= 0) {
-      return tsl::Status(port::error::INVALID_ARGUMENT, "max_seq_length <= 0");
+      return tsl::Status(tsl::error::INVALID_ARGUMENT, "max_seq_length <= 0");
     }
     int dims[] = {batch_size, data_size, 1};
     int strides[] = {dims[1] * dims[2], dims[2], 1};
     TensorDescriptor tensor_desc = CreateTensorDescriptor();
     RETURN_IF_CUDNN_ERROR(cudnnSetTensorNdDescriptor(
@@ -1802,34 +1801,34 @@ tsl::StatusOr<RnnModelDims> ExtractAndCheckRnnForward(
   // check parameters
   if (!(input_h_desc.num_layers() ==
             model_dims.num_layers * model_dims.dir_count &&
         input_h_desc.batch_size() == model_dims.batch_size &&
         input_h_desc.data_size() == model_dims.hidden_size)) {
-    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid input_h shape");
+    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid input_h shape");
   }
   // The LSTM projection will be used if input_h_desc.data_size() <
   // input_c_desc.data_size()
   if (!(input_h_desc.num_layers() == input_c_desc.num_layers() &&
         input_h_desc.batch_size() == input_c_desc.batch_size() &&
         input_h_desc.data_size() <= input_c_desc.data_size())) {
-    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid input_c shape");
+    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid input_c shape");
   }
   if (!(output_desc.max_seq_length() == model_dims.max_seq_length &&
         output_desc.batch_size() == model_dims.batch_size &&
         output_desc.data_size() ==
             model_dims.hidden_size * model_dims.dir_count)) {
-    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid output shape");
+    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid output shape");
   }
   if (!(input_h_desc.num_layers() == output_h_desc.num_layers() &&
         input_h_desc.batch_size() == output_h_desc.batch_size() &&
         input_h_desc.data_size() == output_h_desc.data_size())) {
-    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid output_h shape");
+    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid output_h shape");
   }
   if (!(input_h_desc.num_layers() == output_c_desc.num_layers() &&
         input_h_desc.batch_size() == output_c_desc.batch_size() &&
         input_h_desc.data_size() <= output_c_desc.data_size())) {
-    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid output_c shape");
+    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid output_c shape");
   }
 
   return model_dims;
 }
 
@@ -1847,11 +1846,11 @@ tsl::Status CheckRNNParameterSize(
       /*xDesc=*/input_desc.handles()[0], /*sizeInBytes=*/&params_size_in_bytes,
       /*dataType=*/rnn_desc.data_type()));
 #endif
   if (static_cast<int64_t>(params_size_in_bytes) !=
       rnn_desc.ParamsSizeInBytes()) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "Mismatching RNN parameter size");
   }
   return ::tsl::OkStatus();
 }
 
@@ -1995,11 +1994,11 @@ tsl::Status CudnnSupport::DoRnnForwardImpl(
       timer.reset(new GpuTimer(parent_));
       // The start and stop of the timer should be as close to the Cudnn call as
       // possible. It is still possible for other threads to issue workload on
       // to this stream. So it could take multiple profiling measurements.
       if (!timer->Init() || !timer->Start(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
       }
     }
 
     RETURN_IF_CUDNN_ERROR(cudnnRNNForward(
         /*handle=*/cudnn.handle(), /*rnnDesc=*/rnn_desc.handle(),
@@ -2018,11 +2017,11 @@ tsl::Status CudnnSupport::DoRnnForwardImpl(
         /*reserveSpaceSizeInBytes=*/reserve_space.size(),
         /*reserveSpace=*/reserve_space.opaque()));
 
     if (is_profiling) {
       if (!timer->Stop(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
       }
       auto algo_desc = *rnn_desc.algorithm_config().algorithm();
       output_profile_result->set_algorithm(algo_desc);
       output_profile_result->set_elapsed_time_in_ms(
           timer->GetElapsedMilliseconds());
@@ -2056,11 +2055,11 @@ tsl::Status CudnnSupport::DoRnnForwardImpl(
     timer.reset(new GpuTimer(parent_));
     // The start and stop of the timer should be as close to the Cudnn call as
     // possible. It is still possible for other threads to issue workload on
     // to this stream. So it could take multiple profiling measurements.
     if (!timer->Init() || !timer->Start(AsGpuStream(stream))) {
-      return tsl::Status(port::error::INTERNAL, "Failed to start timer");
+      return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
     }
   }
 
   if (!is_training) {
     if (input_desc.is_var_seq_lengths()) {
@@ -2128,11 +2127,11 @@ tsl::Status CudnnSupport::DoRnnForwardImpl(
     }
   }
 
   if (is_profiling) {
     if (!timer->Stop(AsGpuStream(stream))) {
-      return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
+      return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
     }
     auto algo_desc = *rnn_desc.algorithm_config().algorithm();
     output_profile_result->set_algorithm(algo_desc);
     output_profile_result->set_elapsed_time_in_ms(
         timer->GetElapsedMilliseconds());
@@ -2202,11 +2201,11 @@ tsl::Status CudnnSupport::DoRnnBackwardImpl(
       timer.reset(new GpuTimer(parent_));
       // The start and stop of the timer should be as close to the Cudnn call as
       // possible. It is still possible for other threads to issue workload on
       // to this stream. So it could take multiple profiling measurements.
       if (!timer->Init() || !timer->Start(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
       }
     }
 
     RETURN_IF_CUDNN_ERROR(cudnnRNNBackwardData_v8(
         /*handle=*/cudnn.handle(), /*rnnDesc=*/rnn_desc.handle(),
@@ -2251,11 +2250,11 @@ tsl::Status CudnnSupport::DoRnnBackwardImpl(
           /*reserveSpace=*/reserve_space_data->opaque()));
     }
 
     if (is_profiling) {
       if (!timer->Stop(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
       }
       auto algo_desc = *rnn_desc.algorithm_config().algorithm();
       output_profile_result->set_algorithm(algo_desc);
       output_profile_result->set_elapsed_time_in_ms(
           timer->GetElapsedMilliseconds());
@@ -2273,11 +2272,11 @@ tsl::Status CudnnSupport::DoRnnBackwardImpl(
     timer.reset(new GpuTimer(parent_));
     // The start and stop of the timer should be as close to the Cudnn call as
     // possible. It is still possible for other threads to issue workload on
     // to this stream. So it could take multiple profiling measurements.
     if (!timer->Init() || !timer->Start(AsGpuStream(stream))) {
-      return tsl::Status(port::error::INTERNAL, "Failed to start timer");
+      return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
     }
   }
 
   if (input_desc.is_var_seq_lengths()) {
     RETURN_IF_CUDNN_ERROR(cudnnRNNBackwardDataEx(
@@ -2360,11 +2359,11 @@ tsl::Status CudnnSupport::DoRnnBackwardImpl(
     }
   }
 
   if (is_profiling) {
     if (!timer->Stop(AsGpuStream(stream))) {
-      return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
+      return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
     }
     auto algo_desc = *rnn_desc.algorithm_config().algorithm();
     output_profile_result->set_algorithm(algo_desc);
     output_profile_result->set_elapsed_time_in_ms(
         timer->GetElapsedMilliseconds());
@@ -2402,11 +2401,11 @@ tsl::Status CudnnSupport::DoCtcLossImpl(
       /*algo=*/ctc_loss_algo,
       /*ctcLossDesc=*/ctc_loss_desc.handle(),
       /*workspace=*/scratch_memory.opaque(),
       /*workSpaceSizeInBytes=*/scratch_memory.size()));
 #else
-  return tsl::Status(port::error::INVALID_ARGUMENT,
+  return tsl::Status(tsl::error::INVALID_ARGUMENT,
                      "No supported cudnnCTCLoss when "
                      "CUDNN_VERSION < 7.6.3");
 #endif
 
   return ::tsl::OkStatus();
@@ -2784,11 +2783,11 @@ tsl::StatusOr<cudnnConvolutionFwdAlgo_t> GetCudnnConvolutionForwardAlgo(
         perf_results[r].algo != CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED &&
         perf_results[r].memory <= mem_limit) {
       return perf_results[r].algo;
     }
   }
-  return tsl::Status(port::error::INTERNAL,
+  return tsl::Status(tsl::error::INTERNAL,
                      "cudnnGetConvolutionForwardAlgorithm_v7 returned "
                      "no suitable algorithms. This could be a cudnn bug.");
 #else
   cudnnConvolutionFwdPreference_t preference =
       specify_workspace_limit ? CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT
@@ -2826,11 +2825,11 @@ GetCudnnConvolutionBackwardDataAlgo(const CudnnHandle& cudnn,
             CUDNN_CONVOLUTION_BWD_DATA_ALGO_WINOGRAD_NONFUSED &&
         perf_results[r].memory <= mem_limit) {
       return perf_results[r].algo;
     }
   }
-  return tsl::Status(port::error::INTERNAL,
+  return tsl::Status(tsl::error::INTERNAL,
                      "cudnnGetConvolutionBackwardDataAlgorithm_v7 returned "
                      "no suitable algorithms. This could be a cudnn bug.");
 #else
   cudnnConvolutionBwdDataPreference_t preference =
       specify_workspace_limit
@@ -2868,11 +2867,11 @@ GetCudnnConvolutionBackwardFilterAlgo(const CudnnHandle& cudnn,
             CUDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD_NONFUSED &&
         perf_results[r].memory <= mem_limit) {
       return perf_results[r].algo;
     }
   }
-  return tsl::Status(port::error::INTERNAL,
+  return tsl::Status(tsl::error::INTERNAL,
                      "cudnnGetConvolutionBackwardFilterAlgorithm_v7 returned "
                      "no suitable algorithms. This could be a cudnn bug.");
 #else
   cudnnConvolutionBwdFilterPreference_t preference =
       specify_workspace_limit
@@ -2893,11 +2892,11 @@ tsl::StatusOr<DeviceMemory<uint8_t>> AllocateCudnnConvolutionForwardWorkspace(
     const CudnnTensorDescriptor& output_nd,
     const dnn::AlgorithmDesc& algorithm_desc,
     ScratchAllocator* scratch_allocator) {
   if (IsTensorMathOpSet(conv) != algorithm_desc.tensor_ops_enabled()) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         "Mismatch between cudnn conv and algorithm descriptors.");
   }
 
   // Query the size of the workspace and allocate it.
   size_t size_in_bytes;
@@ -2915,21 +2914,21 @@ tsl::StatusOr<DeviceMemory<uint8_t>> AllocateCudnnConvolutionForwardWorkspace(
 
   int64_t size_in_bytes_int64_t = size_in_bytes;
 
   if (ABSL_PREDICT_FALSE(size_in_bytes_int64_t < 0)) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         "cudnnGetConvolutionForwardWorkspaceSize() returned "
         "negative sizeInBytes value. This could be a cudnn bug.");
   }
 
   if (size_in_bytes_int64_t == 0) {
     return DeviceMemory<uint8_t>();
   }
 
   if (ABSL_PREDICT_FALSE(!scratch_allocator)) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "No scratch allocator provided");
   }
 
   return scratch_allocator->AllocateBytes(size_in_bytes);
 }
@@ -2942,11 +2941,11 @@ AllocateCudnnConvolutionBackwardDataWorkspace(
     const CudnnTensorDescriptor& output_nd,
     const dnn::AlgorithmDesc& algorithm_desc,
     ScratchAllocator* scratch_allocator) {
   if (IsTensorMathOpSet(conv) != algorithm_desc.tensor_ops_enabled()) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         "Mismatch between cudnn conv and algorithm descriptors.");
   }
 
   // Query the size of the workspace and allocate it.
   size_t size_in_bytes;
@@ -2965,21 +2964,21 @@ AllocateCudnnConvolutionBackwardDataWorkspace(
 
   int64_t size_in_bytes_int64_t = size_in_bytes;
 
   if (ABSL_PREDICT_FALSE(size_in_bytes_int64_t < 0)) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         "cudnnGetConvolutionBackwardDataWorkspaceSize() returned "
         "negative sizeInBytes value. This could be a cudnn bug.");
   }
 
   if (size_in_bytes_int64_t == 0) {
     return DeviceMemory<uint8_t>();
   }
 
   if (ABSL_PREDICT_FALSE(!scratch_allocator)) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "No scratch allocator provided");
   }
 
   return scratch_allocator->AllocateBytes(size_in_bytes);
 }
@@ -2992,11 +2991,11 @@ AllocateCudnnConvolutionBackwardFilterWorkspace(
     const CudnnTensorDescriptor& output_nd,
     const dnn::AlgorithmDesc& algorithm_desc,
     ScratchAllocator* scratch_allocator) {
   if (IsTensorMathOpSet(conv) != algorithm_desc.tensor_ops_enabled()) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         "Mismatch between cudnn conv and algorithm descriptors.");
   }
 
   // Query the size of the workspace and allocate it.
   size_t size_in_bytes;
@@ -3015,21 +3014,21 @@ AllocateCudnnConvolutionBackwardFilterWorkspace(
 
   int64_t size_in_bytes_int64_t = size_in_bytes;
 
   if (ABSL_PREDICT_FALSE(size_in_bytes_int64_t < 0)) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         "cudnnGetConvolutionBackwardFilterWorkspaceSize() returned "
         "negative sizeInBytes value. This could be a cudnn bug.");
   }
 
   if (size_in_bytes_int64_t == 0) {
     return DeviceMemory<uint8_t>();
   }
 
   if (ABSL_PREDICT_FALSE(!scratch_allocator)) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "No scratch allocator provided");
   }
 
   return scratch_allocator->AllocateBytes(size_in_bytes);
 }
@@ -3038,11 +3037,11 @@ tsl::StatusOr<bool> UseTensorOps(Stream* stream, dnn::DataType type,
                                  std::optional<dnn::AlgorithmDesc> desc) {
   bool use_tensor_ops;
   if (desc.has_value()) {
     use_tensor_ops = desc->tensor_ops_enabled();
     if (use_tensor_ops && !IsTensorMathEnabled(stream, type)) {
-      return tsl::Status(port::error::INVALID_ARGUMENT,
+      return tsl::Status(tsl::error::INVALID_ARGUMENT,
                          "Algo requests disabled tensor op evaluation.");
     }
   } else {
     use_tensor_ops = IsTensorMathEnabled(stream, type);
   }
@@ -3160,11 +3159,11 @@ tsl::StatusOr<dnn::AlgorithmDesc> GetCudnnConvolutionBackwardDataAlgorithm(
 
   // Failed to allocate workspace for the first algorithm, fall back to the
   // no_scratch algorithm.
   if (!algo_desc.has_value()) {
     return tsl::Status(
-        port::error::INVALID_ARGUMENT,
+        tsl::error::INVALID_ARGUMENT,
         "The primary convolution algorithm failed memory allocation, "
         "while a secondary algorithm is not provided.");
   }
 
   TF_ASSIGN_OR_RETURN(use_tensor_ops,
@@ -3222,11 +3221,11 @@ tsl::StatusOr<dnn::AlgorithmDesc> GetCudnnConvolutionBackwardFilterAlgorithm(
 
   // Failed to allocate workspace for the first algorithm, fall back to the
   // no_scratch algorithm.
   if (!algo_desc.has_value()) {
     return tsl::Status(
-        port::error::INVALID_ARGUMENT,
+        tsl::error::INVALID_ARGUMENT,
         absl::StrCat(
             "The primary convolution algorithm failed memory allocation, "
             "while a secondary algorithm is not provided. Actual error: ",
             scratch_or.status().ToString()));
   }
@@ -4252,21 +4251,21 @@ class CudnnLegacyConvRunner : public dnn::ConvRunner {
       timer.reset(new GpuTimer(parent_));  // NOLINT
       // The start and stop of the timer should be as close to the Cudnn call as
       // possible. It is still possible for other threads to issue workload on
       // to this stream. So it could take multiple profiling measurements.
       if (!timer->Init() || !timer->Start(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
       }
     }
 
     const auto get_fwd_bugs = [&]() -> tsl::Status {
 #if CUDNN_VERSION < 8000
       if (algo_id_ == CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM &&
           ToCudnnDataType(input_type_) == CUDNN_DATA_INT8 &&
           ToCudnnDataType(output_type_) == CUDNN_DATA_FLOAT) {
         return tsl::Status(
-            port::error::FAILED_PRECONDITION,
+            tsl::error::FAILED_PRECONDITION,
             "This configuration potentially produces incorrect results.");
       }
 #else
       (void)output_type_;  // To stop clang-tidy saying it's unused.
 #endif
@@ -4334,11 +4333,11 @@ class CudnnLegacyConvRunner : public dnn::ConvRunner {
                                      static_cast<int>(kind_));
     }
 
     if (is_profiling) {
       if (!timer->Stop(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
       }
       profile_result->set_algorithm(algo);
       profile_result->set_elapsed_time_in_ms(timer->GetElapsedMilliseconds());
       profile_result->set_scratch_size(scratch_memory.size());
     }
@@ -4629,21 +4628,21 @@ class CudnnExecutionPlanRunner<void(Args...)>
       timer.reset(new GpuTimer(parent_));  // NOLINT
       // The start and stop of the timer should be as close to the Cudnn call as
       // possible. It is still possible for other threads to issue workload on
       // to this stream. So it could take multiple profiling measurements.
       if (!timer->Init() || !timer->Start(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
       }
     }
 
     cudnnStatus_t status = cudnnBackendExecute(
         cudnn.handle(), plan_.get_raw_desc(), variantPack.get_raw_desc());
     RETURN_IF_CUDNN_ERROR(status);
 
     if (is_profiling) {
       if (!timer->Stop(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
       }
       TF_ASSIGN_OR_RETURN(auto desc, ToAlgorithmDesc());
       profile_result->set_algorithm(desc);
       profile_result->set_elapsed_time_in_ms(timer->GetElapsedMilliseconds());
       profile_result->set_scratch_size(scratch_memory.size());
@@ -4866,11 +4865,11 @@ tsl::Status CudnnSupport::GetConvolveRunners(
                                                       input_type, &algorithms);
         break;
     }
     if (!got_algos) {
       return tsl::Status(
-          port::error::UNKNOWN,
+          tsl::error::UNKNOWN,
           absl::StrFormat("Listing algorithms failed for kind %d", kind));
     }
 
     for (const auto& algo : algorithms) {
       auto runner_or = ConvolveRunnerFromDesc(
@@ -5035,11 +5034,11 @@ class CudnnLegacyFusedConvRunner : public dnn::FusedConvRunner {
       timer.reset(new GpuTimer(parent_));  // NOLINT
       // The start and stop of the timer should be as close to the Cudnn call as
       // possible. It is still possible for other threads to issue workload on
       // to this stream. So it could take multiple profiling measurements.
       if (!timer->Init() || !timer->Start(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
       }
     }
     auto side_input_data_ptr = (side_input_scale_ == 0)
                                    ? output_data.opaque()
                                    : side_input_data.opaque();
@@ -5063,11 +5062,11 @@ class CudnnLegacyFusedConvRunner : public dnn::FusedConvRunner {
             << "\nactivation_desc.handle() = " << activation_desc_.handle()
             << "\noutput_nd.handle() = " << output_nd_.handle()
             << "\noutput_data.opaque() = " << output_data.opaque();
 
     if (IsTensorMathOpSet(conv_) != tensor_ops_enabled_) {
-      return tsl::Status(port::error::FAILED_PRECONDITION,
+      return tsl::Status(tsl::error::FAILED_PRECONDITION,
                          "Tensor op math type in dnn::AlgorithmDesc does not "
                          "match that of the CudnnConvolutionDescriptor");
     }
 
     // N.B. the scaling parameters alpha1 and alpha2 are pointers to
@@ -5093,11 +5092,11 @@ class CudnnLegacyFusedConvRunner : public dnn::FusedConvRunner {
     }
     RETURN_IF_CUDNN_ERROR(status);
 
     if (profile_result) {
       if (!timer->Stop(AsGpuStream(stream))) {
-        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
+        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
       }
       profile_result->set_algorithm(algo);
       profile_result->set_elapsed_time_in_ms(timer->GetElapsedMilliseconds());
       profile_result->set_scratch_size(scratch_memory.size());
       VLOG(4) << "conv with algorithm " << ToConvForwardAlgo(algo)
@@ -5306,22 +5305,22 @@ tsl::Status CudnnSupport::GetFusedConvolveRunners(
   if (activation_mode != dnn::ActivationMode::kRelu &&
       activation_mode != dnn::ActivationMode::kRelu6 &&
       activation_mode != dnn::ActivationMode::kElu &&
       activation_mode != dnn::ActivationMode::kLeakyRelu &&
       activation_mode != dnn::ActivationMode::kNone) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "CuDNN fusion only supports activations of "
                        "{Relu, Relu6, Elu, <None>}.");
   }
 
   if (!actually_use_cudnn_frontend) {
     std::vector<dnn::AlgorithmDesc> algorithms;
 
     auto cuda_compute_capability = stream->GetCudaComputeCapability();
     if (!GetConvolveAlgorithms(cuda_compute_capability, input_type,
                                &algorithms)) {
-      return tsl::Status(port::error::UNKNOWN,
+      return tsl::Status(tsl::error::UNKNOWN,
                          "Listing fused convolve algorithms failed.");
     }
 
     for (const auto& algo : algorithms) {
       // Only CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM is supported
@@ -5352,11 +5351,11 @@ tsl::Status CudnnSupport::GetFusedConvolveRunners(
   auto op_graph_status = GetCudnnFusedOperationGraph(
       kind, input_type, bias_type, output_type, conv_scale, side_input_scale,
       leakyrelu_alpha, input_descriptor, filter_descriptor, bias_descriptor,
       output_descriptor, convolution_descriptor, activation_mode, cudnn);
   if (!op_graph_status.status().ok()) {
-    return tsl::Status(port::error::INTERNAL,
+    return tsl::Status(tsl::error::INTERNAL,
                        absl::StrCat("Cudnn graph failed to build: ",
                                     op_graph_status.status().ToString()));
   }
   auto op_graph = std::move(op_graph_status).value();
 
@@ -5389,11 +5388,11 @@ tsl::Status CudnnSupport::GetFusedMatmulRunners(
   auto cudnn = cudnn_->GetHandle(parent_, stream);
   auto op_graph_status = GetCudnnFusedMatmulGraph(
       input_type, bias_type, output_type, trans_a, trans_b, m, n, k, lda, ldb,
       ldc, activation_mode, cudnn);
   if (!op_graph_status.status().ok()) {
-    return tsl::Status(port::error::INTERNAL,
+    return tsl::Status(tsl::error::INTERNAL,
                        absl::StrCat("Cudnn graph failed to build: ",
                                     op_graph_status.status().ToString()));
   }
   auto op_graph = std::move(op_graph_status).value();
 
@@ -5683,11 +5682,11 @@ tsl::Status CudnnSupport::DoBatchNormalizationForwardImpl(
 
   auto check_no_side_input_or_activation = [&]() -> tsl::Status {
     if (activation_mode != dnn::ActivationMode::kNone ||
         !side_input.is_null()) {
       return tsl::Status(
-          port::error::INTERNAL,
+          tsl::error::INTERNAL,
           absl::StrCat(
               "Side input and activation are not supported by cuDNN version: ",
               CUDNN_VERSION));
     } else {
       return ::tsl::OkStatus();
@@ -5966,11 +5965,11 @@ tsl::Status CudnnSupport::DoFusedConvolve(
         "go/nvbugs/3326122");
   }
 
   if (activation_mode != dnn::ActivationMode::kRelu &&
       activation_mode != dnn::ActivationMode::kNone) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "cudnnConvolutionBiasActivationForward() only supports "
                        "Relu or None activation.");
   }
 
   CudnnTensorDescriptor conv_input_nd(
@@ -6068,11 +6067,11 @@ tsl::Status CudnnSupport::DoPrepareForCtcLoss(
         /*ctcLossDesc=*/cudnn_ctc_loss_desc.handle(),
         /*sizeInBytes=*/&workspace_size_in_bytes));
   }
   *ctc_loss_algo_id = algo;
 #else
-  return tsl::Status(port::error::INVALID_ARGUMENT,
+  return tsl::Status(tsl::error::INVALID_ARGUMENT,
                      "No supported cudnnGetCTCLossWorkspaceSize when "
                      "CUDNN_VERSION < 7.6.3");
 #endif
   // Allocate the workspace.
   if (workspace_size_in_bytes == 0) {
@@ -6098,11 +6097,11 @@ tsl::Status CudnnSupport::DoCtcLoss(
     const dnn::RnnStateTensorDescriptor& grads_desc,
     DeviceMemoryBase grads_data, DeviceMemory<uint8_t> scratch_memory,
     int ctc_loss_algo_id) {
   // Current cuDNN CTC Loss only supports the float datatype
   if (CUDNN_VERSION < 7603 || element_type != dnn::DataType::kFloat) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "CudnnCtcLossDescriptor is supported only when the "
                        "CUDNN_VERSION >= 7.6.3 and DataType is float");
   }
   CudnnCtcLossDescriptor cudnn_ctc_loss_desc(ToCudnnDataType(element_type));
   const CudnnRnnStateTensorDescriptor& cudnn_probs_desc =
@@ -6380,11 +6379,11 @@ tsl::StatusOr<std::vector<PoolingSplitsSpec>> GetTensorSplits(
   int64_t max_batches_per_split =
       std::numeric_limits<int>::max() / elements_per_batch_input;
 
   if (max_batches_per_split == 0) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrCat(
             "Tensor has too many elements for int32 indexing: batches=",
             num_batches, " elements_per_batch=", elements_per_batch_input,
             "."));
   }
@@ -6440,11 +6439,11 @@ tsl::Status CudnnSupport::DoPoolForward(
   };
 
   auto splits_or =
       GetTensorSplits(input_dimensions, output_dimensions, element_type);
   if (!splits_or.ok()) {
-    return tsl::Status(port::error::INTERNAL, "Cudnn pooling failed to split");
+    return tsl::Status(tsl::error::INTERNAL, "Cudnn pooling failed to split");
   }
   auto splits = std::move(splits_or.value());
 
   dnn::BatchDescriptor input_split = input_dimensions;
   dnn::BatchDescriptor output_split = output_dimensions;
@@ -6509,11 +6508,11 @@ tsl::Status CudnnSupport::DoPoolBackward(
   };
 
   auto splits_or =
       GetTensorSplits(input_dimensions, output_dimensions, element_type);
   if (!splits_or.ok()) {
-    return tsl::Status(port::error::INTERNAL, "Cudnn pooling failed to split");
+    return tsl::Status(tsl::error::INTERNAL, "Cudnn pooling failed to split");
   }
   auto splits = std::move(splits_or.value());
 
   dnn::BatchDescriptor input_split = input_dimensions;
   dnn::BatchDescriptor output_split = output_dimensions;
diff --git a/tensorflow/compiler/xla/stream_executor/cuda/cuda_driver.cc b/tensorflow/compiler/xla/stream_executor/cuda/cuda_driver.cc
index a61a6b1a06718..9c0f6c22e8722 100644
--- a/tensorflow/compiler/xla/stream_executor/cuda/cuda_driver.cc
+++ b/tensorflow/compiler/xla/stream_executor/cuda/cuda_driver.cc
@@ -33,14 +33,14 @@ limitations under the License.
 #include "absl/strings/str_format.h"
 #include "absl/synchronization/mutex.h"
 #include "absl/synchronization/notification.h"
 #include "third_party/gpus/cuda/include/cuda_runtime_api.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_diagnostics.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/tsl/platform/env.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/stacktrace.h"
 #include "tensorflow/tsl/platform/static_threadlocal.h"
 #include "tensorflow/tsl/platform/threadpool.h"
 
 bool FLAGS_gpuexec_cuda_driver_inject_init_error = false;
@@ -265,11 +265,11 @@ static tsl::Status InternalInit() {
   } else {
     LOG(ERROR) << "failed call to cuInit: " << ToString(res);
   }
 
   Diagnostician::LogDiagnosticInformation();
-  return tsl::Status(port::error::ABORTED,
+  return tsl::Status(tsl::error::ABORTED,
                      absl::StrCat("failed call to cuInit: ", ToString(res)));
 }
 
 }  // namespace
 
@@ -398,11 +398,11 @@ bool DeviceOptionsToContextFlags(const DeviceOptions& device_options,
     } else {
       absl::StrAppend(&message, "; could not query total memory");
     }
   }
 
-  return tsl::Status(port::error::INTERNAL, message);
+  return tsl::Status(tsl::error::INTERNAL, message);
 }
 
 /* static */ void GpuDriver::DestroyContext(GpuContext* context) {
   if (context == nullptr) {
     return;
@@ -671,11 +671,11 @@ bool DeviceOptionsToContextFlags(const DeviceOptions& device_options,
   if (result == CUDA_SUCCESS) {
     return device;
   }
 
   return tsl::Status(
-      port::error::INTERNAL,
+      tsl::error::INTERNAL,
       absl::StrCat("failed to get device for context: ", ToString(result)));
 }
 
 /* static */ bool GpuDriver::CreateStream(GpuContext* context, CUstream* stream,
                                           int priority) {
@@ -970,11 +970,11 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
 #endif
 
 /* static */ tsl::Status GpuDriver::DestroyEvent(GpuContext* context,
                                                  CUevent* event) {
   if (*event == nullptr) {
-    return tsl::Status(port::error::INVALID_ARGUMENT,
+    return tsl::Status(tsl::error::INVALID_ARGUMENT,
                        "input event cannot be null");
   }
 
   ScopedActivateContext activated{context};
   RETURN_IF_CUDA_RES_ERROR(cuEventDestroy(*event),
@@ -995,11 +995,11 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
                                                            CUevent event) {
   ScopedActivateContext activated{context};
   CUresult res = cuEventQuery(event);
   if (res != CUDA_SUCCESS && res != CUDA_ERROR_NOT_READY) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat("failed to query event: %s", ToString(res)));
   }
 
   return res;
 }
@@ -1261,15 +1261,15 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
   CUresult res = cuEventCreate(result, cuflags);
 
   if (res == CUDA_SUCCESS) {
     return ::tsl::OkStatus();
   } else if (res == CUDA_ERROR_OUT_OF_MEMORY) {
-    return tsl::Status(port::error::RESOURCE_EXHAUSTED,
+    return tsl::Status(tsl::error::RESOURCE_EXHAUSTED,
                        "could not create CUDA event: out of device memory");
   } else {
     return tsl::Status(
-        port::error::FAILED_PRECONDITION,
+        tsl::error::FAILED_PRECONDITION,
         absl::StrCat("could not create CUDA event: ", ToString(res)));
   }
 }
 
 /* static */ int GpuDriver::GetDeviceCount() {
@@ -1297,18 +1297,18 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
     // everywhere in TF?  Currently this is only used during error
     // handling.  So all is working fine, but TF have a different
     // error then the original one.
     if (context == nullptr) {
       return tsl::Status(
-          port::error::UNAVAILABLE,
+          tsl::error::UNAVAILABLE,
           "Empty context returned while querying context for device pointer");
     }
     return context;
   }
 
   return tsl::Status(
-      port::error::INTERNAL,
+      tsl::error::INTERNAL,
       absl::StrCat("failed to query context for device pointer: ",
                    ToString(result)));
 }
 
 /* static */ tsl::StatusOr<MemorySpace> GpuDriver::GetPointerMemorySpace(
@@ -1322,17 +1322,17 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
         return MemorySpace::kDevice;
       case CU_MEMORYTYPE_HOST:
         return MemorySpace::kHost;
       default:
         return tsl::Status(
-            port::error::INTERNAL,
+            tsl::error::INTERNAL,
             absl::StrCat("unknown memory space provided by CUDA API: ", value));
     }
   }
 
   return tsl::Status(
-      port::error::INTERNAL,
+      tsl::error::INTERNAL,
       absl::StrCat("failed to query device pointer for memory space: ",
                    ToString(result)));
 }
 
 /* static */ tsl::Status GpuDriver::GetPointerAddressRange(CUdeviceptr dptr,
@@ -1344,17 +1344,17 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
   } else if (result == CUDA_ERROR_NOT_FOUND) {
     // We differentiate between "this pointer is unknown" (return here) and
     // "there was an internal error while performing this operation" (return
     // below).
     return tsl::Status(
-        port::error::NOT_FOUND,
+        tsl::error::NOT_FOUND,
         absl::StrFormat("not a device pointer %p; %s",
                         reinterpret_cast<void*>(dptr), ToString(result)));
   }
 
   return tsl::Status(
-      port::error::INTERNAL,
+      tsl::error::INTERNAL,
       absl::StrFormat("failed to get pointer into for device pointer %p; %s",
                       reinterpret_cast<void*>(dptr), ToString(result)));
 }
 
 /* static */ tsl::StatusOr<CUdevice> GpuDriver::GetPointerDevice(
@@ -1375,21 +1375,21 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
 
   CUresult res = cuDeviceGetAttribute(
       cc_major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device);
   if (res != CUDA_SUCCESS) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat(
             "failed to get compute capability major for device: %s; %d",
             ToString(res), device));
   }
 
   res = cuDeviceGetAttribute(
       cc_minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device);
   if (res != CUDA_SUCCESS) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat(
             "failed to get compute capability minor for device: %s; %d",
             ToString(res), device));
   }
 
@@ -1397,17 +1397,17 @@ GpuDriver::CreateMemoryHandle(GpuContext* context, uint64_t bytes) {
 }
 
 /* static */ tsl::Status GpuDriver::GetGpuISAVersion(int* version,
                                                      CUdevice device) {
   return tsl::Status{
-      port::error::INTERNAL,
+      tsl::error::INTERNAL,
       "Feature not supported on CUDA platform (GetGpuISAVersion)"};
 }
 
 /* static */ tsl::Status GpuDriver::GetGpuGCNArchName(CUdevice, std::string*) {
   return tsl::Status{
-      port::error::INTERNAL,
+      tsl::error::INTERNAL,
       "Feature not supported on CUDA platform (GetGpuGCNArchName)"};
 }
 
 // Helper function that turns the integer output of cuDeviceGetAttribute to type
 // T and wraps it in a StatusOr.
@@ -1517,11 +1517,11 @@ static tsl::StatusOr<T> GetSimpleAttribute(CUdevice device,
     CUdevice_attribute attribute, CUdevice device) {
   int val;
   CUresult res = cuDeviceGetAttribute(&val, attribute, device);
   if (res != CUDA_SUCCESS) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat("failed to get device attribute %d for device %d: %s",
                         attribute, device, ToString(res)));
   }
   return val;
 }
@@ -1626,11 +1626,11 @@ static tsl::StatusOr<T> GetSimpleAttribute(CUdevice device,
   ScopedActivateContext activated{from};
   CUresult result = cuCtxEnablePeerAccess(to->context(), 0 /* = flags */);
   if (result != CUDA_SUCCESS &&
       result != CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat("failed to enable peer access from %p to %p: %s", from,
                         to, ToString(result)));
   }
 
   return ::tsl::OkStatus();
diff --git a/tensorflow/compiler/xla/stream_executor/cuda/cuda_fft.cc b/tensorflow/compiler/xla/stream_executor/cuda/cuda_fft.cc
index 72988ab2f5901..d3914e1becca8 100644
--- a/tensorflow/compiler/xla/stream_executor/cuda/cuda_fft.cc
+++ b/tensorflow/compiler/xla/stream_executor/cuda/cuda_fft.cc
@@ -24,11 +24,11 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_helpers.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/device_memory.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/stream.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_internal.h"
diff --git a/tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.cc b/tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.cc
index 77ad104fa74c0..3b01066dc7a1d 100644
--- a/tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.cc
+++ b/tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.cc
@@ -40,21 +40,21 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_event.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_timer.h"
 #include "tensorflow/compiler/xla/stream_executor/kernel_cache_config.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/stream.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_internal.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.h"
 #include "tensorflow/compiler/xla/stream_executor/timer.h"
 #include "tensorflow/tsl/platform/env.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/numbers.h"
 #include "tensorflow/tsl/platform/statusor.h"
 
 // LOG(ERROR) uses a const named ERROR, so a macro with the same name is
 // always unwanted. This happens on Windows that defines such a macro.
@@ -743,11 +743,11 @@ tsl::Status GpuExecutor::WaitForEvent(Stream* stream, Event* event) {
   if (GpuDriver::WaitStreamOnEvent(context_, AsGpuStream(stream)->gpu_stream(),
                                    AsGpuEvent(event)->gpu_event())) {
     return ::tsl::OkStatus();
   } else {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat("error recording waiting for CUDA event on stream %p",
                         stream));
   }
 }
 
diff --git a/tensorflow/compiler/xla/stream_executor/cuda/cuda_platform.cc b/tensorflow/compiler/xla/stream_executor/cuda/cuda_platform.cc
index bd33db289bdca..9be642587685f 100644
--- a/tensorflow/compiler/xla/stream_executor/cuda/cuda_platform.cc
+++ b/tensorflow/compiler/xla/stream_executor/cuda/cuda_platform.cc
@@ -21,12 +21,12 @@ limitations under the License.
 #include "absl/strings/str_cat.h"
 #include "absl/strings/str_format.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_driver.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_platform_id.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/status.h"
 
 namespace stream_executor {
 namespace gpu {
 namespace {
@@ -115,11 +115,11 @@ tsl::StatusOr<StreamExecutor*> CudaPlatform::FirstExecutorForBus(
       return *ExecutorForDevice(i);
     }
   }
 
   return tsl::Status(
-      port::error::NOT_FOUND,
+      tsl::error::NOT_FOUND,
       absl::StrFormat("Executor for bus %d not found.", bus_ordinal));
 }
 
 Platform::Id CudaPlatform::id() const { return cuda::kCudaPlatformId; }
 
@@ -175,11 +175,11 @@ CudaPlatform::GetUncachedExecutor(const StreamExecutorConfig& config) {
       this, std::make_unique<GpuExecutor>(config.plugin_config),
       config.ordinal);
   auto init_status = executor->Init(config.device_options);
   if (!init_status.ok()) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat(
             "failed initializing StreamExecutor for CUDA device ordinal %d: %s",
             config.ordinal, init_status.ToString()));
   }
 
diff --git a/tensorflow/compiler/xla/stream_executor/cuda/cuda_rng.cc b/tensorflow/compiler/xla/stream_executor/cuda/cuda_rng.cc
index 80115f09d3c34..d80ef0b5b4017 100644
--- a/tensorflow/compiler/xla/stream_executor/cuda/cuda_rng.cc
+++ b/tensorflow/compiler/xla/stream_executor/cuda/cuda_rng.cc
@@ -21,11 +21,11 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_helpers.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/cuda/cuda_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/device_memory.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/rng.h"
 #include "tensorflow/tsl/platform/status.h"
 // clang-format off
 #include "third_party/gpus/cuda/include/curand.h"
diff --git a/tensorflow/compiler/xla/stream_executor/host/BUILD b/tensorflow/compiler/xla/stream_executor/host/BUILD
index 8522091da8395..7034425adfdd7 100644
--- a/tensorflow/compiler/xla/stream_executor/host/BUILD
+++ b/tensorflow/compiler/xla/stream_executor/host/BUILD
@@ -44,11 +44,11 @@ cc_library(
         "//tensorflow/compiler/xla/stream_executor:executor_cache",
         "//tensorflow/compiler/xla/stream_executor:multi_platform_manager",
         "//tensorflow/compiler/xla/stream_executor:stream_executor_headers",
         "//tensorflow/compiler/xla/stream_executor/lib",
         "//tensorflow/compiler/xla/stream_executor/platform",
-        "@com_google_absl//absl/base:core_headers",
+        "//tensorflow/tsl/platform:errors",
         "@com_google_absl//absl/memory",
         "@com_google_absl//absl/strings:str_format",
     ],
     alwayslink = True,  # Registers itself with the MultiPlatformManager.
 )
@@ -102,13 +102,12 @@ cc_library(
         ":host_timer",
         "//tensorflow/compiler/xla/stream_executor",
         "//tensorflow/compiler/xla/stream_executor:kernel",
         "//tensorflow/compiler/xla/stream_executor:rng",
         "//tensorflow/compiler/xla/stream_executor:stream_executor_internal",
-        "//tensorflow/compiler/xla/stream_executor:stream_executor_pimpl",
-        "//tensorflow/compiler/xla/stream_executor:timer",
-        "//tensorflow/compiler/xla/stream_executor/lib",
+        "//tensorflow/compiler/xla/stream_executor:stream_executor_pimpl",  # fixdeps: keep
+        "//tensorflow/tsl/platform:errors",
         "//tensorflow/tsl/platform:platform_port",
         "//tensorflow/tsl/platform/profile_utils:profile_utils_cpu_utils",
         "@com_google_absl//absl/functional:any_invocable",
         "@com_google_absl//absl/strings",
         "@com_google_absl//absl/synchronization",
diff --git a/tensorflow/compiler/xla/stream_executor/host/host_gpu_executor.h b/tensorflow/compiler/xla/stream_executor/host/host_gpu_executor.h
index e5a595976ddc5..316fcdc438654 100644
--- a/tensorflow/compiler/xla/stream_executor/host/host_gpu_executor.h
+++ b/tensorflow/compiler/xla/stream_executor/host/host_gpu_executor.h
@@ -23,14 +23,14 @@ limitations under the License.
 
 #include "absl/functional/any_invocable.h"
 #include "tensorflow/compiler/xla/stream_executor/blas.h"
 #include "tensorflow/compiler/xla/stream_executor/host/host_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/host/host_timer.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/rng.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_internal.h"
+#include "tensorflow/tsl/platform/errors.h"
 
 namespace stream_executor {
 namespace host {
 
 // An implementation of StreamExecutor that does no communication or interaction
diff --git a/tensorflow/compiler/xla/stream_executor/host/host_platform.cc b/tensorflow/compiler/xla/stream_executor/host/host_platform.cc
index ac59bed788450..e254ce9709362 100644
--- a/tensorflow/compiler/xla/stream_executor/host/host_platform.cc
+++ b/tensorflow/compiler/xla/stream_executor/host/host_platform.cc
@@ -19,12 +19,12 @@ limitations under the License.
 
 #include "absl/memory/memory.h"
 #include "absl/strings/str_format.h"
 #include "tensorflow/compiler/xla/stream_executor/host/host_gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/host/host_platform_id.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
+#include "tensorflow/tsl/platform/errors.h"
 
 namespace stream_executor {
 namespace host {
 
 HostPlatform::HostPlatform() : name_("Host") {}
@@ -73,11 +73,11 @@ HostPlatform::GetUncachedExecutor(const StreamExecutorConfig& config) {
       this, std::make_unique<HostExecutor>(config.plugin_config),
       config.ordinal);
   auto init_status = executor->Init(config.device_options);
   if (!init_status.ok()) {
     return tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat(
             "failed initializing StreamExecutor for device ordinal %d: %s",
             config.ordinal, init_status.ToString().c_str()));
   }
 
diff --git a/tensorflow/compiler/xla/stream_executor/lib/BUILD b/tensorflow/compiler/xla/stream_executor/lib/BUILD
index fb6972030a115..a450e6338bfa5 100644
--- a/tensorflow/compiler/xla/stream_executor/lib/BUILD
+++ b/tensorflow/compiler/xla/stream_executor/lib/BUILD
@@ -29,11 +29,10 @@ cc_library(
     deps = [
         "@com_google_absl//absl/base:core_headers",
         "@com_google_absl//absl/strings",
         "@com_google_absl//absl/strings:str_format",
         "@com_google_absl//absl/types:span",
-        "//tensorflow/compiler/xla/stream_executor/platform",
         "//tensorflow/tsl/platform:env",
         "//tensorflow/tsl/platform:stacktrace",
         "//tensorflow/tsl/platform:status",
         "//tensorflow/tsl/platform:statusor",
         "//tensorflow/tsl/protobuf:error_codes_proto_impl_cc",
diff --git a/tensorflow/compiler/xla/stream_executor/lib/error.h b/tensorflow/compiler/xla/stream_executor/lib/error.h
deleted file mode 100644
index 1153a4c75c41f..0000000000000
--- a/tensorflow/compiler/xla/stream_executor/lib/error.h
+++ /dev/null
@@ -1,30 +0,0 @@
-/* Copyright 2015 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.
-==============================================================================*/
-
-
-#ifndef TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_ERROR_H_
-#define TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_ERROR_H_
-
-#include "tensorflow/tsl/protobuf/error_codes.pb.h"  // IWYU pragma: export
-
-namespace stream_executor {
-namespace port {
-
-namespace error = tensorflow::error;
-
-}  // namespace port
-}  // namespace stream_executor
-
-#endif  // TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_ERROR_H_
diff --git a/tensorflow/compiler/xla/stream_executor/lib/initialize.h b/tensorflow/compiler/xla/stream_executor/lib/initialize.h
deleted file mode 100644
index 93bec22a8c5e6..0000000000000
--- a/tensorflow/compiler/xla/stream_executor/lib/initialize.h
+++ /dev/null
@@ -1,21 +0,0 @@
-/* Copyright 2015 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.
-==============================================================================*/
-
-#ifndef TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_INITIALIZE_H_
-#define TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_INITIALIZE_H_
-
-#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
-
-#endif  // TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_INITIALIZE_H_
diff --git a/tensorflow/compiler/xla/stream_executor/multi_platform_manager.cc b/tensorflow/compiler/xla/stream_executor/multi_platform_manager.cc
index c60a541a02903..f4322e8ef9486 100644
--- a/tensorflow/compiler/xla/stream_executor/multi_platform_manager.cc
+++ b/tensorflow/compiler/xla/stream_executor/multi_platform_manager.cc
@@ -22,12 +22,11 @@ limitations under the License.
 #include "absl/strings/ascii.h"
 #include "absl/strings/str_format.h"
 #include "absl/strings/str_join.h"
 #include "absl/strings/string_view.h"
 #include "absl/synchronization/mutex.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/tsl/platform/errors.h"
 
 namespace stream_executor {
 namespace {
 
@@ -94,11 +93,11 @@ tsl::Status MultiPlatformManagerImpl::RegisterPlatform(
     std::unique_ptr<Platform> platform) {
   CHECK(platform != nullptr);
   std::string key = absl::AsciiStrToLower(platform->Name());
   absl::MutexLock lock(&mu_);
   if (name_map_.find(key) != name_map_.end()) {
-    return tsl::Status(port::error::INTERNAL,
+    return tsl::Status(tsl::error::INTERNAL,
                        "platform is already registered with name: \"" +
                            platform->Name() + "\"");
   }
   Platform* platform_ptr = platform.get();
   CHECK(id_map_.emplace(platform->id(), platform_ptr).second);
@@ -154,11 +153,11 @@ tsl::StatusOr<Platform*> MultiPlatformManagerImpl::InitializePlatformWithName(
   absl::MutexLock lock(&mu_);
 
   TF_ASSIGN_OR_RETURN(Platform * platform, LookupByNameLocked(target));
   if (platform->Initialized()) {
     return tsl::Status(
-        port::error::FAILED_PRECONDITION,
+        tsl::error::FAILED_PRECONDITION,
         absl::StrCat("platform \"", target, "\" is already initialized"));
   }
 
   TF_RETURN_IF_ERROR(platform->Initialize(options));
 
@@ -170,11 +169,11 @@ tsl::StatusOr<Platform*> MultiPlatformManagerImpl::InitializePlatformWithId(
   absl::MutexLock lock(&mu_);
 
   TF_ASSIGN_OR_RETURN(Platform * platform, LookupByIdLocked(id));
   if (platform->Initialized()) {
     return tsl::Status(
-        port::error::FAILED_PRECONDITION,
+        tsl::error::FAILED_PRECONDITION,
         absl::StrFormat("platform with id %p is already initialized", id));
   }
 
   TF_RETURN_IF_ERROR(platform->Initialize(options));
 
@@ -230,11 +229,11 @@ MultiPlatformManagerImpl::InitializedPlatformNamesWithFilter(
 tsl::StatusOr<Platform*> MultiPlatformManagerImpl::LookupByNameLocked(
     absl::string_view target) {
   auto it = name_map_.find(absl::AsciiStrToLower(target));
   if (it == name_map_.end()) {
     return tsl::Status(
-        port::error::NOT_FOUND,
+        tsl::error::NOT_FOUND,
         absl::StrCat("Could not find registered platform with name: \"", target,
                      "\". Available platform names are: ",
                      absl::StrJoin(InitializedPlatformNamesWithFilter(), " ")));
   }
   return it->second;
@@ -243,11 +242,11 @@ tsl::StatusOr<Platform*> MultiPlatformManagerImpl::LookupByNameLocked(
 tsl::StatusOr<Platform*> MultiPlatformManagerImpl::LookupByIdLocked(
     const Platform::Id& id) {
   auto it = id_map_.find(id);
   if (it == id_map_.end()) {
     return tsl::Status(
-        port::error::NOT_FOUND,
+        tsl::error::NOT_FOUND,
         absl::StrFormat("could not find registered platform with id: %p", id));
   }
   return it->second;
 }
 
diff --git a/tensorflow/compiler/xla/stream_executor/multi_platform_manager.h b/tensorflow/compiler/xla/stream_executor/multi_platform_manager.h
index ea005d973570e..c4ebf8e60991b 100644
--- a/tensorflow/compiler/xla/stream_executor/multi_platform_manager.h
+++ b/tensorflow/compiler/xla/stream_executor/multi_platform_manager.h
@@ -68,12 +68,12 @@ limitations under the License.
 #include <map>
 #include <memory>
 #include <vector>
 
 #include "absl/strings/string_view.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/tsl/platform/status.h"
 #include "tensorflow/tsl/platform/statusor.h"
 
 namespace stream_executor {
diff --git a/tensorflow/compiler/xla/stream_executor/platform.cc b/tensorflow/compiler/xla/stream_executor/platform.cc
index 9e5c22c27627e..aaca9354968ae 100644
--- a/tensorflow/compiler/xla/stream_executor/platform.cc
+++ b/tensorflow/compiler/xla/stream_executor/platform.cc
@@ -14,14 +14,14 @@ limitations under the License.
 ==============================================================================*/
 
 #include "tensorflow/compiler/xla/stream_executor/platform.h"
 
 #include "absl/strings/str_cat.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.h"
+#include "tensorflow/tsl/platform/errors.h"
 
 namespace stream_executor {
 
 std::string PlatformKindString(PlatformKind kind) {
   switch (kind) {
@@ -90,18 +90,18 @@ Platform::~Platform() {}
 bool Platform::Initialized() const { return true; }
 
 tsl::Status Platform::Initialize(
     const std::map<std::string, std::string> &platform_options) {
   if (!platform_options.empty()) {
-    return tsl::Status(port::error::UNIMPLEMENTED,
+    return tsl::Status(tsl::error::UNIMPLEMENTED,
                        "this platform does not support custom initialization");
   }
   return ::tsl::OkStatus();
 }
 
 tsl::Status Platform::ForceExecutorShutdown() {
-  return tsl::Status(port::error::UNIMPLEMENTED,
+  return tsl::Status(tsl::error::UNIMPLEMENTED,
                      "executor shutdown is not supported on this platform");
 }
 
 std::unique_ptr<Platform::PeerAccessMap> Platform::GetPeerAccessMap() {
   auto *map = new PeerAccessMap;
diff --git a/tensorflow/compiler/xla/stream_executor/plugin_registry.cc b/tensorflow/compiler/xla/stream_executor/plugin_registry.cc
index dc682ae5ded93..9b55a25aca2d2 100644
--- a/tensorflow/compiler/xla/stream_executor/plugin_registry.cc
+++ b/tensorflow/compiler/xla/stream_executor/plugin_registry.cc
@@ -17,12 +17,12 @@ limitations under the License.
 
 #include "absl/base/const_init.h"
 #include "absl/strings/str_cat.h"
 #include "absl/strings/str_format.h"
 #include "absl/synchronization/mutex.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/multi_platform_manager.h"
+#include "tensorflow/tsl/platform/errors.h"
 
 namespace stream_executor {
 
 const PluginId kNullPlugin = nullptr;
 
@@ -74,11 +74,11 @@ tsl::Status PluginRegistry::RegisterFactoryInternal(
     std::map<PluginId, FACTORY_TYPE>* factories) {
   absl::MutexLock lock{&GetPluginRegistryMutex()};
 
   if (factories->find(plugin_id) != factories->end()) {
     return tsl::Status(
-        port::error::ALREADY_EXISTS,
+        tsl::error::ALREADY_EXISTS,
         absl::StrFormat("Attempting to register factory for plugin %s when "
                         "one has already been registered",
                         plugin_name));
   }
 
@@ -94,11 +94,11 @@ tsl::StatusOr<FACTORY_TYPE> PluginRegistry::GetFactoryInternal(
   auto iter = factories.find(plugin_id);
   if (iter == factories.end()) {
     iter = generic_factories.find(plugin_id);
     if (iter == generic_factories.end()) {
       return tsl::Status(
-          port::error::NOT_FOUND,
+          tsl::error::NOT_FOUND,
           absl::StrFormat("Plugin ID %p not registered.", plugin_id));
     }
   }
 
   return iter->second;
@@ -215,11 +215,11 @@ bool PluginRegistry::HasFactory(Platform::Id platform_id,
     if (plugin_id == PluginConfig::kDefault) {                                \
       plugin_id = default_factories_[platform_id].FACTORY_VAR;                \
                                                                               \
       if (plugin_id == kNullPlugin) {                                         \
         return tsl::Status(                                                   \
-            port::error::FAILED_PRECONDITION,                                 \
+            tsl::error::FAILED_PRECONDITION,                                  \
             "No suitable " PLUGIN_STRING                                      \
             " plugin registered. Have you linked in a " PLUGIN_STRING         \
             "-providing plugin?");                                            \
       } else {                                                                \
         VLOG(2) << "Selecting default " PLUGIN_STRING " plugin, "             \
@@ -234,11 +234,11 @@ bool PluginRegistry::HasFactory(Platform::Id platform_id,
   template <>                                                                 \
   tsl::StatusOr<PluginRegistry::FACTORY_TYPE> PluginRegistry::GetFactory(     \
       PlatformKind platform_kind, PluginId plugin_id) {                       \
     auto iter = platform_id_by_kind_.find(platform_kind);                     \
     if (iter == platform_id_by_kind_.end()) {                                 \
-      return tsl::Status(port::error::FAILED_PRECONDITION,                    \
+      return tsl::Status(tsl::error::FAILED_PRECONDITION,                     \
                          absl::StrFormat("Platform kind %d not registered.",  \
                                          static_cast<int>(platform_kind)));   \
     }                                                                         \
     return GetFactory<PluginRegistry::FACTORY_TYPE>(iter->second, plugin_id); \
   }
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_blas.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_blas.cc
index ee441e5eedc17..1cf271a8380fd 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_blas.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_blas.cc
@@ -30,12 +30,12 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_activation.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_helpers.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_timer.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/dso_loader.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/scratch_allocator.h"
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_diagnostics.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_diagnostics.cc
index a0f43e02d94ba..e0ccb9be27512 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_diagnostics.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_diagnostics.cc
@@ -33,12 +33,12 @@ limitations under the License.
 #include "absl/strings/numbers.h"
 #include "absl/strings/str_cat.h"
 #include "absl/strings/str_format.h"
 #include "absl/strings/str_split.h"
 #include "absl/strings/strip.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/host_info.h"
 
 namespace stream_executor {
 namespace rocm {
 
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_dnn.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_dnn.cc
index 0d98e247ffef7..62e4c5bc83cfe 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_dnn.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_dnn.cc
@@ -29,21 +29,21 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_activation.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_driver.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_timer.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/dso_loader.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_diagnostics.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/scratch_allocator.h"
 #include "tensorflow/compiler/xla/stream_executor/stream.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.h"
 #include "tensorflow/tsl/platform/env.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/hash.h"
 #include "tensorflow/tsl/util/determinism.h"
 #include "tensorflow/tsl/util/env_var.h"
 
 namespace {
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_driver.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_driver.cc
index 859b767a779a8..4d72ff6590846 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_driver.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_driver.cc
@@ -26,15 +26,15 @@ limitations under the License.
 #include "absl/strings/str_format.h"
 #include "absl/synchronization/mutex.h"
 #include "absl/synchronization/notification.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_diagnostics.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_driver.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_driver_wrapper.h"
 #include "tensorflow/tsl/platform/env.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/numbers.h"
 #include "tensorflow/tsl/platform/stacktrace.h"
 #include "tensorflow/tsl/platform/static_threadlocal.h"
 #include "tensorflow/tsl/platform/threadpool.h"
 
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_fft.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_fft.cc
index e3a64b8bc2cdd..633fced55e158 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_fft.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_fft.cc
@@ -20,12 +20,12 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/device_memory.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_activation.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_helpers.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_stream.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/dso_loader.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_internal.h"
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_gpu_executor.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_gpu_executor.cc
index e28be74f5689d..94dafa1b40943 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_gpu_executor.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_gpu_executor.cc
@@ -26,24 +26,24 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_event.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_stream.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_timer.h"
 #include "tensorflow/compiler/xla/stream_executor/kernel_cache_config.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/dso_loader.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/plugin_registry.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_diagnostics.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_platform_id.h"
 #include "tensorflow/compiler/xla/stream_executor/stream.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_internal.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.h"
 #include "tensorflow/compiler/xla/stream_executor/timer.h"
 #include "tensorflow/tsl/platform/env.h"
+#include "tensorflow/tsl/platform/errors.h"
 
 #ifdef PLATFORMS_GPUS_ROCM_DYNAMIC_LIBROCM_DYNAMIC_LIBROCM_H_
 #error \
     "No driver calls in this file, wrap driver functionality in rocm_driver.cc."
 #endif
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_platform.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_platform.cc
index e7e2566da83c0..5218d5f2f2eee 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_platform.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_platform.cc
@@ -19,13 +19,13 @@ limitations under the License.
 
 #include "absl/base/call_once.h"
 #include "absl/strings/str_format.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_driver.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_executor.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_platform_id.h"
+#include "tensorflow/tsl/platform/errors.h"
 
 namespace stream_executor {
 namespace gpu {
 
 ROCmPlatform::ROCmPlatform()
diff --git a/tensorflow/compiler/xla/stream_executor/rocm/rocm_rng.cc b/tensorflow/compiler/xla/stream_executor/rocm/rocm_rng.cc
index 0a4fbe0a1b4fa..4bd74fe003117 100644
--- a/tensorflow/compiler/xla/stream_executor/rocm/rocm_rng.cc
+++ b/tensorflow/compiler/xla/stream_executor/rocm/rocm_rng.cc
@@ -18,12 +18,12 @@ limitations under the License.
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_activation.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_executor.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_helpers.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_rng.h"
 #include "tensorflow/compiler/xla/stream_executor/gpu/gpu_stream.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/dso_loader.h"
+#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/logging.h"
 #include "tensorflow/compiler/xla/stream_executor/rng.h"
 #include "tensorflow/compiler/xla/stream_executor/rocm/rocm_platform_id.h"
 #include "tensorflow/tsl/platform/env.h"
 
diff --git a/tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.cc b/tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.cc
index c6babbe4919fb..8dc30fd966eb8 100644
--- a/tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.cc
+++ b/tensorflow/compiler/xla/stream_executor/stream_executor_pimpl.cc
@@ -30,15 +30,15 @@ limitations under the License.
 #include "absl/strings/str_cat.h"
 #include "absl/strings/str_format.h"
 #include "absl/synchronization/notification.h"
 #include "tensorflow/compiler/xla/stream_executor/blas.h"
 #include "tensorflow/compiler/xla/stream_executor/fft.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/platform/port.h"
 #include "tensorflow/compiler/xla/stream_executor/rng.h"
 #include "tensorflow/compiler/xla/stream_executor/stream.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor_internal.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/stacktrace.h"
 #include "tensorflow/tsl/platform/statusor.h"
 #include "tensorflow/tsl/platform/threadpool.h"
 #include "tensorflow/tsl/util/env_var.h"
 
@@ -403,11 +403,11 @@ StreamExecutor::createRnnDescriptor(
     dnn::DataType data_type, const dnn::AlgorithmConfig& algorithm_config,
     float dropout, uint64_t seed, ScratchAllocator* state_allocator,
     bool use_padded_io) {
   dnn::DnnSupport* dnn_support = AsDnn();
   if (!dnn_support) {
-    return tsl::Status(port::error::UNKNOWN,
+    return tsl::Status(tsl::error::UNKNOWN,
                        "Fail to find the dnn implementation.");
   }
   return dnn_support->createRnnDescriptor(
       num_layers, hidden_size, input_size, cell_size, batch_size, input_mode,
       direction_mode, rnn_mode, data_type, algorithm_config, dropout, seed,
@@ -418,11 +418,11 @@ tsl::StatusOr<std::unique_ptr<dnn::RnnSequenceTensorDescriptor>>
 StreamExecutor::createRnnSequenceTensorDescriptor(int max_seq_length,
                                                   int batch_size, int data_size,
                                                   dnn::DataType data_type) {
   dnn::DnnSupport* dnn_support = AsDnn();
   if (!dnn_support) {
-    return tsl::Status(port::error::UNKNOWN,
+    return tsl::Status(tsl::error::UNKNOWN,
                        "Fail to find the dnn implementation.");
   }
   return dnn_support->createRnnSequenceTensorDescriptor(
       max_seq_length, batch_size, data_size, data_type);
 }
@@ -432,11 +432,11 @@ StreamExecutor::createRnnSequenceTensorDescriptor(
     int max_seq_length, int batch_size, int data_size,
     const absl::Span<const int>& seq_lengths, bool time_major,
     dnn::DataType data_type) {
   dnn::DnnSupport* dnn_support = AsDnn();
   if (!dnn_support) {
-    return tsl::Status(port::error::UNKNOWN,
+    return tsl::Status(tsl::error::UNKNOWN,
                        "Fail to find the dnn implementation.");
   }
   return dnn_support->createRnnSequenceTensorDescriptor(
       max_seq_length, batch_size, data_size, seq_lengths, time_major,
       data_type);
@@ -446,11 +446,11 @@ tsl::StatusOr<std::unique_ptr<dnn::RnnStateTensorDescriptor>>
 StreamExecutor::createRnnStateTensorDescriptor(int num_layer, int batch_size,
                                                int data_size,
                                                dnn::DataType data_type) {
   dnn::DnnSupport* dnn_support = AsDnn();
   if (!dnn_support) {
-    return tsl::Status(port::error::UNKNOWN,
+    return tsl::Status(tsl::error::UNKNOWN,
                        "Fail to find the dnn implementation.");
   }
   return dnn_support->createRnnStateTensorDescriptor(num_layer, batch_size,
                                                      data_size, data_type);
 }
@@ -544,11 +544,11 @@ tsl::StatusOr<DeviceMemoryBase> StreamExecutor::GetUntypedSymbol(
   if (GetSymbol(symbol_name, module_handle, &opaque, &bytes)) {
     return DeviceMemoryBase(opaque, bytes);
   }
 
   return tsl::Status(
-      port::error::NOT_FOUND,
+      tsl::error::NOT_FOUND,
       absl::StrCat("Check if module containing symbol ", symbol_name,
                    " is loaded (module_handle = ",
                    reinterpret_cast<uintptr_t>(module_handle.id()), ")"));
 }
 
@@ -689,11 +689,11 @@ tsl::Status StreamExecutor::SynchronousMemcpyD2H(
                host_dst);
 
   result = implementation_->SynchronousMemcpy(host_dst, device_src, size);
   if (!result.ok()) {
     result = tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat("failed to synchronously memcpy device-to-host: device "
                         "%p to host %p size %d: %s",
                         device_src.opaque(), host_dst, size,
                         result.ToString()));
   }
@@ -713,11 +713,11 @@ tsl::Status StreamExecutor::SynchronousMemcpyH2D(const void* host_src,
                device_dst);
 
   result = implementation_->SynchronousMemcpy(device_dst, host_src, size);
   if (!result.ok()) {
     result = tsl::Status(
-        port::error::INTERNAL,
+        tsl::error::INTERNAL,
         absl::StrFormat("failed to synchronously memcpy host-to-device: host "
                         "%p to device %p size %d: %s",
                         host_src, device_dst->opaque(), size,
                         result.ToString()));
   }
diff --git a/tensorflow/compiler/xla/stream_executor/tf_allocator_adapter.cc b/tensorflow/compiler/xla/stream_executor/tf_allocator_adapter.cc
index 5dbf5d68dbd11..2dd10d165e1b3 100644
--- a/tensorflow/compiler/xla/stream_executor/tf_allocator_adapter.cc
+++ b/tensorflow/compiler/xla/stream_executor/tf_allocator_adapter.cc
@@ -14,11 +14,10 @@ limitations under the License.
 ==============================================================================*/
 
 #include "tensorflow/compiler/xla/stream_executor/tf_allocator_adapter.h"
 
 #include "absl/synchronization/mutex.h"
-#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
 #include "tensorflow/compiler/xla/stream_executor/stream.h"
 #include "tensorflow/compiler/xla/stream_executor/stream_executor.h"
 #include "tensorflow/tsl/platform/errors.h"
 
 namespace stream_executor {
diff --git a/tensorflow/core/common_runtime/function.cc b/tensorflow/core/common_runtime/function.cc
index 8c6fe48b7494d..df0e098b6ae48 100644
--- a/tensorflow/core/common_runtime/function.cc
+++ b/tensorflow/core/common_runtime/function.cc
@@ -533,11 +533,11 @@ class CallOp : public AsyncOpKernel {
   void ComputeAsync(OpKernelContext* ctx, DoneCallback done) override {
     FunctionLibraryRuntime* lib = ctx->function_library();
     OP_REQUIRES_ASYNC(ctx, lib != nullptr,
                       errors::Internal("No function library is provided."),
                       done);
-    FunctionLibraryRuntime::Options opts;
+    FunctionLibraryRuntime::Options opts(ctx->step_id());
     opts.rendezvous = ctx->rendezvous();
     opts.cancellation_manager = ctx->cancellation_manager();
     opts.step_container = ctx->step_container();
     opts.stats_collector = ctx->stats_collector();
     opts.runner = ctx->runner();
diff --git a/tensorflow/core/data/BUILD b/tensorflow/core/data/BUILD
index bd6823a894258..9028fe8356612 100644
--- a/tensorflow/core/data/BUILD
+++ b/tensorflow/core/data/BUILD
@@ -255,10 +255,11 @@ cc_library(
     srcs = ["tfdataz_metrics.cc"],
     hdrs = ["tfdataz_metrics.h"],
     # copybara:uncomment copts = ["-Wthread-safety-analysis"],
     visibility = ["//visibility:public"],
     deps = [
+        "//tensorflow/core:framework",
         "//tensorflow/core:lib",
         "//tensorflow/core/platform:env",
         "//tensorflow/core/platform:mutex",
         "//tensorflow/core/platform:thread_annotations",
         "@com_google_absl//absl/container:flat_hash_set",
diff --git a/tensorflow/core/data/service/server_lib.cc b/tensorflow/core/data/service/server_lib.cc
index 12d9636ecd11e..24f00fa4ccb33 100644
--- a/tensorflow/core/data/service/server_lib.cc
+++ b/tensorflow/core/data/service/server_lib.cc
@@ -35,11 +35,11 @@ namespace data {
 namespace {
 constexpr char kPortPlaceholder[] = "%port%";
 }
 
 GrpcDataServerBase::GrpcDataServerBase(
-    int port, const std::string& protocol, const std::string server_type,
+    int port, const std::string& protocol, const std::string& server_type,
     std::vector<std::unique_ptr<::grpc::ServerBuilderOption>> options)
     : requested_port_(port),
       protocol_(protocol),
       server_type_(server_type),
       bound_port_(port),
diff --git a/tensorflow/core/data/service/server_lib.h b/tensorflow/core/data/service/server_lib.h
index 279f888d6c3ae..d06b387ba58ad 100644
--- a/tensorflow/core/data/service/server_lib.h
+++ b/tensorflow/core/data/service/server_lib.h
@@ -42,11 +42,11 @@ class GrpcDataServerBase {
   // Constructs a tf.data server with the specified port. If the port is 0, the
   // server will find an available port in `Start()`. The chosen port can be
   // found by calling `BoundPort()`.
   GrpcDataServerBase(
       int requested_port, const std::string& protocol,
-      const std::string server_type,
+      const std::string& server_type,
       std::vector<std::unique_ptr<::grpc::ServerBuilderOption>> options = {});
   virtual ~GrpcDataServerBase() = default;
 
   // Starts the server running asynchronously.
   Status Start();
diff --git a/tensorflow/core/data/service/snapshot/BUILD b/tensorflow/core/data/service/snapshot/BUILD
index 8a053910dd1e8..22a8e3d03ceeb 100644
--- a/tensorflow/core/data/service/snapshot/BUILD
+++ b/tensorflow/core/data/service/snapshot/BUILD
@@ -82,11 +82,13 @@ cc_library(
     name = "path_utils",
     srcs = ["path_utils.cc"],
     hdrs = ["path_utils.h"],
     compatible_with = get_compatible_with_portable(),
     deps = [
+        "//tensorflow/tsl/platform:errors",
         "//tensorflow/tsl/platform:path",
+        "//tensorflow/tsl/platform:statusor",
         "@com_google_absl//absl/strings",
     ],
 )
 
 tf_cc_test(
@@ -94,10 +96,12 @@ tf_cc_test(
     srcs = ["path_utils_test.cc"],
     deps = [
         ":path_utils",
         "//tensorflow/core:test",
         "//tensorflow/core:test_main",
+        "//tensorflow/tsl/platform:status_matchers",
+        "//tensorflow/tsl/protobuf:protos_all_cc",
     ],
 )
 
 cc_library(
     name = "snapshot_split_provider",
diff --git a/tensorflow/core/data/service/snapshot/path_utils.cc b/tensorflow/core/data/service/snapshot/path_utils.cc
index 050c904de0cde..04b1cedf4fe2f 100644
--- a/tensorflow/core/data/service/snapshot/path_utils.cc
+++ b/tensorflow/core/data/service/snapshot/path_utils.cc
@@ -13,14 +13,19 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 #include "tensorflow/core/data/service/snapshot/path_utils.h"
 
 #include <string>
+#include <utility>
+#include <vector>
 
 #include "absl/strings/str_cat.h"
+#include "absl/strings/str_split.h"
 #include "absl/strings/string_view.h"
+#include "tensorflow/tsl/platform/errors.h"
 #include "tensorflow/tsl/platform/path.h"
+#include "tensorflow/tsl/platform/statusor.h"
 
 namespace tensorflow {
 namespace data {
 namespace {
 
@@ -63,10 +68,32 @@ std::string SplitPath(absl::string_view snapshot_path, int64_t stream_index,
   return tsl::io::JoinPath(
       SourceDirectory(snapshot_path, stream_index, source_id),
       absl::StrCat("split_", local_index, "_", global_index));
 }
 
+tsl::StatusOr<std::pair<int64_t, int64_t>> SplitIndex(
+    absl::string_view split_path) {
+  std::vector<std::string> tokens = absl::StrSplit(split_path, '_');
+  int64_t local_split_index = 0, global_split_index = 0;
+  if (tokens.size() != 3 || tokens[0] != "split" ||
+      !absl::SimpleAtoi(tokens[1], &local_split_index) ||
+      local_split_index < 0 ||
+      !absl::SimpleAtoi(tokens[2], &global_split_index) ||
+      global_split_index < 0) {
+    return tsl::errors::InvalidArgument(
+        "Invalid split file name: ", split_path,
+        ". Expected split_<local_split_index>_<global_split_index>.");
+  }
+  if (local_split_index > global_split_index) {
+    return tsl::errors::InvalidArgument(
+        "Invalid split file name: ", split_path, ". The local split index ",
+        local_split_index, " exceeds the global split index ",
+        global_split_index, ".");
+  }
+  return std::make_pair(local_split_index, global_split_index);
+}
+
 std::string SnapshotMetadataFilePath(absl::string_view snapshot_path_) {
   return tsl::io::JoinPath(snapshot_path_, kSnapshotMetadataFileName);
 }
 
 std::string DatasetDefFilePath(absl::string_view snapshot_path_) {
diff --git a/tensorflow/core/data/service/snapshot/path_utils.h b/tensorflow/core/data/service/snapshot/path_utils.h
index 58d5a565cbb9a..387d4d12093ea 100644
--- a/tensorflow/core/data/service/snapshot/path_utils.h
+++ b/tensorflow/core/data/service/snapshot/path_utils.h
@@ -14,12 +14,14 @@ limitations under the License.
 ==============================================================================*/
 #ifndef TENSORFLOW_CORE_DATA_SERVICE_SNAPSHOT_PATH_UTILS_H_
 #define TENSORFLOW_CORE_DATA_SERVICE_SNAPSHOT_PATH_UTILS_H_
 
 #include <string>
+#include <utility>
 
 #include "absl/strings/string_view.h"
+#include "tensorflow/tsl/platform/statusor.h"
 
 namespace tensorflow {
 namespace data {
 
 // Returns the directory path for the assigned streams of a snapshot.
@@ -43,10 +45,16 @@ std::string SourceDirectory(absl::string_view snapshot_path,
 // of a snapshot.
 std::string SplitPath(absl::string_view snapshot_path, int64_t stream_index,
                       int64_t source_id, int64_t local_index,
                       int64_t global_index);
 
+// Returns a pair of {local_split_index, global_split_index} of the split. The
+// expected format of `split_path` is:
+// split_<local_split_index>_<global_split_index>
+tsl::StatusOr<std::pair<int64_t, int64_t>> SplitIndex(
+    absl::string_view split_path);
+
 // Returns the path of the DONE file of a snapshot stream.
 std::string StreamDoneFilePath(absl::string_view snapshot_path,
                                int64_t stream_index);
 
 // Returns the path of the serialized metadata for a snapshot.
diff --git a/tensorflow/core/data/service/snapshot/path_utils_test.cc b/tensorflow/core/data/service/snapshot/path_utils_test.cc
index e84c359d510e1..7bd27bbd85eef 100644
--- a/tensorflow/core/data/service/snapshot/path_utils_test.cc
+++ b/tensorflow/core/data/service/snapshot/path_utils_test.cc
@@ -12,17 +12,23 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 #include "tensorflow/core/data/service/snapshot/path_utils.h"
 
+#include "tensorflow/tsl/platform/status_matchers.h"
 #include "tensorflow/tsl/platform/test.h"
+#include "tensorflow/tsl/protobuf/error_codes.pb.h"
 
 namespace tensorflow {
 namespace data {
 namespace {
 
+using ::testing::HasSubstr;
 using ::testing::MatchesRegex;
+using ::testing::Pair;
+using tsl::testing::IsOkAndHolds;
+using tsl::testing::StatusIs;
 
 TEST(PathUtilsTest, StreamsDirectory) {
   EXPECT_THAT(StreamsDirectory("/path/to/snapshot"),
               MatchesRegex("/path/to/snapshot.streams"));
 }
@@ -49,10 +55,38 @@ TEST(PathUtilsTest, SplitPath) {
                 /*local_index=*/2, /*global_index=*/3),
       MatchesRegex(
           "/path/to/snapshot.streams.stream_0.splits.source_1.split_2_3"));
 }
 
+TEST(PathUtilsTest, SplitIndex) {
+  EXPECT_THAT(SplitIndex("split_0_1"), IsOkAndHolds(Pair(0, 1)));
+}
+
+TEST(PathUtilsTest, InvalidSplitFile) {
+  EXPECT_THAT(
+      SplitIndex(""),
+      StatusIs(error::INVALID_ARGUMENT,
+               HasSubstr(
+                   "Expected split_<local_split_index>_<global_split_index>")));
+  EXPECT_THAT(
+      SplitIndex("split_123"),
+      StatusIs(error::INVALID_ARGUMENT,
+               HasSubstr(
+                   "Expected split_<local_split_index>_<global_split_index>")));
+  EXPECT_THAT(
+      SplitIndex("split_-1_(-1)"),
+      StatusIs(error::INVALID_ARGUMENT,
+               HasSubstr(
+                   "Expected split_<local_split_index>_<global_split_index>")));
+  EXPECT_THAT(
+      SplitIndex("split_5_0"),
+      StatusIs(
+          error::INVALID_ARGUMENT,
+          HasSubstr(
+              "The local split index 5 exceeds the global split index 0")));
+}
+
 TEST(PathUtilsTest, StreamDoneFilePath) {
   EXPECT_THAT(StreamDoneFilePath("/path/to/snapshot", /*stream_index=*/0),
               MatchesRegex("/path/to/snapshot.streams.stream_0.DONE"));
 }
 
diff --git a/tensorflow/core/data/service/snapshot/snapshot_manager.cc b/tensorflow/core/data/service/snapshot/snapshot_manager.cc
index 9fe47df0ee6b7..2f24c14648580 100644
--- a/tensorflow/core/data/service/snapshot/snapshot_manager.cc
+++ b/tensorflow/core/data/service/snapshot/snapshot_manager.cc
@@ -185,25 +185,12 @@ Status SnapshotManager::ReadOnDiskSource(
   for (const auto& split_filename : split_filenames) {
     std::string split_path = io::JoinPath(source_path, split_filename);
 
     // `split_filename` must have this format:
     // "split_<local_split_index>_<global_split_index>".
-    std::vector<std::string> tokens = absl::StrSplit(split_filename, '_');
-    int64_t local_split_index;
-    int64_t global_split_index;
-    if (tokens.size() != 3 ||
-        !absl::SimpleAtoi(tokens[1], &local_split_index) ||
-        local_split_index < 0 ||
-        !absl::SimpleAtoi(tokens[2], &global_split_index) ||
-        global_split_index < 0) {
-      return InvalidArgument("can't parse the name of ", split_path);
-    }
-    if (local_split_index > global_split_index) {
-      return InvalidArgument(
-          "found conflict between local split index and global split index in ",
-          "name of ", split_path);
-    }
+    TF_ASSIGN_OR_RETURN(auto split_index, SplitIndex(split_filename));
+    auto [local_split_index, global_split_index] = split_index;
     if (local_split_index > split_filenames.size() - 1) {
       return InvalidArgument(
           "found conflict between the number of splits and name of ",
           split_path);
     }
diff --git a/tensorflow/core/data/tfdataz_metrics.cc b/tensorflow/core/data/tfdataz_metrics.cc
index 496a8a49b6d90..5516f5f9a4f43 100644
--- a/tensorflow/core/data/tfdataz_metrics.cc
+++ b/tensorflow/core/data/tfdataz_metrics.cc
@@ -86,12 +86,13 @@ absl::Duration ApproximateLatencyEstimator::GetAverageLatency(Duration duration)
       static_cast<double>(latency_count_counter_ -
                           latency_count_[PrevSlot(static_cast<int>(duration))]);
   return absl::Duration(absl::Microseconds(interval_latency)) / interval_count;
 }
 
-TfDatazMetricsCollector::TfDatazMetricsCollector(const Env& env)
-    : latency_estimator_(env) {}
+TfDatazMetricsCollector::TfDatazMetricsCollector(const Env& env,
+                                                 IteratorBase* iterator)
+    : iterator_(iterator), latency_estimator_(env) {}
 
 void TfDatazMetricsCollector::RecordGetNextLatency(
     int64_t get_next_latency_usec) {
   if (get_next_latency_usec > 0) {
     latency_estimator_.AddLatency(get_next_latency_usec);
@@ -111,10 +112,14 @@ absl::Duration TfDatazMetricsCollector::GetAverageLatencyForLastFiveMinutes() {
 absl::Duration TfDatazMetricsCollector::GetAverageLatencyForLastSixtyMinutes() {
   return latency_estimator_.GetAverageLatency(
       ApproximateLatencyEstimator::Duration::kSixtyMinutes);
 }
 
+int64_t TfDatazMetricsCollector::GetIteratorTotalMemoryUsage() {
+  return iterator_->TotalBufferedBytes();
+}
+
 namespace {
 static mutex* get_tfdataz_metrics_registry_lock() {
   static mutex tfdataz_metrics_registry_lock(LINKER_INITIALIZED);
   return &tfdataz_metrics_registry_lock;
 }
diff --git a/tensorflow/core/data/tfdataz_metrics.h b/tensorflow/core/data/tfdataz_metrics.h
index 8a44e6ebaa5fe..bb59a73f8d04f 100644
--- a/tensorflow/core/data/tfdataz_metrics.h
+++ b/tensorflow/core/data/tfdataz_metrics.h
@@ -23,10 +23,11 @@ limitations under the License.
 #include <utility>
 #include <vector>
 
 #include "absl/container/flat_hash_set.h"
 #include "absl/time/time.h"
+#include "tensorflow/core/framework/dataset.h"
 #include "tensorflow/core/platform/env.h"
 #include "tensorflow/core/platform/mutex.h"
 #include "tensorflow/core/platform/thread_annotations.h"
 
 namespace tensorflow {
@@ -93,11 +94,11 @@ class TfDatazMetricsCollector {
  public:
   // Constructs a `TfDatazMetricsCollector`.
   // We only collect metrics for CPU devices. This is a heuristic to avoid
   // collecting metrics for device-side iterators created by the multi-device
   // iterator mechanism.
-  explicit TfDatazMetricsCollector(const Env& env);
+  TfDatazMetricsCollector(const Env& env, IteratorBase* iterator);
 
   // Records `GetNext` call latency.
   void RecordGetNextLatency(int64_t get_next_latency_usec);
 
   // Returns the average `GetNext` latency for past 1 minute.
@@ -107,11 +108,17 @@ class TfDatazMetricsCollector {
   absl::Duration GetAverageLatencyForLastFiveMinutes();
 
   // Returns the average `GetNext` latency for past 60 minutes.
   absl::Duration GetAverageLatencyForLastSixtyMinutes();
 
+  // Returns the total memory (in bytes) used by the iterator.
+  // Total memory used by the iterator includes the total number of bytes
+  // buffered in all nodes in the subtree.
+  int64_t GetIteratorTotalMemoryUsage();
+
  private:
+  IteratorBase* iterator_;  // not owned
   ApproximateLatencyEstimator latency_estimator_;
 };
 
 // Thread-safe global registry for the /tfdataz metrics. All callers to
 // `TfDatazMetricsRegistry` use the same instance to register and deregister
diff --git a/tensorflow/core/data/tfdataz_metrics_test.cc b/tensorflow/core/data/tfdataz_metrics_test.cc
index 02f02658965dd..2d3d75946e4d9 100644
--- a/tensorflow/core/data/tfdataz_metrics_test.cc
+++ b/tensorflow/core/data/tfdataz_metrics_test.cc
@@ -16,11 +16,11 @@ limitations under the License.
 
 #include <memory>
 #include <utility>
 
 #include "absl/time/time.h"
-#include "tensorflow/core/framework/types.h"
+#include "tensorflow/core/framework/dataset.h"
 #include "tensorflow/core/platform/env.h"
 #include "tensorflow/core/platform/test.h"
 #include "tensorflow/core/util/fake_clock_env.h"
 
 namespace tensorflow {
@@ -39,18 +39,20 @@ static int64_t k61MinutesInMicros =
 
 class TfDatazMetricsTest : public ::testing::Test {
  protected:
   void SetUp() override {
     env_ = std::make_unique<FakeClockEnv>(Env::Default());
-    tfdataz_metrics_ = std::make_unique<TfDatazMetricsCollector>(*env_);
+    tfdataz_metrics_ =
+        std::make_unique<TfDatazMetricsCollector>(*env_, iterator_.get());
   }
 
   void TearDown() override {
     env_.reset();
     tfdataz_metrics_.reset();
   }
 
+  std::unique_ptr<IteratorBase> iterator_;
   std::unique_ptr<FakeClockEnv> env_;
   std::unique_ptr<TfDatazMetricsCollector> tfdataz_metrics_;
 };
 
 TEST_F(TfDatazMetricsTest, RecordGetNextLatency) {
@@ -182,28 +184,30 @@ class ScopedTfDataMetricsRegistration {
  private:
   std::shared_ptr<TfDatazMetricsCollector> collector_;
 };
 
 TEST(TfDatazMetricsRegistryTest, Register) {
-  auto collector_one =
-      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
-  auto collector_two =
-      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
+  std::unique_ptr<IteratorBase> iterator;
+  auto collector_one = std::make_shared<TfDatazMetricsCollector>(
+      *Env::Default(), iterator.get());
+  auto collector_two = std::make_shared<TfDatazMetricsCollector>(
+      *Env::Default(), iterator.get());
 
   ScopedTfDataMetricsRegistration scoped_registration_one(collector_one);
   ScopedTfDataMetricsRegistration scoped_registration_two(collector_two);
 
   EXPECT_EQ(TfDatazMetricsRegistry::GetIteratorMetricCollectors().size(), 2);
 }
 
 TEST(TfDatazMetricsRegistryTest, Deregister) {
-  auto collector_one =
-      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
-  auto collector_two =
-      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
-  auto collector_three =
-      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
+  std::unique_ptr<IteratorBase> iterator;
+  auto collector_one = std::make_shared<TfDatazMetricsCollector>(
+      *Env::Default(), iterator.get());
+  auto collector_two = std::make_shared<TfDatazMetricsCollector>(
+      *Env::Default(), iterator.get());
+  auto collector_three = std::make_shared<TfDatazMetricsCollector>(
+      *Env::Default(), iterator.get());
   ScopedTfDataMetricsRegistration scoped_registration_one(collector_one);
   ScopedTfDataMetricsRegistration scoped_registration_two(collector_two);
   ScopedTfDataMetricsRegistration scoped_registration_three(collector_three);
   EXPECT_EQ(TfDatazMetricsRegistry::GetIteratorMetricCollectors().size(), 3);
 
diff --git a/tensorflow/core/distributed_runtime/eager/remote_mgr.cc b/tensorflow/core/distributed_runtime/eager/remote_mgr.cc
index 9e0977e0e14aa..070edf7f5a936 100644
--- a/tensorflow/core/distributed_runtime/eager/remote_mgr.cc
+++ b/tensorflow/core/distributed_runtime/eager/remote_mgr.cc
@@ -31,11 +31,11 @@ namespace {
 Status WithErrorSourcePayload(Status error) {
   core::platform::ErrorSourceProto error_source_proto;
   error_source_proto.set_error_source(
       core::platform::ErrorSourceProto::EAGER_REMOTE_MGR);
   error.SetPayload(tensorflow::kErrorSource,
-                   error_source_proto.SerializeAsString());
+                   absl::Cord(error_source_proto.SerializeAsString()));
   return error;
 }
 }  // namespace
 
 namespace eager {
diff --git a/tensorflow/core/distributed_runtime/integration_test/coordination_test_opkernel_registration.cc b/tensorflow/core/distributed_runtime/integration_test/coordination_test_opkernel_registration.cc
index 33b4f7fe6bfae..9b04a419af657 100644
--- a/tensorflow/core/distributed_runtime/integration_test/coordination_test_opkernel_registration.cc
+++ b/tensorflow/core/distributed_runtime/integration_test/coordination_test_opkernel_registration.cc
@@ -147,11 +147,12 @@ class TestReportErrorToClusterOp : public OpKernel {
                            "initialized properly."));
       return;
     }
     tensorflow::Status s(static_cast<tensorflow::error::Code>(error_code),
                          error_message);
-    s.SetPayload(tsl::CoordinationErrorPayloadKey(), "testing error payload");
+    s.SetPayload(tsl::CoordinationErrorPayloadKey(),
+                 absl::Cord("testing error payload"));
     OP_REQUIRES_OK(ctx, coord_agent->ReportError(s));
   }
 };
 REGISTER_KERNEL_BUILDER(Name("TestReportErrorToCluster").Device(DEVICE_DEFAULT),
                         TestReportErrorToClusterOp);
diff --git a/tensorflow/core/framework/dataset.h b/tensorflow/core/framework/dataset.h
index f895b610217ed..14fb62a6ef122 100644
--- a/tensorflow/core/framework/dataset.h
+++ b/tensorflow/core/framework/dataset.h
@@ -957,10 +957,17 @@ class IteratorBase : public Checkpointable {
     VLOG(1) << "Restored " << prefix() << " in "
             << (EnvTime::NowMicros() - start_us) << "us";
     return OkStatus();
   }
 
+  // Returns the total number of bytes buffered by the iterator across all nodes
+  // in the subtree for which autotuning is enabled.
+  int64_t TotalBufferedBytes() const {
+    if (node_) return node_->TotalBufferedBytes();
+    return 0;
+  }
+
  protected:
   // Returns a node that models this iterator.
   virtual std::shared_ptr<model::Node> CreateNode(
       IteratorContext* ctx, model::Node::Args args) const = 0;
 
diff --git a/tensorflow/core/framework/op_requires.h b/tensorflow/core/framework/op_requires.h
index a0baed6f41793..c5fb7796ecf6e 100644
--- a/tensorflow/core/framework/op_requires.h
+++ b/tensorflow/core/framework/op_requires.h
@@ -60,11 +60,11 @@ namespace tensorflow {
 #define OP_REQUIRES_OK_OR_SET_PAYLOAD(CTX, PAYLOAD_KEY, PAYLOAD_VALUE, STATUS) \
   do {                                                                         \
     if (!TF_PREDICT_TRUE(STATUS.ok())) {                                       \
       CheckNotInComputeAsync((CTX), "OP_REQUIRES_OK_ASYNC");                   \
       if (!PAYLOAD_VALUE.empty()) {                                            \
-        STATUS.SetPayload(PAYLOAD_KEY, PAYLOAD_VALUE);                         \
+        STATUS.SetPayload(PAYLOAD_KEY, absl::Cord(PAYLOAD_VALUE));             \
       }                                                                        \
       (CTX)->CtxFailureWithWarning(__FILE__, __LINE__, STATUS);                \
       return;                                                                  \
     }                                                                          \
   } while (0)
diff --git a/tensorflow/core/kernels/data/iterator_ops.cc b/tensorflow/core/kernels/data/iterator_ops.cc
index 65f9367b74038..0e0007ce8d66e 100644
--- a/tensorflow/core/kernels/data/iterator_ops.cc
+++ b/tensorflow/core/kernels/data/iterator_ops.cc
@@ -91,12 +91,10 @@ IteratorResource::IteratorResource(
       iterator_state_(std::make_shared<State>(std::move(flib_def),
                                               std::move(pflr), flr,
                                               /*iterator=*/nullptr)),
       output_dtypes_(output_dtypes),
       output_shapes_(output_shapes) {
-  tf_dataz_metrics_collector_ = std::make_shared<TfDatazMetricsCollector>(*env);
-  TfDatazMetricsRegistry::Register(tf_dataz_metrics_collector_);
   VLOG(2) << "creating iterator resource";
 }
 
 IteratorResource::~IteratorResource() {
   TfDatazMetricsRegistry::Deregister(tf_dataz_metrics_collector_);
@@ -272,10 +270,13 @@ Status IteratorResource::SetIteratorFromDataset(OpKernelContext* ctx,
       VerifyShapesCompatible(output_shapes_, iterator->output_shapes()));
   new_state->DowncastAndSetIteratorAndDataset(std::move(iterator), dataset);
   new_state->MergeCheckpoint(iter_ctx.checkpoint());
   mutex_lock l(mu_);
   std::swap(iterator_state_, new_state);
+  tf_dataz_metrics_collector_ =
+      std::make_shared<TfDatazMetricsCollector>(env_, iterator.get());
+  TfDatazMetricsRegistry::Register(tf_dataz_metrics_collector_);
   return OkStatus();
 }
 
 void IteratorResource::State::DowncastAndSetIteratorAndDataset(
     std::unique_ptr<IteratorBase> it, const DatasetBase* dataset) {
diff --git a/tensorflow/core/kernels/function_ops.cc b/tensorflow/core/kernels/function_ops.cc
index bf63ff82578a6..b46ad01dedae4 100644
--- a/tensorflow/core/kernels/function_ops.cc
+++ b/tensorflow/core/kernels/function_ops.cc
@@ -236,11 +236,11 @@ class SymbolicGradientOp : public AsyncOpKernel {
 
     FunctionLibraryRuntime::Handle handle;
     OP_REQUIRES_OK_ASYNC(
         ctx, lib->Instantiate(kGradientOp, AttrSlice(def()), &handle), done);
 
-    FunctionLibraryRuntime::Options opts;
+    FunctionLibraryRuntime::Options opts(ctx->step_id());
     opts.rendezvous = ctx->rendezvous();
     opts.cancellation_manager = ctx->cancellation_manager();
     opts.collective_executor = ctx->collective_executor();
     opts.runner = ctx->runner();
     opts.run_all_kernels_inline = ctx->run_all_kernels_inline();
diff --git a/tensorflow/core/kernels/functional_ops.cc b/tensorflow/core/kernels/functional_ops.cc
index 5f338f125839a..b20360873e2a3 100644
--- a/tensorflow/core/kernels/functional_ops.cc
+++ b/tensorflow/core/kernels/functional_ops.cc
@@ -158,11 +158,12 @@ class IfOp : public AsyncOpKernel {
           ctx_(ctx),
           cond_(cond),
           then_handle_(then_handle),
           else_handle_(else_handle),
           done_(std::move(done)),
-          lib_(CHECK_NOTNULL(ctx_->function_library())) {
+          lib_(CHECK_NOTNULL(ctx_->function_library())),
+          opts_(ctx->step_id()) {
       SetRunOptions(ctx_, &opts_, true /* always_collect_stats */);
       for (int i = 1; i < ctx_->num_inputs(); ++i) {
         args_.push_back(ctx_->input(i));
       }
     }
@@ -284,11 +285,12 @@ class CaseOp : public AsyncOpKernel {
         : kernel_(kernel),
           ctx_(ctx),
           branch_(branch),
           branch_handles_(branch_handles),
           done_(std::move(done)),
-          lib_(CHECK_NOTNULL(ctx_->function_library())) {
+          lib_(CHECK_NOTNULL(ctx_->function_library())),
+          opts_(ctx->step_id()) {
       SetRunOptions(ctx_, &opts_, true /* always_collect_stats */);
       for (int i = 1; i < ctx_->num_inputs(); ++i) {
         args_.push_back(ctx_->input(i));
       }
     }
@@ -505,11 +507,12 @@ class WhileOp : public AsyncOpKernel {
         : kernel_(kernel),
           ctx_(ctx),
           cond_handle_(cond_handle),
           body_handle_(body_handle),
           done_(std::move(done)),
-          lib_(CHECK_NOTNULL(ctx_->function_library())) {
+          lib_(CHECK_NOTNULL(ctx_->function_library())),
+          opts_(ctx->step_id()) {
       SetRunOptions(ctx_, &opts_, false /* always_collect_stats */);
       GetArgsFromContext(ctx, &args_, &loop_var_types_);
       body_frame_ =
           std::make_unique<BodyFuncCallFrame>(&args_, &rets_, loop_var_types_);
     }
@@ -749,10 +752,11 @@ class ForOp : public AsyncOpKernel {
     State(ForOp* kernel, OpKernelContext* ctx, DoneCallback done)
         : kernel_(kernel),
           ctx_(ctx),
           done_(std::move(done)),
           lib_(CHECK_NOTNULL(ctx_->function_library())),
+          opts_(ctx->step_id()),
           args_(1 + ctx_->num_inputs() - 3) {
       args_[0] = Tensor(DT_INT32, {});
       iter_ = &args_[0].scalar<int32>()();
 
       const int32_t num_loop_inputs = ctx_->num_inputs() - 3;
diff --git a/tensorflow/core/kernels/rnn/gru_ops.cc b/tensorflow/core/kernels/rnn/gru_ops.cc
index 702e618804da5..93aa88ca06733 100644
--- a/tensorflow/core/kernels/rnn/gru_ops.cc
+++ b/tensorflow/core/kernels/rnn/gru_ops.cc
@@ -47,65 +47,72 @@ class GRUCellBlockOp : public OpKernel {
     OP_REQUIRES_OK(ctx, ctx->input("b_ru", &b_ru_tensor));
 
     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;
     OP_REQUIRES_OK(
         ctx, ctx->allocate_output("r", TensorShape({batch_size, cell_size}),
@@ -202,99 +209,114 @@ class GRUBlockCellGradOp : public OpKernel {
     OP_REQUIRES_OK(ctx, ctx->input("c", &c_tensor));
 
     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. ",
                     batch_size));
     OP_REQUIRES(ctx, r_tensor->dim_size(1) == cell_size,
                 errors::InvalidArgument(
                     "r.dims(1) != cell_size: ", r_tensor->dim_size(1), " vs. ",
                     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. ",
                     batch_size));
     OP_REQUIRES(ctx, u_tensor->dim_size(1) == cell_size,
                 errors::InvalidArgument(
                     "u.dims(1) != cell_size: ", u_tensor->dim_size(1), " vs. ",
                     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. ",
                     batch_size));
     OP_REQUIRES(ctx, c_tensor->dim_size(1) == cell_size,
                 errors::InvalidArgument(
                     "c.dims(1) != cell_size: ", c_tensor->dim_size(1), " vs. ",
                     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),
                     " vs. ", batch_size));
     OP_REQUIRES(ctx, d_h_tensor->dim_size(1) == cell_size,
diff --git a/tensorflow/core/lib/core/status_test.cc b/tensorflow/core/lib/core/status_test.cc
index 9c93d2c1b3d2d..e53b1b9aac12d 100644
--- a/tensorflow/core/lib/core/status_test.cc
+++ b/tensorflow/core/lib/core/status_test.cc
@@ -175,27 +175,27 @@ TEST(StatusGroup, AggregateWithMultipleErrorStatus) {
                                 aborted.error_message()));
 }
 
 TEST(Status, InvalidPayloadGetsIgnored) {
   Status s = Status();
-  s.SetPayload("Invalid", "Invalid Val");
+  s.SetPayload("Invalid", absl::Cord("Invalid Val"));
   ASSERT_FALSE(s.GetPayload("Invalid").has_value());
   bool is_err_erased = s.ErasePayload("Invalid");
   ASSERT_EQ(is_err_erased, false);
 }
 
 TEST(Status, SetPayloadSetsOrUpdatesIt) {
   Status s(error::INTERNAL, "Error message");
-  s.SetPayload("Error key", "Original");
+  s.SetPayload("Error key", absl::Cord("Original"));
   ASSERT_EQ(s.GetPayload("Error key"), absl::Cord("Original"));
-  s.SetPayload("Error key", "Updated");
+  s.SetPayload("Error key", absl::Cord("Updated"));
   ASSERT_EQ(s.GetPayload("Error key"), absl::Cord("Updated"));
 }
 
 TEST(Status, ErasePayloadRemovesIt) {
   Status s(error::INTERNAL, "Error message");
-  s.SetPayload("Error key", "Original");
+  s.SetPayload("Error key", absl::Cord("Original"));
 
   bool is_err_erased = s.ErasePayload("Error key");
   ASSERT_EQ(is_err_erased, true);
   is_err_erased = s.ErasePayload("Error key");
   ASSERT_EQ(is_err_erased, false);
diff --git a/tensorflow/core/platform/error_payloads.cc b/tensorflow/core/platform/error_payloads.cc
index ecc0ede7e1059..f208ce7d57a17 100644
--- a/tensorflow/core/platform/error_payloads.cc
+++ b/tensorflow/core/platform/error_payloads.cc
@@ -25,10 +25,10 @@ void OkOrSetErrorCounterPayload(
   if (!status.ok() &&
       !status.GetPayload(tensorflow::kErrorSource).has_value()) {
     ErrorSourceProto error_source_proto;
     error_source_proto.set_error_source(error_source);
     status.SetPayload(tensorflow::kErrorSource,
-                      error_source_proto.SerializeAsString());
+                      absl::Cord(error_source_proto.SerializeAsString()));
   }
 }
 
 }  // namespace tsl
diff --git a/tensorflow/core/tpu/kernels/tpu_compile_op_common.cc b/tensorflow/core/tpu/kernels/tpu_compile_op_common.cc
index cd16fad2adb40..d6c9fef379901 100644
--- a/tensorflow/core/tpu/kernels/tpu_compile_op_common.cc
+++ b/tensorflow/core/tpu/kernels/tpu_compile_op_common.cc
@@ -406,11 +406,11 @@ Status TpuCompileOpKernelCommon::ComputeInternal(OpKernelContext* ctx) {
       }
     }
     SerializeToTString(proto, &output.scalar<tstring>()());
     ctx->set_output(0, output);
     status.SetPayload(TpuCompileInterface::kTpuCompileErrorPayloadKey,
-                      output.scalar<tstring>()());
+                      absl::Cord(output.scalar<tstring>()()));
   }
 
   if (status.ok()) {
     for (int i = 0; i < num_cores_with_compiled_programs; ++i) {
       Tensor output(DT_STRING, TensorShape({3}));
diff --git a/tensorflow/core/tpu/kernels/tpu_functional_ops.cc b/tensorflow/core/tpu/kernels/tpu_functional_ops.cc
index 3132238dbdc64..4cf5538ae9c16 100644
--- a/tensorflow/core/tpu/kernels/tpu_functional_ops.cc
+++ b/tensorflow/core/tpu/kernels/tpu_functional_ops.cc
@@ -1408,11 +1408,11 @@ Status TPUPartitionedCallOp::InitializeVarOnTPU(
       strings::StrCat(ndef->name(), "_init_ord_", device_ordinal);
 
   TF_RETURN_IF_ERROR(
       InstantiatePartition(*init_graph, fname, device, &fhandle, nullptr));
 
-  FunctionLibraryRuntime::Options opts;
+  FunctionLibraryRuntime::Options opts(ctx->step_id());
   opts.step_container = ctx->step_container();
   opts.cancellation_manager = ctx->cancellation_manager();
   opts.stats_collector = ctx->stats_collector();
 
   // Blocking on threads in the same thread pool is disallowed because
@@ -1567,11 +1567,11 @@ Status TPUPartitionedCallOp::InitializeShardedVarOnTPU(
     TF_RETURN_IF_ERROR(InstantiatePartition(*subgraph, function_name, target,
                                             &handle, nullptr));
     functions.push_back(DeviceAndFHandle{.device = target, .handle = handle});
   }
 
-  FunctionLibraryRuntime::Options opts;
+  FunctionLibraryRuntime::Options opts(ctx->step_id());
 
   // Blocking on threads in the same thread pool is disallowed because
   // concurrent warm-up requests can exhaust the default thread pool.
   // Create a new thread pool to initialize variables on TPU.
   std::function<void(std::function<void()>)> runner =
@@ -2700,11 +2700,11 @@ void TPUPartitionedCallOp::ExecuteLocalFunction(
 
 void TPUPartitionedCallOp::ExecuteFunctions(
     const std::vector<DeviceAndFHandle>& functions, OpKernelContext* ctx,
     int device_ordinal, int64_t ordinal_selector_req_id, DoneCallback done) {
   profiler::TraceMe trace_me("TPUPartitionedCallOp-ExecuteFunctions");
-  FunctionLibraryRuntime::Options opts;
+  FunctionLibraryRuntime::Options opts(ctx->step_id());
   opts.step_container = ctx->step_container();
   opts.stats_collector = ctx->stats_collector();
   // TODO(akshayka): Consider selecting a runner on a per-device basis,
   // i.e., using device-specific threadpools when available.
   opts.runner = ctx->runner();
diff --git a/tensorflow/core/tpu/tpu_embedding_errors.cc b/tensorflow/core/tpu/tpu_embedding_errors.cc
index f4c3759ab03fc..53efce986f174 100644
--- a/tensorflow/core/tpu/tpu_embedding_errors.cc
+++ b/tensorflow/core/tpu/tpu_embedding_errors.cc
@@ -27,11 +27,12 @@ Status AppendTpuEmbeddingErrorPayload(Status obj) {
   } else {
     const std::string error_message =
         absl::StrCat(kTpuEmbeddingErrorMessage, ". ", obj.error_message());
     Status status(obj.code(), error_message);
     TPUEmbeddingError error_payload;
-    status.SetPayload(kTpuEmbeddingErrorUrl, error_payload.SerializeAsString());
+    status.SetPayload(kTpuEmbeddingErrorUrl,
+                      absl::Cord(error_payload.SerializeAsString()));
     return status;
   }
 }
 
 bool HasTpuEmbeddingErrorPayload(const Status& status) {
diff --git a/tensorflow/core/tpu/tpu_embedding_errors.h b/tensorflow/core/tpu/tpu_embedding_errors.h
index a9b084ec8d84e..9c5d308b336b2 100644
--- a/tensorflow/core/tpu/tpu_embedding_errors.h
+++ b/tensorflow/core/tpu/tpu_embedding_errors.h
@@ -48,11 +48,12 @@ StatusOr<T> AppendTpuEmbeddingErrorPayload(StatusOr<T> obj) {
   } else {
     const std::string error_message = absl::StrCat(
         kTpuEmbeddingErrorMessage, ". ", obj.status().error_message());
     Status status(obj.status().code(), error_message);
     TPUEmbeddingError error_payload;
-    status.SetPayload(kTpuEmbeddingErrorUrl, error_payload.SerializeAsString());
+    status.SetPayload(kTpuEmbeddingErrorUrl,
+                      absl::Cord(error_payload.SerializeAsString()));
     return status;
   }
 }
 
 // Returns true if the tensorflow::Status obj has a payload of type
diff --git a/tensorflow/core/util/zen_util.h b/tensorflow/core/util/zen_util.h
index 47463bf6c4834..4910d3f4a5b68 100644
--- a/tensorflow/core/util/zen_util.h
+++ b/tensorflow/core/util/zen_util.h
@@ -21,22 +21,22 @@ limitations under the License.
 #include "absl/base/call_once.h"
 #include "tensorflow/core/util/env_var.h"
 
 namespace tensorflow {
 
-int64_t GetMempool() {
+inline int64_t GetMempool() {
   static absl::once_flag once;
   static int64_t mempool = 1;
   absl::call_once(once, [&] {
     TF_CHECK_OK(
         ReadInt64FromEnvVar("ZENDNN_ENABLE_MEMPOOL", mempool, &mempool));
     return mempool;
   });
   return mempool;
 }
 
-bool IsBlockedFormatEnabled() {
+inline bool IsBlockedFormatEnabled() {
   static absl::once_flag once;
   static bool blocked_format = false;
   absl::call_once(once, [&] {
     TF_CHECK_OK(ReadBoolFromEnvVar("ZENDNN_BLOCKED_FORMAT", blocked_format,
                                    &blocked_format));
diff --git a/tensorflow/lite/delegates/xnnpack/README.md b/tensorflow/lite/delegates/xnnpack/README.md
index e5e8458a70d59..130dd6fe8901b 100644
--- a/tensorflow/lite/delegates/xnnpack/README.md
+++ b/tensorflow/lite/delegates/xnnpack/README.md
@@ -452,27 +452,38 @@ Below is the list of currently supported floating-point operators:
 * Input, filter, bias (if present) and output tensors must be in 32-bit
   floating-point format.
 * Output size, filter and bias (if present) must be static (use
   `kTfLiteMmapRo` allocation type).
 
-### Floating-Point (IEEE FP16) Operators (experimental)
+### Floating-Point (IEEE FP16) Operators
 
 XNNPACK supports half-precision (using IEEE FP16 format) inference for a subset
 of floating-point operators. XNNPACK automatically enables half-precision
 inference when the following conditions are met:
 
 * XNNPACK runs on hardware that natively supports computations in IEEE FP16
-format. Currently, this hardware is limited to ARM64 devices with ARMv8.2 FP16
-arithmetics extension, and includes Android phones starting with Pixel 3,
-Galaxy S9 (Snapdragon SoC), Galaxy S10 (Exynos SoC), iOS devices with A11 or
-newer SoCs, and all Apple Silicon Macs.
+format. Currently, this hardware is limited to ARM & ARM64 devices with
+ARMv8.2 FP16 arithmetics extension, and includes Android phones starting with
+Pixel 3, Galaxy S9 (Snapdragon SoC), Galaxy S10 (Exynos SoC), iOS devices with
+A11 or newer SoCs, all Apple Silicon Macs, and Windows ARM64 laptops based with
+Snapdragon 850 SoC or newer.
 
 * IEEE FP16 inference is supported for every floating-point operator in the
 model.
 
 * The model's "reduced_precision_support" metadata indicates that the model
-is compatible with FP16 inference.
+is compatible with FP16 inference. The metadata can be added during model
+conversion using the `_experimental_supported_accumulation_type` attribute
+of the [tf.lite.TargetSpec](https://www.tensorflow.org/api_docs/python/tf/lite/TargetSpec)
+object:
+
+```python
+converter.optimizations = [tf.lite.Optimize.DEFAULT]
+...
+converter.target_spec.supported_types = [tf.float16]
+converter.target_spec._experimental_supported_accumulation_type = tf.dtypes.float16
+```
 
 When the above conditions are met, XNNPACK replace FP32 operators with their
 FP16 equivalents, and insert additional operators to convert model inputs
 from FP32 to FP16 and convert model outputs back from FP16 to FP32. If the
 above conditions are not met, XNNPACK will perform model inference with FP32
@@ -484,11 +495,11 @@ and in particular for testing end-to-end accuracy of model when FP16 inference
 is used. Forcing FP16 inference has several effects:
 
 * Besides ARM64 devices with ARMv8.2 FP16 arithmetics extension, forced FP16
 inference is supported on x86/x86-64 devices with AVX2 extension in emulation
 mode: all elementary floating-point operations are computed in FP32, then
-converted to FP16 and back to FP32. Note that such simulation is not exactly
+converted to FP16 and back to FP32. Note that such simulation is not bit-exact
 equivalent to native FP16 inference, but simulates the effects of restricted
 mantissa precision and exponent range in the native FP16 arithmetics.
 
 * On devices that support neither the native FP16 arithmetics (ARM64 devices
 with ARMv8.2 FP16 arithmetics extension), nor emulation (x86/x86-64 devices with
@@ -510,175 +521,14 @@ TfLiteXNNPackDelegateOptions xnnpack_options =
 xnnpack_options.flags |= TFLITE_XNNPACK_DELEGATE_FLAG_FORCE_FP16;
 TfLiteDelegate* xnnpack_delegate =
     TfLiteXNNPackDelegateCreate(&xnnpack_options);
 ```
 
-Below is the list of operators supported in IEEE FP16 inference:
-
-#### `ABS`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `ADD`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `AVERAGE_POOL_2D`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `CEIL`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `CONV_2D`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `CONCATENATION`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `DEPTH_TO_SPACE`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `DEPTHWISE_CONV_2D`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `DIV`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `FLOOR`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `FULLY_CONNECTED`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `HARD_SWISH`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `LEAKY_RELU`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `LOGISTIC`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `MAX_POOL_2D`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `MAXIMUM`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `MEAN`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `MINIMUM`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `MUL`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `NEG`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `PAD`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `PRELU`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `RELU`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `RELU6`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `RELU_N1_TO_1`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `RESHAPE`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `RESIZE_BILINEAR`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `ROUND`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SLICE`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SOFTMAX`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SPACE_TO_DEPTH`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SPLIT`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SQRT`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SQUARE`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SQUARED_DIFFERENCE`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `STRIDED_SLICE`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `SUB`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).
-
-#### `TRANSPOSE`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
-
-#### `TRANSPOSE_CONV`
-
-* Must satisfy constraints on the floating-point (FP32) operator.
+XNNPACK has full feature parity between FP32 and FP16 operators: all operators
+that are supported for FP32 inference are also supported for FP16 inference,
+and vice versa. In particular, sparse inference operators are supported for FP16
+inference on ARM processors.
 
 ### Quantized Operators
 
 By default, quantized inference in XNNPACK delegate is disabled, and XNNPACK is
 used only for floating-point models. Support for quantized inference in XNNPACK
@@ -853,11 +703,12 @@ Below is the list of currently supported quantized operators:
 
 ### Sparse Inference
 
 XNNPACK backend supports sparse inference for CNN models described in the
 [Fast Sparse ConvNets](https://arxiv.org/abs/1911.09723) paper. Sparse
-inference is restricted to subgraphs with the following operators:
+inference is restricted to subgraphs with the following floating-point
+operators:
 
 * Sparse subgraph must store its weights in sparse representation (using
   `DENSIFY` operators in the TensorFlow Lite schema).
 * Sparse subgraph must start with a 3x3 stride-2 `CONV_2D` operator with
   padding 1 on each side, no dilation, and 3 input channels.
diff --git a/tensorflow/python/data/experimental/kernel_tests/service/snapshot_ft_test.py b/tensorflow/python/data/experimental/kernel_tests/service/snapshot_ft_test.py
index 2df9faf26636c..02e63b87d325b 100644
--- a/tensorflow/python/data/experimental/kernel_tests/service/snapshot_ft_test.py
+++ b/tensorflow/python/data/experimental/kernel_tests/service/snapshot_ft_test.py
@@ -129,18 +129,20 @@ class SnapshotFtTest(data_service_test_base.TestBase, parameterized.TestCase):
       )
   )
   def testSnapshotRecoveryFailsWithBadSplitNames(self, bad_split_filename):
     cluster, _ = self.setup()
     write_file(os.path.join(self.source_dir(), bad_split_filename))
-    with self.assertRaisesRegex(ValueError, "can't parse"):
+    with self.assertRaisesRegex(
+        ValueError, "Expected split_<local_split_index>_<global_split_index>"):
       cluster.restart_dispatcher()
 
   @combinations.generate(test_base.eager_only_combinations())
   def testSnapshotRecoveryFailsWithOutOfOrderSplitName(self):
     cluster, _ = self.setup()
     write_file(os.path.join(self.source_dir(), "split_1_0"))
-    with self.assertRaisesRegex(ValueError, "found conflict"):
+    with self.assertRaisesRegex(
+        ValueError, "The local split index 1 exceeds the global split index 0"):
       cluster.restart_dispatcher()
 
   @combinations.generate(test_base.eager_only_combinations())
   def testSnapshotRecoveryFailsWithOutOfBoundsSplitName(self):
     cluster, _ = self.setup()
diff --git a/tensorflow/python/framework/BUILD b/tensorflow/python/framework/BUILD
index 50aa9e7b748b3..bd8375563a5dd 100644
--- a/tensorflow/python/framework/BUILD
+++ b/tensorflow/python/framework/BUILD
@@ -1092,10 +1092,12 @@ tf_py_test(
 
 py_library(
     name = "tensor_conversion_registry",
     srcs = ["tensor_conversion_registry.py"],
     srcs_version = "PY3",
+    # TODO(b/266747022): remove extra visibility
+    visibility = visibility + ["//learning/brain/experimental:__subpackages__"],
     deps = [
         "//tensorflow/python/eager:context",
     ],
 )
 
diff --git a/tensorflow/python/framework/errors_test_helper.cc b/tensorflow/python/framework/errors_test_helper.cc
index 833eb0e44843b..270e21f24101c 100644
--- a/tensorflow/python/framework/errors_test_helper.cc
+++ b/tensorflow/python/framework/errors_test_helper.cc
@@ -19,12 +19,12 @@ limitations under the License.
 namespace tensorflow {
 PYBIND11_MODULE(_errors_test_helper, m) {
   m.def("TestRaiseFromStatus", [](int code) {
     tensorflow::Status status(static_cast<tensorflow::error::Code>(code),
                               "test message");
-    status.SetPayload("key1", "value1");
-    status.SetPayload("key2", "value2");
+    status.SetPayload("key1", absl::Cord("value1"));
+    status.SetPayload("key2", absl::Cord("value2"));
     MaybeRaiseRegisteredFromStatus(status);
     return 0;
   });
 
   m.def("TestRaiseFromTFStatus", [](int code) {
diff --git a/tensorflow/python/tpu/tpu_strategy_util.py b/tensorflow/python/tpu/tpu_strategy_util.py
index 9daf746b99621..dd74c23f5c7e5 100644
--- a/tensorflow/python/tpu/tpu_strategy_util.py
+++ b/tensorflow/python/tpu/tpu_strategy_util.py
@@ -20,10 +20,12 @@ from tensorflow.core.protobuf import config_pb2
 from tensorflow.python.client import session as session_lib
 from tensorflow.python.distribute.cluster_resolver.tpu_cluster_resolver import TPUClusterResolver
 from tensorflow.python.eager import context
 from tensorflow.python.eager import monitoring
 from tensorflow.python.eager.def_function import function
+from tensorflow.python.eager.def_function import functions_run_eagerly
+from tensorflow.python.eager.def_function import run_functions_eagerly
 from tensorflow.python.framework import device
 from tensorflow.python.framework import errors
 from tensorflow.python.framework import ops
 from tensorflow.python.platform import tf_logging as logging
 from tensorflow.python.tpu import topology
@@ -109,20 +111,31 @@ def initialize_tpu_system(cluster_resolver=None):
           tpu_cancellation_closes_chips=False)
 
     # The TPU_SYSTEM device must match the device used in tpu.initialize_system
     # exactly, otherwise you can get errors if there are multiple TPU_SYSTEM
     # devices available.
+    run_eagerly = functions_run_eagerly()
+    if run_eagerly:
+      logging.warning(
+          "It looks like tf.function behavior was disabled, perhaps using"
+          " tf.config.run_functions_eagerly."
+          " tf.tpu.experimental.initialize_tpu_system requires tf.function to"
+          " work. This primitive will override the disable."
+      )
+    run_functions_eagerly(False)
     try:
       with ops.device(tpu._tpu_system_device_name(job)):  # pylint: disable=protected-access
         output = _tpu_init_fn()
       context.async_wait()
     except errors.InvalidArgumentError as e:
       raise errors.NotFoundError(
           None, None,
           "TPUs not found in the cluster. Failed in initialization: "
           + str(e))
-
+    finally:
+      if run_eagerly is not None:
+        run_functions_eagerly(run_eagerly)
     # Clear out the eager context caches since the memory is invalid now.
     context.context()._initialize_logical_devices()  # pylint: disable=protected-access
 
     serialized_topology = output.numpy()
   elif not ops.executing_eagerly_outside_functions():
@@ -219,12 +232,25 @@ def shutdown_tpu_system(cluster_resolver=None):
       tpu.shutdown_system(job=job)
 
     # The TPU_SYSTEM device must match the device used in tpu.shutdown_system
     # exactly, otherwise you can get errors if there are multiple TPU_SYSTEM
     # devices available.
-    with ops.device(tpu._tpu_system_device_name(job)):  # pylint: disable=protected-access
-      _tpu_shutdown_fn()
+    run_eagerly = functions_run_eagerly()
+    if run_eagerly:
+      logging.warning(
+          "It looks like tf.function behavior was disabled, perhaps using"
+          " tf.config.run_functions_eagerly."
+          " tf.tpu.experimental.shutdown_tpu_system requires tf.function to"
+          " work. This primitive will override the disable."
+      )
+    run_functions_eagerly(False)
+    try:
+      with ops.device(tpu._tpu_system_device_name(job)):  # pylint: disable=protected-access
+        _tpu_shutdown_fn()
+    finally:
+      if run_eagerly is not None:
+        run_functions_eagerly(run_eagerly)
 
     # Clear out the eager context caches since the memory is invalid now.
     logging.info("Clearing out eager caches")
     context.context()._clear_caches()  # pylint: disable=protected-access
     context.context().clear_kernel_cache()
diff --git a/tensorflow/tools/ci_build/release/requirements_common.txt b/tensorflow/tools/ci_build/release/requirements_common.txt
index f333f0614c42f..b4b785fc53d75 100644
--- a/tensorflow/tools/ci_build/release/requirements_common.txt
+++ b/tensorflow/tools/ci_build/release/requirements_common.txt
@@ -6,18 +6,18 @@ astunparse ~= 1.6.3
 flatbuffers ~= 2.0
 google_pasta ~= 0.2
 h5py ~= 3.8.0  # Earliest version for Python 3.11
 # TODO(b/262592253): Support older versions of NumPy for Python 3.10 and lower
 # to support TFX. Remove when Apache Beam upgrades to newer NumPy.
-numpy ~= 1.21.4; python_version < '3.11'
+numpy ~= 1.22.0; python_version < '3.11'
 numpy ~= 1.23.2; python_version >= '3.11' # Earliest version for Python 3.11
 opt_einsum ~= 3.3.0
 protobuf ~= 3.19.3  # NOTE: Earliest version for Python 3.10
 six ~= 1.16.0
 termcolor ~= 2.1.1
 typing_extensions ~= 3.10.0.0
-wheel ~= 0.36.2
+wheel ~= 0.38.1
 wrapt ~= 1.14.1
 
 # We need to pin the gast dependency exactly
 gast == 0.4.0
 
@@ -35,6 +35,6 @@ portpicker ~= 1.4.0
 scipy ~= 1.7.2; python_version < '3.11'
 scipy ~= 1.9.2; python_version >= '3.11' # Earliest version for Python 3.11
 
 # This is usually vendored in setuptools but ensure it gets installed in CI anyway
 # No bound here, we prefer the one in setuptools
-packaging
\ No newline at end of file
+packaging
diff --git a/tensorflow/tools/ci_build/release/requirements_mac.txt b/tensorflow/tools/ci_build/release/requirements_mac.txt
index 026844feb6463..0e0f832d088ea 100644
--- a/tensorflow/tools/ci_build/release/requirements_mac.txt
+++ b/tensorflow/tools/ci_build/release/requirements_mac.txt
@@ -1,9 +1,9 @@
 -r requirements_common.txt
 
 # Dependencies only required for Mac
-certifi ~= 2020.12.5
+certifi ~= 2022.12.07
 
 # Install build related dependencies
 twine ~= 3.6.0
 setuptools
 
diff --git a/tensorflow/tools/tf_sig_build_dockerfiles/devel.requirements.txt b/tensorflow/tools/tf_sig_build_dockerfiles/devel.requirements.txt
index e93f7f8105af9..f87eaff1228c1 100644
--- a/tensorflow/tools/tf_sig_build_dockerfiles/devel.requirements.txt
+++ b/tensorflow/tools/tf_sig_build_dockerfiles/devel.requirements.txt
@@ -9,19 +9,19 @@ astunparse ~= 1.6.3
 flatbuffers ~= 2.0
 google_pasta ~= 0.2
 h5py ~= 3.8.0 # Earliest version for Python 3.11
 # TODO(b/262592253): Support older versions of NumPy for Python 3.10 and lower
 # to support TFX. Remove when Apache Beam upgrades to newer NumPy.
-numpy ~= 1.21.4; python_version < '3.11'
+numpy ~= 1.22.0; python_version < '3.11'
 numpy ~= 1.23.2; python_version >= '3.11' # Earliest version for Python 3.11
 opt_einsum ~= 3.3.0
 packaging ~= 21.3
 protobuf ~= 3.20.1
 six ~= 1.16.0
 termcolor ~= 2.1.1
 typing_extensions ~= 3.10.0.0
-wheel ~= 0.36.2
+wheel ~= 0.38.1
 wrapt ~= 1.14.1
 # We need to pin the gast dependency exactly
 gast == 0.4.0
 # Finally, install tensorboard and estimator and keras
 # Note that here we want the latest version that matches TF major.minor version
@@ -46,6 +46,6 @@ PyYAML ~= 6.0
 auditwheel ~= 5.0.0
 twine ~= 3.6.0
 # For user tool scripts
 junitparser ~= 2.2.0
 lxml ~= 4.9.1
-pylint ~= 2.13.9
\ No newline at end of file
+pylint ~= 2.13.9
diff --git a/tensorflow/tsl/c/tsl_status.cc b/tensorflow/tsl/c/tsl_status.cc
index 82fdd7018e72e..0bddcc7e62c31 100644
--- a/tensorflow/tsl/c/tsl_status.cc
+++ b/tensorflow/tsl/c/tsl_status.cc
@@ -34,11 +34,11 @@ void TSL_SetStatus(TSL_Status* s, TSL_Code code, const char* msg) {
   }
   s->status = Status(static_cast<Code>(code), tsl::StringPiece(msg));
 }
 
 void TSL_SetPayload(TSL_Status* s, const char* key, const char* value) {
-  s->status.SetPayload(key, value);
+  s->status.SetPayload(key, absl::Cord(absl::string_view(value)));
 }
 
 void TSL_SetStatusFromIOError(TSL_Status* s, int error_code,
                               const char* context) {
   // TODO(b/139060984): Handle windows when changing its filesystem
diff --git a/tensorflow/tsl/c/tsl_status_helper_test.cc b/tensorflow/tsl/c/tsl_status_helper_test.cc
index e1496312f2e9c..dc0970d89f152 100644
--- a/tensorflow/tsl/c/tsl_status_helper_test.cc
+++ b/tensorflow/tsl/c/tsl_status_helper_test.cc
@@ -22,12 +22,12 @@ namespace tsl {
 namespace {
 
 TEST(StatusHelper, TestStatusHelper) {
   TSL_Status* s = TSL_NewStatus();
   Status cc_status(errors::InvalidArgument("some error"));
-  cc_status.SetPayload("key1", "value1");
-  cc_status.SetPayload("key2", "value2");
+  cc_status.SetPayload("key1", absl::Cord("value1"));
+  cc_status.SetPayload("key2", absl::Cord("value2"));
   Set_TSL_Status_from_Status(s, cc_status);
   ASSERT_EQ(TSL_INVALID_ARGUMENT, TSL_GetCode(s));
   ASSERT_EQ(std::string("some error"), TSL_Message(s));
 
   Status another_cc_status(StatusFromTSL_Status(s));
diff --git a/tensorflow/tsl/distributed_runtime/coordination/coordination_service_error_util.h b/tensorflow/tsl/distributed_runtime/coordination/coordination_service_error_util.h
index 829c221da8439..90db528558930 100644
--- a/tensorflow/tsl/distributed_runtime/coordination/coordination_service_error_util.h
+++ b/tensorflow/tsl/distributed_runtime/coordination/coordination_service_error_util.h
@@ -27,11 +27,11 @@ constexpr absl::string_view CoordinationErrorPayloadKey() {
 }
 
 // Mark error as a coordination service error (as opposed to RPC
 // errors).
 inline Status MakeCoordinationError(Status s) {
-  s.SetPayload(CoordinationErrorPayloadKey(), "");
+  s.SetPayload(CoordinationErrorPayloadKey(), absl::Cord(""));
   return s;
 }
 
 // Mark error as a coordination service error (as opposed to RPC
 // errors), and indicate error origin.
@@ -41,18 +41,20 @@ inline Status MakeCoordinationError(Status s,
                                     const tensorflow::CoordinatedTask& origin,
                                     bool is_reported_error = false) {
   tensorflow::CoordinationServiceError error;
   *error.mutable_source_task() = origin;
   error.set_is_reported_error(is_reported_error);
-  s.SetPayload(CoordinationErrorPayloadKey(), error.SerializeAsString());
+  s.SetPayload(CoordinationErrorPayloadKey(),
+               absl::Cord(error.SerializeAsString()));
   return s;
 }
 
 // Mark error as a coordination service error with payload.
 inline Status MakeCoordinationError(
     Status s, const tensorflow::CoordinationServiceError& payload) {
-  s.SetPayload(CoordinationErrorPayloadKey(), payload.SerializeAsString());
+  s.SetPayload(CoordinationErrorPayloadKey(),
+               absl::Cord(payload.SerializeAsString()));
   return s;
 }
 }  // namespace tsl
 
 #endif  // TENSORFLOW_TSL_DISTRIBUTED_RUNTIME_COORDINATION_COORDINATION_SERVICE_ERROR_UTIL_H_
diff --git a/tensorflow/tsl/distributed_runtime/rpc/grpc_util.h b/tensorflow/tsl/distributed_runtime/rpc/grpc_util.h
index e7bedfe500f60..a60ded995d766 100644
--- a/tensorflow/tsl/distributed_runtime/rpc/grpc_util.h
+++ b/tensorflow/tsl/distributed_runtime/rpc/grpc_util.h
@@ -69,16 +69,16 @@ inline std::string SerializePayloads(const Status& s) {
 
 inline void InsertSerializedPayloads(Status& s, std::string payloads) {
   tensorflow::distributed_runtime::GrpcPayloadContainer container;
   if (container.ParseFromString(payloads)) {
     for (const auto& key_val : container.payloads()) {
-      s.SetPayload(key_val.first, key_val.second);
+      s.SetPayload(key_val.first, absl::Cord(key_val.second));
     }
   } else {
     s.SetPayload(kGrpcPayloadsLost,
-                 tensorflow::distributed_runtime::GrpcPayloadsLost()
-                     .SerializeAsString());
+                 absl::Cord(tensorflow::distributed_runtime::GrpcPayloadsLost()
+                                .SerializeAsString()));
   }
 }
 
 inline Status FromGrpcStatus(const ::grpc::Status& s) {
   if (s.ok()) {
diff --git a/tensorflow/tsl/distributed_runtime/rpc/grpc_util_test.cc b/tensorflow/tsl/distributed_runtime/rpc/grpc_util_test.cc
index 46dbe7dc431eb..47f8efccb735f 100644
--- a/tensorflow/tsl/distributed_runtime/rpc/grpc_util_test.cc
+++ b/tensorflow/tsl/distributed_runtime/rpc/grpc_util_test.cc
@@ -69,11 +69,11 @@ TestRequest MakeProto(int size) {
   return proto;
 }
 
 TEST(PayloadSerialization, PayloadsAreTransmitted) {
   Status status = errors::InvalidArgument("invalid arg message");
-  status.SetPayload("a", "\\xFF\\x02\\x03");
+  status.SetPayload("a", absl::Cord("\\xFF\\x02\\x03"));
   Status status_recovered = FromGrpcStatus(ToGrpcStatus(status));
 
   ASSERT_TRUE(status_recovered.GetPayload("a").has_value());
   EXPECT_EQ(status_recovered.GetPayload("a").value(), "\\xFF\\x02\\x03");
 }
diff --git a/tensorflow/tsl/platform/errors.h b/tensorflow/tsl/platform/errors.h
index a7b29fea5c92b..02dfc74ebaba7 100644
--- a/tensorflow/tsl/platform/errors.h
+++ b/tensorflow/tsl/platform/errors.h
@@ -19,10 +19,11 @@ limitations under the License.
 #include <sstream>
 #include <string>
 #include <utility>
 
 #include "absl/base/attributes.h"
+#include "absl/strings/cord.h"
 #include "absl/strings/str_join.h"
 #include "tensorflow/tsl/platform/logging.h"
 #include "tensorflow/tsl/platform/macros.h"
 #include "tensorflow/tsl/platform/status.h"
 #include "tensorflow/tsl/platform/str_util.h"
@@ -100,19 +101,19 @@ inline std::unordered_map<std::string, std::string> GetPayloads(
 // payloads if they exist with the same key.
 inline void InsertPayloads(
     ::tsl::Status& status,
     const std::unordered_map<std::string, std::string>& payloads) {
   for (const auto& payload : payloads) {
-    status.SetPayload(payload.first, payload.second);
+    status.SetPayload(payload.first, absl::Cord(payload.second));
   }
 }
 
 // Copies all payloads from one Status to another. Will overwrite existing
 // payloads in the destination if they exist with the same key.
 inline void CopyPayloads(const ::tsl::Status& from, ::tsl::Status& to) {
   from.ForEachPayload([&to](tsl::StringPiece key, tsl::StringPiece value) {
-    to.SetPayload(key, value);
+    to.SetPayload(key, absl::Cord(value));
   });
 }
 
 // Creates a new status with the given code, message and payloads.
 inline ::tsl::Status Create(
diff --git a/tensorflow/tsl/platform/test.h b/tensorflow/tsl/platform/test.h
index 0a577b847d8b5..ef0f57cb218fb 100644
--- a/tensorflow/tsl/platform/test.h
+++ b/tensorflow/tsl/platform/test.h
@@ -83,9 +83,16 @@ int RandomSeed();
 
 // Returns an unused port number, for use in multi-process testing.
 // NOTE: This function is not thread-safe.
 int PickUnusedPortOrDie();
 
+// Constant which is false internally and true in open source.
+#ifdef PLATFORM_GOOGLE
+inline constexpr bool kIsOpenSource = false;
+#else
+inline constexpr bool kIsOpenSource = true;
+#endif  // PLATFORM_GOOGLE
+
 }  // namespace testing
 }  // namespace tsl
 
 #endif  // TENSORFLOW_TSL_PLATFORM_TEST_H_
diff --git a/tensorflow/tsl/profiler/lib/BUILD b/tensorflow/tsl/profiler/lib/BUILD
index 00037bd97d348..ed41b7dfa9d7d 100644
--- a/tensorflow/tsl/profiler/lib/BUILD
+++ b/tensorflow/tsl/profiler/lib/BUILD
@@ -211,10 +211,11 @@ cc_library(
     hdrs = ["traceme_encode.h"],
     visibility = ["//visibility:public"],
     deps = [
         "//tensorflow/tsl/platform:logging",
         "//tensorflow/tsl/platform:macros",
+        "@com_google_absl//absl/base:core_headers",
         "@com_google_absl//absl/strings",
     ],
 )
 
 tsl_cc_test(
@@ -224,10 +225,11 @@ tsl_cc_test(
         ":traceme_encode",
         "//tensorflow/tsl/platform",
         "//tensorflow/tsl/platform:test",
         "//tensorflow/tsl/platform:test_main",
         "@com_google_absl//absl/strings",
+        "@com_google_absl//absl/strings:str_format",
     ],
 )
 
 tf_profiler_pybind_cc_library_wrapper(
     name = "traceme_for_pybind",
diff --git a/tensorflow/tsl/profiler/lib/traceme_encode.h b/tensorflow/tsl/profiler/lib/traceme_encode.h
index 1e94f32e55c9b..dac5f08515a74 100644
--- a/tensorflow/tsl/profiler/lib/traceme_encode.h
+++ b/tensorflow/tsl/profiler/lib/traceme_encode.h
@@ -18,10 +18,11 @@ limitations under the License.
 #include <string.h>
 
 #include <initializer_list>
 #include <string>
 
+#include "absl/base/attributes.h"
 #include "absl/strings/match.h"
 #include "absl/strings/str_cat.h"
 #include "absl/strings/string_view.h"
 #include "tensorflow/tsl/platform/logging.h"
 #include "tensorflow/tsl/platform/macros.h"
@@ -29,18 +30,25 @@ limitations under the License.
 namespace tsl {
 namespace profiler {
 
 // An argument passed to TraceMeEncode.
 struct TraceMeArg {
-  // This constructor is required because absl::AlphaNum is non-copyable.
-  template <typename Value>
-  TraceMeArg(absl::string_view k, Value v) : key(k), value(v) {}
+  // String conversions of value types are supported via AlphaNum. We keep a
+  // reference to the AlphaNum's internal buffer here, so it must remain valid
+  // for the lifetime of this object. We cannot store it by value because it is
+  // not safe to construct an AlphaNum as a member of a class, particularly when
+  // AbslStringify is being used (it may reference default arguments that are on
+  // the caller's stack, if we constructed it here those default arguments would
+  // be destroyed before they are used).
+  TraceMeArg(absl::string_view k,
+             const absl::AlphaNum& v ABSL_ATTRIBUTE_LIFETIME_BOUND)
+      : key(k), value(v.Piece()) {}
 
   TF_DISALLOW_COPY_AND_ASSIGN(TraceMeArg);
 
   absl::string_view key;
-  absl::AlphaNum value;
+  absl::string_view value;
 };
 
 namespace traceme_internal {
 
 // Copies the contents of str to the address pointed by out.
@@ -72,11 +80,11 @@ TF_ATTRIBUTE_ALWAYS_INLINE inline std::string AppendArgs(
     char* out = begin + old_size;
     *out++ = '#';
     for (const auto& arg : args) {
       out = Append(out, arg.key);
       *out++ = '=';
-      out = Append(out, arg.value.Piece());
+      out = Append(out, arg.value);
       *out++ = ',';
     }
     *(out - 1) = '#';
     DCHECK_EQ(out, begin + new_size);
   }
diff --git a/tensorflow/tsl/profiler/lib/traceme_encode_test.cc b/tensorflow/tsl/profiler/lib/traceme_encode_test.cc
index f7fdf53ad3c8a..d058e27bc097f 100644
--- a/tensorflow/tsl/profiler/lib/traceme_encode_test.cc
+++ b/tensorflow/tsl/profiler/lib/traceme_encode_test.cc
@@ -15,10 +15,11 @@ limitations under the License.
 #include "tensorflow/tsl/profiler/lib/traceme_encode.h"
 
 #include <string>
 
 #include "absl/strings/str_cat.h"
+#include "absl/strings/str_format.h"
 #include "tensorflow/tsl/platform/platform.h"
 #include "tensorflow/tsl/platform/test.h"
 
 namespace tsl {
 namespace profiler {
@@ -51,10 +52,31 @@ TEST(TraceMeEncodeTest, TemporaryStringTest) {
                                      absl::StrCat("World:", 2020)}}),
             "Hello#context=World:2020#");
 }
 #endif
 
+// This can be removed when the absl version has been updated to include
+// AbslStringify for open source builds.
+#if defined(PLATFORM_GOOGLE)
+
+struct Point {
+  template <typename Sink>
+  friend void AbslStringify(Sink& sink, const Point& p) {
+    absl::Format(&sink, "(%d, %d)", p.x, p.y);
+  }
+
+  int x;
+  int y;
+};
+
+TEST(TraceMeEncodeTest, AbslStringifyTest) {
+  EXPECT_EQ(TraceMeEncode("Plot", {{"point", Point{10, 20}}}),
+            "Plot#point=(10, 20)#");
+}
+
+#endif
+
 TEST(TraceMeEncodeTest, NoNameTest) {
   EXPECT_EQ(TraceMeEncode({{"context", "World"}, {"request_id", 42}}),
             "#context=World,request_id=42#");
 }
 
diff --git a/third_party/tf_runtime/workspace.bzl b/third_party/tf_runtime/workspace.bzl
index 18d6a2b6ca334..2117290386637 100644
--- a/third_party/tf_runtime/workspace.bzl
+++ b/third_party/tf_runtime/workspace.bzl
@@ -4,12 +4,12 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls")
 
 def repo():
     """Imports TFRT."""
 
     # Attention: tools parse and update these lines.
-    TFRT_COMMIT = "c1248a015d23949afa2471bb21f6f52850aead7d"
-    TFRT_SHA256 = "8cdd8ea905478ac4ffd36ffb39cebe288d3b840d71a02d418bc6a8a760f92af8"
+    TFRT_COMMIT = "c653281a1a23c0c3d41536a983c7d10fcc5b1fbf"
+    TFRT_SHA256 = "3d1edd27c4e36d9cfc9493aef7088489babb370d2a7955bab3545acfbb024ccf"
 
     tf_http_archive(
         name = "tf_runtime",
         sha256 = TFRT_SHA256,
         strip_prefix = "runtime-{commit}".format(commit = TFRT_COMMIT),

Vulnerable code: build:mkl_aarch64_threadpool --define=build_with_acl=true
        "//tensorflow/compiler/xla/stream_executor/lib",
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
  error.SetPayload(kEvalErrorDetailUrl, error_payload);
        "//tensorflow/compiler/xla/stream_executor/lib",
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
      return tsl::Status(port::error::UNKNOWN, oss.str());              \
      return tsl::Status(port::error::UNKNOWN, oss.str());              \
      return tsl::Status(port::error::INTERNAL, error);
  return tsl::Status(port::error::INTERNAL,
      return tsl::Status(port::error::INVALID_ARGUMENT,
      return tsl::Status(port::error::INVALID_ARGUMENT, "max_seq_length <= 0");
      return tsl::Status(port::error::INVALID_ARGUMENT, "max_seq_length <= 0");
    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid input_h shape");
    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid input_c shape");
    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid output shape");
    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid output_h shape");
    return tsl::Status(port::error::INVALID_ARGUMENT, "Invalid output_c shape");
    return tsl::Status(port::error::INVALID_ARGUMENT,
        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
      return tsl::Status(port::error::INTERNAL, "Failed to start timer");
      return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
      return tsl::Status(port::error::INTERNAL, "Failed to start timer");
      return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
  return tsl::Status(port::error::INVALID_ARGUMENT,
  return tsl::Status(port::error::INTERNAL,
  return tsl::Status(port::error::INTERNAL,
  return tsl::Status(port::error::INTERNAL,
        port::error::INTERNAL,
        port::error::INTERNAL,
    return tsl::Status(port::error::INVALID_ARGUMENT,
        port::error::INTERNAL,
        port::error::INTERNAL,
    return tsl::Status(port::error::INVALID_ARGUMENT,
        port::error::INTERNAL,
        port::error::INTERNAL,
    return tsl::Status(port::error::INVALID_ARGUMENT,
      return tsl::Status(port::error::INVALID_ARGUMENT,
        port::error::INVALID_ARGUMENT,
        port::error::INVALID_ARGUMENT,
        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
            port::error::FAILED_PRECONDITION,
        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
          port::error::UNKNOWN,
        return tsl::Status(port::error::INTERNAL, "Failed to start timer");
      return tsl::Status(port::error::FAILED_PRECONDITION,
        return tsl::Status(port::error::INTERNAL, "Failed to stop timer");
    return tsl::Status(port::error::INVALID_ARGUMENT,
      return tsl::Status(port::error::UNKNOWN,
    return tsl::Status(port::error::INTERNAL,
    return tsl::Status(port::error::INTERNAL,
          port::error::INTERNAL,
    return tsl::Status(port::error::INVALID_ARGUMENT,
  return tsl::Status(port::error::INVALID_ARGUMENT,
    return tsl::Status(port::error::INVALID_ARGUMENT,
        port::error::INTERNAL,
    return tsl::Status(port::error::INTERNAL, "Cudnn pooling failed to split");
    return tsl::Status(port::error::INTERNAL, "Cudnn pooling failed to split");
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
  return tsl::Status(port::error::ABORTED,
  return tsl::Status(port::error::INTERNAL, message);
      port::error::INTERNAL,
    return tsl::Status(port::error::INVALID_ARGUMENT,
        port::error::INTERNAL,
    return tsl::Status(port::error::RESOURCE_EXHAUSTED,
        port::error::FAILED_PRECONDITION,
          port::error::UNAVAILABLE,
      port::error::INTERNAL,
            port::error::INTERNAL,
      port::error::INTERNAL,
        port::error::NOT_FOUND,
      port::error::INTERNAL,
        port::error::INTERNAL,
        port::error::INTERNAL,
      port::error::INTERNAL,
      port::error::INTERNAL,
        port::error::INTERNAL,
        port::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
        port::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
      port::error::NOT_FOUND,
        port::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
        "@com_google_absl//absl/base:core_headers",
        "//tensorflow/compiler/xla/stream_executor:stream_executor_pimpl",
        "//tensorflow/compiler/xla/stream_executor:timer",
        "//tensorflow/compiler/xla/stream_executor/lib",
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
        port::error::INTERNAL,
        "//tensorflow/compiler/xla/stream_executor/platform",
/* Copyright 2015 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.
==============================================================================*/


#ifndef TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_ERROR_H_
#define TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_ERROR_H_

#include "tensorflow/tsl/protobuf/error_codes.pb.h"  // IWYU pragma: export

namespace stream_executor {
namespace port {

namespace error = tensorflow::error;

}  // namespace port
}  // namespace stream_executor

#endif  // TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_ERROR_H_
/* Copyright 2015 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.
==============================================================================*/

#ifndef TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_INITIALIZE_H_
#define TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_INITIALIZE_H_

#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"

#endif  // TENSORFLOW_COMPILER_XLA_STREAM_EXECUTOR_LIB_INITIALIZE_H_
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
    return tsl::Status(port::error::INTERNAL,
        port::error::FAILED_PRECONDITION,
        port::error::FAILED_PRECONDITION,
        port::error::NOT_FOUND,
        port::error::NOT_FOUND,
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
    return tsl::Status(port::error::UNIMPLEMENTED,
  return tsl::Status(port::error::UNIMPLEMENTED,
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
        port::error::ALREADY_EXISTS,
          port::error::NOT_FOUND,
            port::error::FAILED_PRECONDITION,                                 \
      return tsl::Status(port::error::FAILED_PRECONDITION,                    \
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
    return tsl::Status(port::error::UNKNOWN,
    return tsl::Status(port::error::UNKNOWN,
    return tsl::Status(port::error::UNKNOWN,
    return tsl::Status(port::error::UNKNOWN,
      port::error::NOT_FOUND,
        port::error::INTERNAL,
        port::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/lib/error.h"
    FunctionLibraryRuntime::Options opts;
    int port, const std::string& protocol, const std::string server_type,
      const std::string server_type,
    std::vector<std::string> tokens = absl::StrSplit(split_filename, '_');
    int64_t local_split_index;
    int64_t global_split_index;
    if (tokens.size() != 3 ||
        !absl::SimpleAtoi(tokens[1], &local_split_index) ||
        local_split_index < 0 ||
        !absl::SimpleAtoi(tokens[2], &global_split_index) ||
        global_split_index < 0) {
      return InvalidArgument("can't parse the name of ", split_path);
    }
    if (local_split_index > global_split_index) {
      return InvalidArgument(
          "found conflict between local split index and global split index in ",
          "name of ", split_path);
    }
TfDatazMetricsCollector::TfDatazMetricsCollector(const Env& env)
    : latency_estimator_(env) {}
  explicit TfDatazMetricsCollector(const Env& env);
#include "tensorflow/core/framework/types.h"
    tfdataz_metrics_ = std::make_unique<TfDatazMetricsCollector>(*env_);
  auto collector_one =
      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
  auto collector_two =
      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
  auto collector_one =
      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
  auto collector_two =
      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
  auto collector_three =
      std::make_shared<TfDatazMetricsCollector>(*Env::Default());
                   error_source_proto.SerializeAsString());
    s.SetPayload(tsl::CoordinationErrorPayloadKey(), "testing error payload");
        STATUS.SetPayload(PAYLOAD_KEY, PAYLOAD_VALUE);                         \
  tf_dataz_metrics_collector_ = std::make_shared<TfDatazMetricsCollector>(*env);
  TfDatazMetricsRegistry::Register(tf_dataz_metrics_collector_);
    FunctionLibraryRuntime::Options opts;
          lib_(CHECK_NOTNULL(ctx_->function_library())) {
          lib_(CHECK_NOTNULL(ctx_->function_library())) {
          lib_(CHECK_NOTNULL(ctx_->function_library())) {
    const int64_t cell_size = h_prev_tensor->dim_size(1);

    // Sanity checks for input shapes.
    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));


    OP_REQUIRES(ctx, b_ru_tensor->dims() == 1,
                errors::InvalidArgument("Rank of b_ru must be 1",
                                        b_ru_tensor->dims(), " vs. 1", 1));
    OP_REQUIRES(ctx, b_c_tensor->dims() == 1,
                errors::InvalidArgument("Rank of b_c must be 1",
                                        b_c_tensor->dims(), " vs. 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]
    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));


    OP_REQUIRES(ctx, b_ru_tensor->dims() == 1,
                errors::InvalidArgument("Rank of b_ru must be 1",
                                        b_ru_tensor->dims(), " vs. 1"));

    OP_REQUIRES(ctx, b_c_tensor->dims() == 1,
                errors::InvalidArgument("Rank of b_c must be 1 ",
                                        b_c_tensor->dims(), " vs. 1"));

  s.SetPayload("Invalid", "Invalid Val");
  s.SetPayload("Error key", "Original");
  s.SetPayload("Error key", "Updated");
  s.SetPayload("Error key", "Original");
                      error_source_proto.SerializeAsString());
                      output.scalar<tstring>()());
  FunctionLibraryRuntime::Options opts;
  FunctionLibraryRuntime::Options opts;
  FunctionLibraryRuntime::Options opts;
    status.SetPayload(kTpuEmbeddingErrorUrl, error_payload.SerializeAsString());
    status.SetPayload(kTpuEmbeddingErrorUrl, error_payload.SerializeAsString());
int64_t GetMempool() {
bool IsBlockedFormatEnabled() {
### Floating-Point (IEEE FP16) Operators (experimental)
format. Currently, this hardware is limited to ARM64 devices with ARMv8.2 FP16
arithmetics extension, and includes Android phones starting with Pixel 3,
Galaxy S9 (Snapdragon SoC), Galaxy S10 (Exynos SoC), iOS devices with A11 or
newer SoCs, and all Apple Silicon Macs.
is compatible with FP16 inference.
converted to FP16 and back to FP32. Note that such simulation is not exactly
Below is the list of operators supported in IEEE FP16 inference:

#### `ABS`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `ADD`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `AVERAGE_POOL_2D`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `CEIL`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `CONV_2D`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `CONCATENATION`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `DEPTH_TO_SPACE`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `DEPTHWISE_CONV_2D`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `DIV`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `FLOOR`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `FULLY_CONNECTED`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `HARD_SWISH`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `LEAKY_RELU`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `LOGISTIC`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `MAX_POOL_2D`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `MAXIMUM`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `MEAN`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `MINIMUM`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `MUL`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `NEG`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `PAD`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `PRELU`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `RELU`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `RELU6`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `RELU_N1_TO_1`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `RESHAPE`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `RESIZE_BILINEAR`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `ROUND`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SLICE`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SOFTMAX`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SPACE_TO_DEPTH`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SPLIT`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SQRT`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SQUARE`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SQUARED_DIFFERENCE`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `STRIDED_SLICE`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `SUB`

* Must satisfy constraints on the floating-point (FP32) operator.
* Neither of the inputs can be static (use `kTfLiteMmapRo` allocation type).

#### `TRANSPOSE`

* Must satisfy constraints on the floating-point (FP32) operator.

#### `TRANSPOSE_CONV`

* Must satisfy constraints on the floating-point (FP32) operator.
inference is restricted to subgraphs with the following operators:
    with self.assertRaisesRegex(ValueError, "can't parse"):
    with self.assertRaisesRegex(ValueError, "found conflict"):
    status.SetPayload("key1", "value1");
    status.SetPayload("key2", "value2");

    with ops.device(tpu._tpu_system_device_name(job)):  # pylint: disable=protected-access
      _tpu_shutdown_fn()
numpy ~= 1.21.4; python_version < '3.11'
wheel ~= 0.36.2
packaging
certifi ~= 2020.12.5
numpy ~= 1.21.4; python_version < '3.11'
wheel ~= 0.36.2
pylint ~= 2.13.9
  s->status.SetPayload(key, value);
  cc_status.SetPayload("key1", "value1");
  cc_status.SetPayload("key2", "value2");
  s.SetPayload(CoordinationErrorPayloadKey(), "");
  s.SetPayload(CoordinationErrorPayloadKey(), error.SerializeAsString());
  s.SetPayload(CoordinationErrorPayloadKey(), payload.SerializeAsString());
      s.SetPayload(key_val.first, key_val.second);
                 tensorflow::distributed_runtime::GrpcPayloadsLost()
                     .SerializeAsString());
  status.SetPayload("a", "\\xFF\\x02\\x03");
    status.SetPayload(payload.first, payload.second);
    to.SetPayload(key, value);
  // This constructor is required because absl::AlphaNum is non-copyable.
  template <typename Value>
  TraceMeArg(absl::string_view k, Value v) : key(k), value(v) {}
  absl::AlphaNum value;
      out = Append(out, arg.value.Piece());
    TFRT_COMMIT = "c1248a015d23949afa2471bb21f6f52850aead7d"
    TFRT_SHA256 = "8cdd8ea905478ac4ffd36ffb39cebe288d3b840d71a02d418bc6a8a760f92af8"
Fixed code:         "//learning/brain/tfrt/tpu_plugin:__pkg__",
        "//tensorflow/compiler/xla/stream_executor/platform",
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
  error.SetPayload(kEvalErrorDetailUrl, absl::Cord(error_payload));
        "//tensorflow/compiler/xla/mlir/tools/mlir_replay/public:compiler_trace_proto_cc_impl",
  if (config.epilogue() != gpu::GemmBackendConfig::DEFAULT) {
    props.emplace_back(
        "epilogue", gpu::GemmBackendConfig::Epilogue_Name(config.epilogue()));
  }
        "//tensorflow/tsl/platform:errors",
        "//tensorflow/tsl/platform:errors",
        "//tensorflow/tsl/platform:errors",
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
      return tsl::Status(tsl::error::UNKNOWN, oss.str());               \
      return tsl::Status(tsl::error::UNKNOWN, oss.str());               \
      return tsl::Status(tsl::error::INTERNAL, error);
  return tsl::Status(tsl::error::INTERNAL,
      return tsl::Status(tsl::error::INVALID_ARGUMENT,
      return tsl::Status(tsl::error::INVALID_ARGUMENT, "max_seq_length <= 0");
      return tsl::Status(tsl::error::INVALID_ARGUMENT, "max_seq_length <= 0");
    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid input_h shape");
    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid input_c shape");
    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid output shape");
    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid output_h shape");
    return tsl::Status(tsl::error::INVALID_ARGUMENT, "Invalid output_c shape");
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
      return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
      return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
      return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
      return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
  return tsl::Status(tsl::error::INVALID_ARGUMENT,
  return tsl::Status(tsl::error::INTERNAL,
  return tsl::Status(tsl::error::INTERNAL,
  return tsl::Status(tsl::error::INTERNAL,
        tsl::error::INTERNAL,
        tsl::error::INTERNAL,
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
        tsl::error::INTERNAL,
        tsl::error::INTERNAL,
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
        tsl::error::INTERNAL,
        tsl::error::INTERNAL,
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
      return tsl::Status(tsl::error::INVALID_ARGUMENT,
        tsl::error::INVALID_ARGUMENT,
        tsl::error::INVALID_ARGUMENT,
        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
            tsl::error::FAILED_PRECONDITION,
        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
          tsl::error::UNKNOWN,
        return tsl::Status(tsl::error::INTERNAL, "Failed to start timer");
      return tsl::Status(tsl::error::FAILED_PRECONDITION,
        return tsl::Status(tsl::error::INTERNAL, "Failed to stop timer");
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
      return tsl::Status(tsl::error::UNKNOWN,
    return tsl::Status(tsl::error::INTERNAL,
    return tsl::Status(tsl::error::INTERNAL,
          tsl::error::INTERNAL,
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
  return tsl::Status(tsl::error::INVALID_ARGUMENT,
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
        tsl::error::INTERNAL,
    return tsl::Status(tsl::error::INTERNAL, "Cudnn pooling failed to split");
    return tsl::Status(tsl::error::INTERNAL, "Cudnn pooling failed to split");
#include "tensorflow/tsl/platform/errors.h"
  return tsl::Status(tsl::error::ABORTED,
  return tsl::Status(tsl::error::INTERNAL, message);
      tsl::error::INTERNAL,
    return tsl::Status(tsl::error::INVALID_ARGUMENT,
        tsl::error::INTERNAL,
    return tsl::Status(tsl::error::RESOURCE_EXHAUSTED,
        tsl::error::FAILED_PRECONDITION,
          tsl::error::UNAVAILABLE,
      tsl::error::INTERNAL,
            tsl::error::INTERNAL,
      tsl::error::INTERNAL,
        tsl::error::NOT_FOUND,
      tsl::error::INTERNAL,
        tsl::error::INTERNAL,
        tsl::error::INTERNAL,
      tsl::error::INTERNAL,
      tsl::error::INTERNAL,
        tsl::error::INTERNAL,
        tsl::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
        tsl::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
      tsl::error::NOT_FOUND,
        tsl::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
        "//tensorflow/tsl/platform:errors",
        "//tensorflow/compiler/xla/stream_executor:stream_executor_pimpl",  # fixdeps: keep
        "//tensorflow/tsl/platform:errors",
#include "tensorflow/tsl/platform/errors.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
        tsl::error::INTERNAL,
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
    return tsl::Status(tsl::error::INTERNAL,
        tsl::error::FAILED_PRECONDITION,
        tsl::error::FAILED_PRECONDITION,
        tsl::error::NOT_FOUND,
        tsl::error::NOT_FOUND,
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
    return tsl::Status(tsl::error::UNIMPLEMENTED,
  return tsl::Status(tsl::error::UNIMPLEMENTED,
#include "tensorflow/tsl/platform/errors.h"
        tsl::error::ALREADY_EXISTS,
          tsl::error::NOT_FOUND,
            tsl::error::FAILED_PRECONDITION,                                  \
      return tsl::Status(tsl::error::FAILED_PRECONDITION,                     \
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
#include "tensorflow/tsl/platform/errors.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
#include "tensorflow/compiler/xla/stream_executor/platform/initialize.h"
#include "tensorflow/tsl/platform/errors.h"
    return tsl::Status(tsl::error::UNKNOWN,
    return tsl::Status(tsl::error::UNKNOWN,
    return tsl::Status(tsl::error::UNKNOWN,
    return tsl::Status(tsl::error::UNKNOWN,
      tsl::error::NOT_FOUND,
        tsl::error::INTERNAL,
        tsl::error::INTERNAL,
    FunctionLibraryRuntime::Options opts(ctx->step_id());
        "//tensorflow/core:framework",
    int port, const std::string& protocol, const std::string& server_type,
      const std::string& server_type,
        "//tensorflow/tsl/platform:errors",
        "//tensorflow/tsl/platform:statusor",
        "//tensorflow/tsl/platform:status_matchers",
        "//tensorflow/tsl/protobuf:protos_all_cc",
#include <utility>
#include <vector>
#include "absl/strings/str_split.h"
#include "tensorflow/tsl/platform/errors.h"
#include "tensorflow/tsl/platform/statusor.h"
tsl::StatusOr<std::pair<int64_t, int64_t>> SplitIndex(
    absl::string_view split_path) {
  std::vector<std::string> tokens = absl::StrSplit(split_path, '_');
  int64_t local_split_index = 0, global_split_index = 0;
  if (tokens.size() != 3 || tokens[0] != "split" ||
      !absl::SimpleAtoi(tokens[1], &local_split_index) ||
      local_split_index < 0 ||
      !absl::SimpleAtoi(tokens[2], &global_split_index) ||
      global_split_index < 0) {
    return tsl::errors::InvalidArgument(
        "Invalid split file name: ", split_path,
        ". Expected split_<local_split_index>_<global_split_index>.");
  }
  if (local_split_index > global_split_index) {
    return tsl::errors::InvalidArgument(
        "Invalid split file name: ", split_path, ". The local split index ",
        local_split_index, " exceeds the global split index ",
        global_split_index, ".");
  }
  return std::make_pair(local_split_index, global_split_index);
}

#include <utility>
#include "tensorflow/tsl/platform/statusor.h"
// Returns a pair of {local_split_index, global_split_index} of the split. The
// expected format of `split_path` is:
// split_<local_split_index>_<global_split_index>
tsl::StatusOr<std::pair<int64_t, int64_t>> SplitIndex(
    absl::string_view split_path);

#include "tensorflow/tsl/platform/status_matchers.h"
#include "tensorflow/tsl/protobuf/error_codes.pb.h"
using ::testing::HasSubstr;
using ::testing::Pair;
using tsl::testing::IsOkAndHolds;
using tsl::testing::StatusIs;
TEST(PathUtilsTest, SplitIndex) {
  EXPECT_THAT(SplitIndex("split_0_1"), IsOkAndHolds(Pair(0, 1)));
}

TEST(PathUtilsTest, InvalidSplitFile) {
  EXPECT_THAT(
      SplitIndex(""),
      StatusIs(error::INVALID_ARGUMENT,
               HasSubstr(
                   "Expected split_<local_split_index>_<global_split_index>")));
  EXPECT_THAT(
      SplitIndex("split_123"),
      StatusIs(error::INVALID_ARGUMENT,
               HasSubstr(
                   "Expected split_<local_split_index>_<global_split_index>")));
  EXPECT_THAT(
      SplitIndex("split_-1_(-1)"),
      StatusIs(error::INVALID_ARGUMENT,
               HasSubstr(
                   "Expected split_<local_split_index>_<global_split_index>")));
  EXPECT_THAT(
      SplitIndex("split_5_0"),
      StatusIs(
          error::INVALID_ARGUMENT,
          HasSubstr(
              "The local split index 5 exceeds the global split index 0")));
}

    TF_ASSIGN_OR_RETURN(auto split_index, SplitIndex(split_filename));
    auto [local_split_index, global_split_index] = split_index;
TfDatazMetricsCollector::TfDatazMetricsCollector(const Env& env,
                                                 IteratorBase* iterator)
    : iterator_(iterator), latency_estimator_(env) {}
int64_t TfDatazMetricsCollector::GetIteratorTotalMemoryUsage() {
  return iterator_->TotalBufferedBytes();
}

#include "tensorflow/core/framework/dataset.h"
  TfDatazMetricsCollector(const Env& env, IteratorBase* iterator);
  // Returns the total memory (in bytes) used by the iterator.
  // Total memory used by the iterator includes the total number of bytes
  // buffered in all nodes in the subtree.
  int64_t GetIteratorTotalMemoryUsage();

  IteratorBase* iterator_;  // not owned
#include "tensorflow/core/framework/dataset.h"
    tfdataz_metrics_ =
        std::make_unique<TfDatazMetricsCollector>(*env_, iterator_.get());
  std::unique_ptr<IteratorBase> iterator_;
  std::unique_ptr<IteratorBase> iterator;
  auto collector_one = std::make_shared<TfDatazMetricsCollector>(
      *Env::Default(), iterator.get());
  auto collector_two = std::make_shared<TfDatazMetricsCollector>(
      *Env::Default(), iterator.get());
  std::unique_ptr<IteratorBase> iterator;
  auto collector_one = std::make_shared<TfDatazMetricsCollector>(
      *Env::Default(), iterator.get());
  auto collector_two = std::make_shared<TfDatazMetricsCollector>(
      *Env::Default(), iterator.get());
  auto collector_three = std::make_shared<TfDatazMetricsCollector>(
      *Env::Default(), iterator.get());
                   absl::Cord(error_source_proto.SerializeAsString()));
    s.SetPayload(tsl::CoordinationErrorPayloadKey(),
                 absl::Cord("testing error payload"));
  // Returns the total number of bytes buffered by the iterator across all nodes
  // in the subtree for which autotuning is enabled.
  int64_t TotalBufferedBytes() const {
    if (node_) return node_->TotalBufferedBytes();
    return 0;
  }

        STATUS.SetPayload(PAYLOAD_KEY, absl::Cord(PAYLOAD_VALUE));             \
  tf_dataz_metrics_collector_ =
      std::make_shared<TfDatazMetricsCollector>(env_, iterator.get());
  TfDatazMetricsRegistry::Register(tf_dataz_metrics_collector_);
    FunctionLibraryRuntime::Options opts(ctx->step_id());
          lib_(CHECK_NOTNULL(ctx_->function_library())),
          opts_(ctx->step_id()) {
          lib_(CHECK_NOTNULL(ctx_->function_library())),
          opts_(ctx->step_id()) {
          lib_(CHECK_NOTNULL(ctx_->function_library())),
          opts_(ctx->step_id()) {
          opts_(ctx->step_id()),
    // 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"));
    OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(h_prev_tensor->shape()),
                errors::InvalidArgument("Rank of h_prev must be 2, got ",
                                        h_prev_tensor->dims()));
    const int64_t cell_size = h_prev_tensor->dim_size(1);
    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, TensorShapeUtils::IsMatrix(w_c_tensor->shape()),
                errors::InvalidArgument("Rank of w_c must be 2, got ",
                                        w_c_tensor->dims()));
    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, TensorShapeUtils::IsVector(b_c_tensor->shape()),
                errors::InvalidArgument("Rank of b_c must be 1, got ",
                                        b_c_tensor->dims()));
    // 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()));
    // 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()));
    const int64_t cell_size = h_prev_tensor->dim_size(1);
    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, TensorShapeUtils::IsMatrix(w_c_tensor->shape()),
                errors::InvalidArgument("Rank of w_c must be 2, got ",
                                        w_c_tensor->dims()));
    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, TensorShapeUtils::IsVector(b_c_tensor->shape()),
                errors::InvalidArgument("Rank of b_c must be 1, got ",
                                        b_c_tensor->dims()));
    OP_REQUIRES(
        ctx, TensorShapeUtils::IsMatrix(r_tensor->shape()),
        errors::InvalidArgument("Rank of r must be 2, got ", r_tensor->dims()));
    OP_REQUIRES(
        ctx, TensorShapeUtils::IsMatrix(u_tensor->shape()),
        errors::InvalidArgument("Rank of u must be 2, got ", u_tensor->dims()));
    OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(c_tensor->shape()),
                errors::InvalidArgument("Rank of w_c must be 2, got ",
                                        c_tensor->dims()));
    OP_REQUIRES(ctx, TensorShapeUtils::IsMatrix(d_h_tensor->shape()),
                errors::InvalidArgument("Rank of d_h must be 2, got ",
                                        d_h_tensor->dims()));
  s.SetPayload("Invalid", absl::Cord("Invalid Val"));
  s.SetPayload("Error key", absl::Cord("Original"));
  s.SetPayload("Error key", absl::Cord("Updated"));
  s.SetPayload("Error key", absl::Cord("Original"));
                      absl::Cord(error_source_proto.SerializeAsString()));
                      absl::Cord(output.scalar<tstring>()()));
  FunctionLibraryRuntime::Options opts(ctx->step_id());
  FunctionLibraryRuntime::Options opts(ctx->step_id());
  FunctionLibraryRuntime::Options opts(ctx->step_id());
    status.SetPayload(kTpuEmbeddingErrorUrl,
                      absl::Cord(error_payload.SerializeAsString()));
    status.SetPayload(kTpuEmbeddingErrorUrl,
                      absl::Cord(error_payload.SerializeAsString()));
inline int64_t GetMempool() {
inline bool IsBlockedFormatEnabled() {
### Floating-Point (IEEE FP16) Operators
format. Currently, this hardware is limited to ARM & ARM64 devices with
ARMv8.2 FP16 arithmetics extension, and includes Android phones starting with
Pixel 3, Galaxy S9 (Snapdragon SoC), Galaxy S10 (Exynos SoC), iOS devices with
A11 or newer SoCs, all Apple Silicon Macs, and Windows ARM64 laptops based with
Snapdragon 850 SoC or newer.
is compatible with FP16 inference. The metadata can be added during model
conversion using the `_experimental_supported_accumulation_type` attribute
of the [tf.lite.TargetSpec](https://www.tensorflow.org/api_docs/python/tf/lite/TargetSpec)
object:

```python
converter.optimizations = [tf.lite.Optimize.DEFAULT]
...
converter.target_spec.supported_types = [tf.float16]
converter.target_spec._experimental_supported_accumulation_type = tf.dtypes.float16
```
converted to FP16 and back to FP32. Note that such simulation is not bit-exact
XNNPACK has full feature parity between FP32 and FP16 operators: all operators
that are supported for FP32 inference are also supported for FP16 inference,
and vice versa. In particular, sparse inference operators are supported for FP16
inference on ARM processors.
inference is restricted to subgraphs with the following floating-point
operators:
    with self.assertRaisesRegex(
        ValueError, "Expected split_<local_split_index>_<global_split_index>"):
    with self.assertRaisesRegex(
        ValueError, "The local split index 1 exceeds the global split index 0"):
    # TODO(b/266747022): remove extra visibility
    visibility = visibility + ["//learning/brain/experimental:__subpackages__"],
    status.SetPayload("key1", absl::Cord("value1"));
    status.SetPayload("key2", absl::Cord("value2"));
from tensorflow.python.eager.def_function import functions_run_eagerly
from tensorflow.python.eager.def_function import run_functions_eagerly
    run_eagerly = functions_run_eagerly()
    if run_eagerly:
      logging.warning(
          "It looks like tf.function behavior was disabled, perhaps using"
          " tf.config.run_functions_eagerly."
          " tf.tpu.experimental.initialize_tpu_system requires tf.function to"
          " work. This primitive will override the disable."
      )
    run_functions_eagerly(False)
    finally:
      if run_eagerly is not None:
        run_functions_eagerly(run_eagerly)
    run_eagerly = functions_run_eagerly()
    if run_eagerly:
      logging.warning(
          "It looks like tf.function behavior was disabled, perhaps using"
          " tf.config.run_functions_eagerly."
          " tf.tpu.experimental.shutdown_tpu_system requires tf.function to"
          " work. This primitive will override the disable."
      )
    run_functions_eagerly(False)
    try:
      with ops.device(tpu._tpu_system_device_name(job)):  # pylint: disable=protected-access
        _tpu_shutdown_fn()
    finally:
      if run_eagerly is not None:
        run_functions_eagerly(run_eagerly)
numpy ~= 1.22.0; python_version < '3.11'
wheel ~= 0.38.1
packaging
certifi ~= 2022.12.07
numpy ~= 1.22.0; python_version < '3.11'
wheel ~= 0.38.1
pylint ~= 2.13.9
  s->status.SetPayload(key, absl::Cord(absl::string_view(value)));
  cc_status.SetPayload("key1", absl::Cord("value1"));
  cc_status.SetPayload("key2", absl::Cord("value2"));
  s.SetPayload(CoordinationErrorPayloadKey(), absl::Cord(""));
  s.SetPayload(CoordinationErrorPayloadKey(),
               absl::Cord(error.SerializeAsString()));
  s.SetPayload(CoordinationErrorPayloadKey(),
               absl::Cord(payload.SerializeAsString()));
      s.SetPayload(key_val.first, absl::Cord(key_val.second));
                 absl::Cord(tensorflow::distributed_runtime::GrpcPayloadsLost()
                                .SerializeAsString()));
  status.SetPayload("a", absl::Cord("\\xFF\\x02\\x03"));
#include "absl/strings/cord.h"
    status.SetPayload(payload.first, absl::Cord(payload.second));
    to.SetPayload(key, absl::Cord(value));
// Constant which is false internally and true in open source.
#ifdef PLATFORM_GOOGLE
inline constexpr bool kIsOpenSource = false;
#else
inline constexpr bool kIsOpenSource = true;
#endif  // PLATFORM_GOOGLE

        "@com_google_absl//absl/base:core_headers",
        "@com_google_absl//absl/strings:str_format",
#include "absl/base/attributes.h"
  // String conversions of value types are supported via AlphaNum. We keep a
  // reference to the AlphaNum's internal buffer here, so it must remain valid
  // for the lifetime of this object. We cannot store it by value because it is
  // not safe to construct an AlphaNum as a member of a class, particularly when
  // AbslStringify is being used (it may reference default arguments that are on
  // the caller's stack, if we constructed it here those default arguments would
  // be destroyed before they are used).
  TraceMeArg(absl::string_view k,
             const absl::AlphaNum& v ABSL_ATTRIBUTE_LIFETIME_BOUND)
      : key(k), value(v.Piece()) {}
  absl::string_view value;
      out = Append(out, arg.value);
#include "absl/strings/str_format.h"
// This can be removed when the absl version has been updated to include
// AbslStringify for open source builds.
#if defined(PLATFORM_GOOGLE)

struct Point {
  template <typename Sink>
  friend void AbslStringify(Sink& sink, const Point& p) {
    absl::Format(&sink, "(%d, %d)", p.x, p.y);
  }

  int x;
  int y;
};

TEST(TraceMeEncodeTest, AbslStringifyTest) {
  EXPECT_EQ(TraceMeEncode("Plot", {{"point", Point{10, 20}}}),
            "Plot#point=(10, 20)#");
}

#endif

    TFRT_COMMIT = "c653281a1a23c0c3d41536a983c7d10fcc5b1fbf"
    TFRT_SHA256 = "3d1edd27c4e36d9cfc9493aef7088489babb370d2a7955bab3545acfbb024ccf"