Remove executable stack from xla_extension.so.

Fix jax-ml/jax#26781.

PiperOrigin-RevId: 732235187

tensorflow/lite/tools/cmake/modules/xnnpack.cmake

@@ -23,7 +23,7 @@ OverridableFetchContent_Declare(
   xnnpack
   GIT_REPOSITORY https://github.com/google/XNNPACK
   # Sync with tensorflow/workspace2.bzl
-  GIT_TAG 24794834234a7926d2f553d34e84204c8ac99dfd
+  GIT_TAG 5b4978cae19292232a27bdf0f495819bf5297167
   GIT_PROGRESS TRUE
   PREFIX "${CMAKE_BINARY_DIR}"
   SOURCE_DIR "${CMAKE_BINARY_DIR}/xnnpack"

third_party/stablehlo/temporary.patch

@@ -31,18 +31,6 @@ diff --ruN a/stablehlo/stablehlo/conversions/tosa/transforms/StablehloLegalizeTo
      return success();
    }
  };
-diff --ruN a/stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir b/stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir
---- stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir
-+++ stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir
-@@ -57,7 +57,7 @@
- // -----
-
- func.func @constant_with_dynamic_shape() -> tensor<1x?xf32, #stablehlo.bounds<?, 5>> {
--  // expected-error@+2 {{elements literal type must have static shape}}
-+  // expected-error@below {{elements literal type must have static shape}}
-   %c = stablehlo.constant dense<1> : tensor<1x?xf32, #stablehlo.bounds<?, 5>>
-   return %c : tensor<1x?xf32, #stablehlo.bounds<?, 5>>
- }
 diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir b/stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir
 --- stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir
 +++ stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir
@@ -62,39 +50,6 @@ diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplific
  }
 
  // -----
-diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir b/stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
---- stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
-+++ stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
-@@ -127,3 +127,15 @@
-     %4 = stablehlo.uniform_quantize %3 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
-     func.return %4: tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
- }
-+
-+// -----
-+
-+// CHECK-LABEL @failed_to_match_zero_defining_op
-+// CHECK{LITERAL}: %cst = stablehlo.constant dense<0.000000e+00> : tensor<2xf32>
-+// CHECK-NEXT: %0 = stablehlo.uniform_quantize %cst : (tensor<2xf32>) -> tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+// CHECK-NEXT: return %0 : tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+func.func @failed_to_match_zero_defining_op() -> tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>> {
-+  %0 = stablehlo.constant dense<0.000000e+00> : tensor<2xf32>
-+  %1 = stablehlo.uniform_quantize %0 : (tensor<2xf32>) -> tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+  return %1 : tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+}
-diff --ruN a/stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp b/stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
---- stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
-+++ stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
-@@ -63,6 +63,10 @@
-
-     // Collect quantized operands and result types to rewrite.
-     // All operands and results must be quantized
-+    if (computeOp->getNumOperands() == 0)
-+      return rewriter.notifyMatchFailure(computeOp,
-+                                         "requires non empty operands");
-+
-     llvm::SmallVector<Value> quantizedComputeOpOperands;
-     for (const Value& operand : computeOp->getOperands()) {
-       auto* definingOp = operand.getDefiningOp();
 diff --ruN a/stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp b/stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp
 --- stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp
 +++ stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp

third_party/stablehlo/workspace.bzl

@@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls")
 
 def repo():
     # LINT.IfChange
-    STABLEHLO_COMMIT = "5e9b356b928955a308d87746b1b32294f1f3eebe"
-    STABLEHLO_SHA256 = "6d47ba59d7433911d6adef12ddd48d0ee4d8a1c960d7fc21f84268be28cb7860"
+    STABLEHLO_COMMIT = "03597b1e592129f0c79e99e5ed65dac7ebee240f"
+    STABLEHLO_SHA256 = "0424d8bdb8f367efc28bc02f24229ba1b1c98a0dae47215fef53177038ba58e7"
     # LINT.ThenChange(Google-internal path)
 
     tf_http_archive(

third_party/xla/third_party/stablehlo/temporary.patch

@@ -31,18 +31,6 @@ diff --ruN a/stablehlo/stablehlo/conversions/tosa/transforms/StablehloLegalizeTo
      return success();
    }
  };
