class MHLO_ComparisonDirectionValue<string enumStr> :
  ConstantAttr<MHLO_ComparisonDirectionAttr, "::mlir::mhlo::ComparisonDirection::" # enumStr>;

class CHLO_ComparisonDirectionValue<string enumStr> :
  ConstantAttr<CHLO_ComparisonDirectionAttr, "::mlir::chlo::ComparisonDirection::" # enumStr>;

// TODO(b/228291745): Assert that $x and $y have the same shape.

      absl::Span<const xla::Literal* const> arguments) {
    // Replace args by mhlo.constant since the lowering passes can't lower
    // tf.Const.
    TF_ASSIGN_OR_RETURN(
        auto module_op,
        ReplaceFuncArgsByConstant(program, arguments, /*use_mhlo_const=*/true));

    // Run the Convert TF Quant Types, TF Quant -> MHLO Quant and MHLO Quant ->
    // MHLO int passes.
    PassManager pm(module_op->getContext());

           "Whether to serialize TPUCompileMetadataProto metadata in 'tf._TPUCompileMlir' op as a proto debug string">
  ];

  let dependentDialects = [
    "mlir::mhlo::MhloDialect",
    "mlir::tf_device::TensorFlowDeviceDialect"
  ];
  let constructor = "mlir::TFTPU::CreateTPURewritePass()";
}

// The TF dialect uses some TF types that are illegal in the MHLO dialect and
// some generic types that are legal in MHLO. This pass legalizes TF types into
// types that are legal in MHLO. For example, TF::Qint8Type is converted to i8.
// Rewrites here should run before TF to MHLO op legalizations are run.

#include <memory>
#include <string>
#include <utility>

    mlir::registerAllPasses();
    mlir::registerTensorFlowPasses();
    mlir::TFDevice::registerTensorFlowDevicePasses();
    mlir::mhlo::registerAllMhloPasses();
    // These are in compiler/mlir/xla and not part of the above MHLO
    // passes.
    mlir::mhlo::registerTfXlaPasses();
    mlir::mhlo::registerLegalizeTFPass();
    mlir::quant::stablehlo::registerBridgePasses();
    mlir::tosa::registerLegalizeTosaPasses();

#include "xla/client/xla_computation.h"
#include "xla/mlir_hlo/mhlo/IR/hlo_ops.h"
#include "xla/shape_util.h"
#include "xla/xla_data.pb.h"
#include "tensorflow/core/framework/op_kernel.h"
#include "tsl/lib/core/status_test_util.h"
#include "tsl/platform/errors.h"
#include "tsl/platform/status.h"
#include "tsl/platform/statusor.h"

namespace mlir {
namespace mhlo {

using ::mlir::LogicalResult;
using ::mlir::ModuleOp;

#include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h"
#include "tensorflow/compiler/mlir/tf2xla/transforms/utils.h"
#include "xla/mlir_hlo/mhlo/IR/hlo_ops.h"
#include "xla/mlir_hlo/utils/convert_op_folder.h"
#include "xla/mlir_hlo/utils/hlo_utils.h"
#include "xla/xla_data.pb.h"

namespace mlir {
namespace mhlo {

namespace {

constexpr absl::string_view kGroupSizeAttrName =
    "tf2xla.collective_info.group_size";

    // CHECK: %[[RQ:.*]] = stablehlo.uniform_quantize %[[DOT]]
    // CHECK: return %[[RQ]]

    %0 = stablehlo.dot_general %arg0, %arg1, batching_dims = [0, 1] x [0, 1], contracting_dims = [3] x [2] {mhlo.frontend_attributes = {grad_x = "false", grad_y = "false"}} : (tensor<2x3x1x1024xf32>, tensor<2x3x1024x3xf32>) -> tensor<2x3x1x3xf32>
    return %0 : tensor<2x3x1x3xf32>
  }
}

// -----

  module attributes {tf.versions = {bad_consumers = [], min_consumer = 0 : i32, producer = 268 : i32}} {
  func.func @main(%arg0: tensor<?xi32, #mhlo.type_extensions<bounds = [3]>> ) -> (tensor<?xi32, #mhlo.type_extensions<bounds = [3]>>) {
    func.return %arg0 : tensor<?xi32, #mhlo.type_extensions<bounds = [3]>>
  }
})";

  auto compilation_result = CompileMlirModule(
      kMlirModuleWithBoundedDynamicArgStr,

    module @module___inference_tpu_function_41 attributes {tf.versions = {bad_consumers = [], min_consumer = 0 : i32, producer = 1617 : i32}} {
      func.func @main(%arg0: tensor<2x2xf32>) -> (tensor<2x2xf32> {mhlo.sharding = "\08\03\1A\02\02\01\22\02\00\01"}) {
        %0 = tf_executor.graph {