],
    deps = [
        ":xla_cpu_device",
        ":xla_cpu_jit",
        "//tensorflow/compiler/plugin",
    ] + if_cuda_or_rocm([
        ":xla_gpu_device",
        ":xla_gpu_jit",
    ]) + if_with_tpu_support([
        ":xla_tpu_device",
        ":xla_tpu_jit",
    ]),
    alwayslink = 1,
)

cc_library(
    name = "xla_cpu_jit",
    visibility = ["//visibility:public"],
    deps = [

  explicit Tf2XlaRewriterTestPeer(mlir::Operation* op)
      : op_builder_(op),
        empty_rewriter_(op_builder_),
        tf2xla_rewriter_(op, empty_rewriter_,
                         /*device_type=*/"XLA_CPU_JIT") {}

  absl::StatusOr<TupleOp> ImportXlaComputationIntoModule(
      XlaComputation& computation) {
    return tf2xla_rewriter_.ImportXlaComputation(computation);
  }

 private:
  OpBuilder op_builder_;

//     eliminates the use of resource variables.
//   . Legalizes the operations to XLA HLO operations.
//   . Canonicalizes the XLA HLO operations.
//
// device_type: XLA JIT device to use for compilation such as "XLA_CPU_JIT",
//   "XLA_GPU_JIT" or "XLA_TPU_JIT".
// use_tuple_args: when this is true, always create a tuple argument for the
//   entry computation.
// enable_op_fallback: when this is true, prefer tf2xla fallback kernels over

  mhlo::PopulateLegalizeTfPatterns(context, &patterns);
  TF::PopulateTFLoweringBeforeHLOPatterns(context, &patterns);
  mhlo::Tf2XlaTypeConverter converter;
  mhlo::PopulateLegalizeTfWithTf2XlaPatterns(
      "XLA_CPU_JIT", patterns, context, converter, /*prefer_tf2xla=*/false);
  stablehlo::StablehloToHloTypeConverter hlo_converter;
  chlo::populateChloToHloPatterns(context, &hlo_converter, &patterns);

static constexpr char kMlirToGraphCompilationCheckName[] =
    "mlir-to-graph-compilation-check";
// Use CPU arbitrarily in order to check that a graph compiles at all
static constexpr char kArbitraryDeviceName[] = "XLA_CPU_JIT";

namespace {
inline absl::string_view StringRefToView(llvm::StringRef ref) {
  return {ref.data(), ref.size()};
}
}  // namespace

static OwningOpRef<mlir::ModuleOp> GraphdefToMlirTranslateFunction(

      ParseArgumentShapes(mlir::StringRefToView(input_shapes), arg_shapes);
  if (!args_status.ok()) {
    LOG(ERROR) << args_status;
    return mlir::failure();
  }

  auto device_type = "XLA_CPU_JIT";
  llvm::MutableArrayRef<std::unique_ptr<mlir::Pass>>
      custom_legalization_passes{};
  XlaCompilationResult compilation_result;
  auto compilation_status =
      via_builder

// RUN: tf-opt "-xla-legalize-tf=device-type=XLA_CPU_JIT legalize-chlo=false use-tf2xla-fallback=true prefer-tf2xla=true" %s | FileCheck %s
// RUN: tf-opt "-xla-legalize-tf=device-type=XLA_CPU_JIT legalize-chlo=false prefer-tf2xla=true" %s | FileCheck --check-prefix NOFALLBACK %s

module attributes {tf.versions = {bad_consumers = [], min_consumer = 0 : i32, producer = 268 : i32}} {

// CHECK-LABEL: @abs
func.func @abs(%arg0: tensor<2xf32>) -> tensor<2xf32> {

  absl::StatusOr<XlaCompiler::CompilationResult> compile_result =
      LegalizeMlirToHlo(function_to_hlo_args, metadata_proto, use_tuple_args,
                        /*device_type=*/"XLA_CPU_JIT",
                        custom_legalization_passes,
                        /*shape_determination_fns=*/{}, arg_shapes,
                        &arg_core_mapping, &per_core_arg_shapes, client);

// RUN: tf-opt "-xla-legalize-tf=device-type=XLA_CPU_JIT prefer-tf2xla=true use-tf2xla-fallback=true" %s -verify-diagnostics -mlir-disable-threading  | FileCheck %s

module attributes {tf.versions = {bad_consumers = [], min_consumer = 0 : i32, producer = 268 : i32}} {
  // CHECK-LABEL: binary_op
  func.func @binary_op(%arg0: tensor<2xf32>, %arg1: tensor<2xf32>) -> tensor<2xf32> {
    // CHECK: mhlo.atan2

      "Convert TF dialect to HLO dialect (used for compilation in bridge).",
      [](mlir::OpPassManager& pm) {
        tensorflow::CreateConvertMlirToXlaHloPipeline(
            pm, /*device_type=*/"XLA_CPU_JIT", /*enable_op_fallback=*/false,
            /*custom_legalization_passes=*/{});
      });
}