-diff --ruN a/stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir b/stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir
---- stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir
-+++ stablehlo/stablehlo/tests/ops_stablehlo_bounded_dynamism.mlir
-@@ -57,7 +57,7 @@
- // -----
-
- func.func @constant_with_dynamic_shape() -> tensor<1x?xf32, #stablehlo.bounds<?, 5>> {
--  // expected-error@+2 {{elements literal type must have static shape}}
-+  // expected-error@below {{elements literal type must have static shape}}
-   %c = stablehlo.constant dense<1> : tensor<1x?xf32, #stablehlo.bounds<?, 5>>
-   return %c : tensor<1x?xf32, #stablehlo.bounds<?, 5>>
- }
 diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir b/stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir
 --- stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir
 +++ stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplification.mlir
@@ -62,39 +50,6 @@ diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_aggressive_simplific
  }
 
  // -----
-diff --ruN a/stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir b/stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
---- stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
-+++ stablehlo/stablehlo/tests/transforms/stablehlo_legalize_qdq_to_quantized_op.mlir
-@@ -127,3 +127,15 @@
-     %4 = stablehlo.uniform_quantize %3 : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
-     func.return %4: tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>
- }
-+
-+// -----
-+
-+// CHECK-LABEL @failed_to_match_zero_defining_op
-+// CHECK{LITERAL}: %cst = stablehlo.constant dense<0.000000e+00> : tensor<2xf32>
-+// CHECK-NEXT: %0 = stablehlo.uniform_quantize %cst : (tensor<2xf32>) -> tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+// CHECK-NEXT: return %0 : tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+func.func @failed_to_match_zero_defining_op() -> tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>> {
-+  %0 = stablehlo.constant dense<0.000000e+00> : tensor<2xf32>
-+  %1 = stablehlo.uniform_quantize %0 : (tensor<2xf32>) -> tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+  return %1 : tensor<2x!quant.uniform<u8:f32, 3.400000e+01:16>>
-+}
-diff --ruN a/stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp b/stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
---- stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
-+++ stablehlo/stablehlo/transforms/StablehloLegalizeQDQToQuantizedOp.cpp
-@@ -63,6 +63,10 @@
-
-     // Collect quantized operands and result types to rewrite.
-     // All operands and results must be quantized
-+    if (computeOp->getNumOperands() == 0)
-+      return rewriter.notifyMatchFailure(computeOp,
-+                                         "requires non empty operands");
-+
-     llvm::SmallVector<Value> quantizedComputeOpOperands;
-     for (const Value& operand : computeOp->getOperands()) {
-       auto* definingOp = operand.getDefiningOp();
 diff --ruN a/stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp b/stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp
 --- stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp
 +++ stablehlo/stablehlo/transforms/optimization/StablehloAggressiveSimplification.cpp

third_party/xla/third_party/stablehlo/workspace.bzl

@@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls")
 
 def repo():
     # LINT.IfChange
-    STABLEHLO_COMMIT = "5e9b356b928955a308d87746b1b32294f1f3eebe"
-    STABLEHLO_SHA256 = "6d47ba59d7433911d6adef12ddd48d0ee4d8a1c960d7fc21f84268be28cb7860"
+    STABLEHLO_COMMIT = "03597b1e592129f0c79e99e5ed65dac7ebee240f"
+    STABLEHLO_SHA256 = "0424d8bdb8f367efc28bc02f24229ba1b1c98a0dae47215fef53177038ba58e7"
     # LINT.ThenChange(Google-internal path)
 
     tf_http_archive(

third_party/xla/tsl_workspace2.bzl

@@ -115,9 +115,9 @@ def _tf_repositories():
     # LINT.IfChange
     tf_http_archive(
         name = "XNNPACK",
-        sha256 = "face04056299ca22e2dbbf122a27aef289443dc7b1ad7e33a52714d6acc084eb",
-        strip_prefix = "XNNPACK-e55b3998cadb320188759aaada27328cbacc3253",
-        urls = tf_mirror_urls("https://github.com/google/XNNPACK/archive/e55b3998cadb320188759aaada27328cbacc3253.zip"),
+        sha256 = "9e290e7b094134bdda0cad4ef4b89625fbde3c4b8e8f5dc84044c0f2e55b875a",
+        strip_prefix = "XNNPACK-5b4978cae19292232a27bdf0f495819bf5297167",
+        urls = tf_mirror_urls("https://github.com/google/XNNPACK/archive/5b4978cae19292232a27bdf0f495819bf5297167.zip"),
     )
     # LINT.ThenChange(//tensorflow/lite/tools/cmake/modules/xnnpack.cmake)
 

third_party/xla/workspace2.bzl

@@ -45,9 +45,9 @@ def _tf_repositories():
     # LINT.IfChange
     tf_http_archive(
         name = "XNNPACK",
-        sha256 = "face04056299ca22e2dbbf122a27aef289443dc7b1ad7e33a52714d6acc084eb",
-        strip_prefix = "XNNPACK-e55b3998cadb320188759aaada27328cbacc3253",
-        urls = tf_mirror_urls("https://github.com/google/XNNPACK/archive/e55b3998cadb320188759aaada27328cbacc3253.zip"),
+        sha256 = "9e290e7b094134bdda0cad4ef4b89625fbde3c4b8e8f5dc84044c0f2e55b875a",
+        strip_prefix = "XNNPACK-5b4978cae19292232a27bdf0f495819bf5297167",
+        urls = tf_mirror_urls("https://github.com/google/XNNPACK/archive/5b4978cae19292232a27bdf0f495819bf5297167.zip"),
     )
     # LINT.ThenChange(//tensorflow/lite/tools/cmake/modules/xnnpack.cmake)
 

third_party/xla/xla/hlo/evaluator/hlo_evaluator.h

@@ -462,11 +462,13 @@ class HloEvaluator : public ConstDfsHloVisitorWithDefault {
   bool use_fast_path_reduce_ = true;
 
  private:
-  template <typename ReturnT, typename NativeT>
+  template <typename ReturnT, typename NativeT, typename UnaryOp>
   static absl::StatusOr<Literal> ElementWiseUnaryOpImpl(
-      const HloInstruction* instruction,
-      const std::function<ReturnT(NativeT)>& unary_op,
+      const HloInstruction* instruction, UnaryOp&& unary_op,
       const Literal& operand_literal) {
+    static_assert(std::is_invocable_r_v<ReturnT, UnaryOp, NativeT>,
+                  "Invalid UnaryOp signature");
+
     const Shape& shape = instruction->shape();
     const auto* operand = instruction->operand(0);
     TF_RET_CHECK(ShapeUtil::SameDimensions(shape, operand->shape()));

third_party/xla/xla/hlo/evaluator/hlo_evaluator_test.cc

@@ -6257,12 +6257,15 @@ BENCHMARK(BM_ReducePrecisely);
 static void BM_UnaryOp(benchmark::State& state) {
   int64_t d = state.range(0);
 
-  auto input = LiteralUtil::CreateFull({d, d}, 1.0f);
+  std::unique_ptr<HloInstruction> input =
+      HloInstruction::CreateConstant(LiteralUtil::CreateFull({d, d}, 1.0f));
+
+  std::unique_ptr<HloInstruction> unary = HloInstruction::CreateUnary(
+      ShapeUtil::MakeShape(F32, {d, d}), HloOpcode::kExp, input.get());
 
   HloEvaluator evaluator;
   for (auto s : state) {
-    CHECK_OK(
-        evaluator.EvaluateElementwiseUnaryOp(HloOpcode::kExp, input).status());
+    CHECK_OK(evaluator.Evaluate(unary.get()).status());
   }
 }
 
@@ -6276,13 +6279,17 @@ BENCHMARK(BM_UnaryOp)
 static void BM_BinaryOp(benchmark::State& state) {
   int64_t d = state.range(0);
 
-  auto lhs = LiteralUtil::CreateFull({d, d}, 1.0f);
-  auto rhs = LiteralUtil::CreateFull({d, d}, 2.0f);
+  std::unique_ptr<HloInstruction> lhs =
+      HloInstruction::CreateConstant(LiteralUtil::CreateFull({d, d}, 1.0f));
+  std::unique_ptr<HloInstruction> rhs =
+      HloInstruction::CreateConstant(LiteralUtil::CreateFull({d, d}, 2.0f));
+
+  std::unique_ptr<HloInstruction> binary = HloInstruction::CreateBinary(
+      ShapeUtil::MakeShape(F32, {d, d}), HloOpcode::kAdd, lhs.get(), rhs.get());
 
   HloEvaluator evaluator;
   for (auto s : state) {
-    CHECK_OK(evaluator.EvaluateElementwiseBinaryOp(HloOpcode::kAdd, lhs, rhs)
-                 .status());
+    CHECK_OK(evaluator.Evaluate(binary.get()).status());
   }
 }
 

third_party/xla/xla/hlo/evaluator/hlo_evaluator_typed_visitor.h

@@ -141,25 +141,29 @@ class HloEvaluatorTypedVisitor : public ConstDfsHloVisitorWithDefault {
  public:
   explicit HloEvaluatorTypedVisitor(HloEvaluator* p) : parent_(p) {}
 
-  // The following higher-order functions convert a function with ElementwiseT
-  // to a function with ReturnT.
-  std::function<ReturnT(ReturnT)> ConvertUnaryFunction(
-      const std::function<ElementwiseT(ElementwiseT)>& unary_op) {
+  // Converts a UnaryOp from a unary function on ElementwiseT to a unary
+  // function on ReturnT types.
+  template <typename UnaryOp>
+  auto ConvertUnaryFunction(const UnaryOp& unary_op) {
     return [&unary_op](ReturnT arg) {
       return static_cast<ReturnT>(unary_op(static_cast<ElementwiseT>(arg)));
     };
   }
-  std::function<ReturnT(ReturnT, ReturnT)> ConvertBinaryFunction(
-      const std::function<ElementwiseT(ElementwiseT, ElementwiseT)>&
-          binary_op) {
+
+  // Converts a BinaryOp from a binary function on ElementwiseT to a binary
+  // function on ReturnT types.
+  template <typename BinaryOp>
+  auto ConvertBinaryFunction(const BinaryOp& binary_op) {
     return [&binary_op](ReturnT arg1, ReturnT arg2) {
       return static_cast<ReturnT>(binary_op(static_cast<ElementwiseT>(arg1),
                                             static_cast<ElementwiseT>(arg2)));
     };
   }
-  std::function<ReturnT(ReturnT, ReturnT, ReturnT)> ConvertTernaryFunction(
-      const std::function<ElementwiseT(ElementwiseT, ElementwiseT,
-                                       ElementwiseT)>& ternary_op) {
+
+  // Converts a TernaryOp from a ternary function on ElementwiseT to a ternary
+  // function on ReturnT types.
+  template <typename TernaryOp>
+  auto ConvertTernaryFunction(const TernaryOp& ternary_op) {
     return [&ternary_op](ReturnT arg1, ReturnT arg2, ReturnT arg3) {
       return static_cast<ReturnT>(ternary_op(static_cast<ElementwiseT>(arg1),
                                              static_cast<ElementwiseT>(arg2),
@@ -748,10 +752,9 @@ class HloEvaluatorTypedVisitor : public ConstDfsHloVisitorWithDefault {
         }
         return std::min(high, std::max(value, low));
       };
-      TF_ASSIGN_OR_RETURN(
-          parent_->evaluated_[clamp],
-          ElementwiseTernaryOp(clamp,
-                               std::move(ConvertTernaryFunction(clamp_op))));
+      TF_ASSIGN_OR_RETURN(parent_->evaluated_[clamp],
+                          (ElementwiseTernaryOp<ReturnT, ReturnT, ReturnT>(
+                              clamp, ConvertTernaryFunction(clamp_op))));
       return absl::OkStatus();
     }
     return UnsupportedTypeError(clamp);
@@ -760,15 +763,12 @@ class HloEvaluatorTypedVisitor : public ConstDfsHloVisitorWithDefault {
   absl::Status HandleSelect(const HloInstruction* select) override {
     CHECK(!ShapeUtil::IsScalar(select->operand(0)->shape()));
     CHECK(select->shape().IsArray());
-    std::function<ReturnT(bool, ReturnT, ReturnT)> select_op =
-        [](bool pred, ReturnT on_true, ReturnT on_false) {
-          if (pred) {
-            return on_true;
-          }
-          return on_false;
-        };
+    auto select_op = [](bool pred, ReturnT on_true, ReturnT on_false) {
+      return pred ? on_true : on_false;
+    };
     TF_ASSIGN_OR_RETURN(parent_->evaluated_[select],
-                        ElementwiseTernaryOp(select, std::move(select_op)));
+                        (ElementwiseTernaryOp<bool, ReturnT, ReturnT>(
+                            select, std::move(select_op))));
     return absl::OkStatus();
   }
 
@@ -1646,23 +1646,29 @@ class HloEvaluatorTypedVisitor : public ConstDfsHloVisitorWithDefault {
   }
 
  private:
-  absl::StatusOr<Literal> ElementWiseUnaryOp(
-      const HloInstruction* instruction,
-      const std::function<ElementwiseT(ElementwiseT)>& unary_op) {
+  template <typename UnaryOp>
+  absl::StatusOr<Literal> ElementWiseUnaryOp(const HloInstruction* instruction,
+                                             UnaryOp&& unary_op) {
+    static_assert(std::is_invocable_r_v<ElementwiseT, UnaryOp, ElementwiseT>,
+                  "Invalid UnaryOp signature");
+
     const Literal& operand_literal =
         parent_->GetEvaluatedLiteralFor(instruction->operand(0));
     TF_ASSIGN_OR_RETURN(
         auto result_literal,
         (HloEvaluator::ElementWiseUnaryOpImpl<ReturnT, ReturnT>(
             instruction, ConvertUnaryFunction(unary_op), operand_literal)));
 
-    return std::move(result_literal);
+    return result_literal;
   }
 
-  absl::StatusOr<Literal> ElementWiseBinaryOp(
-      const HloInstruction* instruction,
-      const std::function<ElementwiseT(ElementwiseT, ElementwiseT)>&
-          binary_op) {
+  template <typename BinaryOp>
+  absl::StatusOr<Literal> ElementWiseBinaryOp(const HloInstruction* instruction,
+                                              BinaryOp&& binary_op) {
+    static_assert(std::is_invocable_r_v<ElementwiseT, BinaryOp, ElementwiseT,
+                                        ElementwiseT>,
+                  "Invalid BinaryOp signature");
+
     const auto& shape = instruction->shape();
     const auto* lhs = instruction->operand(0);
     const auto* rhs = instruction->operand(1);
@@ -1694,13 +1700,18 @@ class HloEvaluatorTypedVisitor : public ConstDfsHloVisitorWithDefault {
                 rhs_literal.Get<ReturnT>(multi_index));
           }));
     }
-    return std::move(result);
+
+    return result;
   }
 
-  template <typename LhsType, typename RhsType, typename EhsType>
+  template <typename LhsType, typename RhsType, typename EhsType,
+            typename TernaryOp>
   absl::StatusOr<Literal> ElementwiseTernaryOp(
-      const HloInstruction* instruction,
-      const std::function<ReturnT(LhsType, RhsType, EhsType)>& ternary_op) {
+      const HloInstruction* instruction, TernaryOp&& ternary_op) {
+    static_assert(
+        std::is_invocable_r_v<ReturnT, TernaryOp, LhsType, RhsType, EhsType>,
+        "Invalid TernaryOp signature");
+
     const auto& shape = instruction->shape();
     const auto* lhs = instruction->operand(0);
     const auto* rhs = instruction->operand(1);
@@ -1739,7 +1750,7 @@ class HloEvaluatorTypedVisitor : public ConstDfsHloVisitorWithDefault {
           }));
     }
 
-    return std::move(result);
+    return result;
   }
 
   template <typename NativeT>

third_party/xla/xla/hlo/tools/BUILD

@@ -97,20 +97,6 @@ xla_cc_binary(
     ],
 )
 
-xla_cc_binary(
-    name = "convert",
-    srcs = ["convert.cc"],
-    deps = [
-        "//xla/service:hlo_module_util",
-        "//xla/service:hlo_proto_util",
-        "//xla/tsl/util:command_line_flags",
-        "//xla/tsl/util:fixed_option_set_flag",
-        "@local_tsl//tsl/platform:logging",
-        "@local_tsl//tsl/platform:platform_port",
-        "@local_tsl//tsl/platform:protobuf",
-    ],
-)
-
 xla_cc_binary(
     name = "hlo_proto_to_json",
     srcs = ["hlo_proto_to_json.cc"],