function_(function),
      platform_info_(XlaPlatformInfoFromDevice(ctx->device())),
      has_ref_vars_(has_ref_vars) {}

void XlaLocalLaunchBase::ComputeAsync(OpKernelContext* ctx, DoneCallback done) {
  VLOG(1) << "XlaLocalLaunchOpBase::Compute "
          << Canonicalize(function_.name(), AttrSlice(&function_.attr()));
  xla_launch_counter->GetCell(platform_info_.device_type().type_string())
      ->IncrementBy(1);

      GetAllocator(ctx->device(), stream, platform_info_);
  se::DeviceMemoryAllocator* allocator = allocator_ptr.get();
  XlaComputationLaunchContext launch_context(
      client, allocator, client->default_device_ordinal(),
      /*allocate_xla_tensors=*/platform_info_.xla_device_metadata() != nullptr,
      platform_info_.xla_device_metadata()
          ? platform_info_.xla_device_metadata()->UseMultipleStreams()

// vanilla TensorFlow op as long as the bridge supports it.
class XlaCompileOnDemandOp : public OpKernel {
 public:
  explicit XlaCompileOnDemandOp(OpKernelConstruction* ctx)
      : OpKernel(ctx),
        platform_info_(XlaPlatformInfoFromDevice(ctx->device())) {}
  void Compute(OpKernelContext* ctx) override;

 private:
  Status Compile(const std::vector<XlaCompiler::Argument>& args,
                 OpKernelContext* ctx,

  // Indexes of compile-time constant inputs
  const std::vector<int> constants_;
  // Indexes of resource inputs
  const std::vector<int> resources_;

  const NameAttrList function_;
  const XlaPlatformInfo platform_info_;

  bool has_ref_vars_;
};

// XlaLocalLaunchOp is used to replace a region of the TensorFlow graph
// which will be compiled and executed using XLA.  The XlaLocalLaunchOp is

Status GetCompilationDeviceTypeAndPjRtClient(
    const XlaPlatformInfo& platform_info, FunctionLibraryRuntime* flr,
    DeviceType* compilation_device_type, xla::PjRtClient** pjrt_client) {
  DeviceType device_type = platform_info.device_type();

  if (platform_info.xla_device_metadata()) {
    VLOG(2) << "Building PjRtDeviceCompiler using "
               "platform_info.xla_device_metadata().";

    *compilation_device_type =

  XlaPlatformInfo platform_info = XlaPlatformInfoFromDevice(device);

  TF_ASSERT_OK_AND_ASSIGN(
      DeviceType compilation_device_type,
      GetCompilationDeviceType(platform_info.device_type()));

  XlaDeviceCompiler* xla_device_compiler = nullptr;
  TF_EXPECT_OK(BuildXlaDeviceCompiler(device, device_setup_.flr(),
                                      platform_info, compilation_device_type,

  std::shared_ptr<se::DeviceMemoryAllocator> custom_allocator;
  XlaPlatformInfo platform_info(
      compilation_device_type, platform_id, xla_device_metadata.get(),
      /*pjrt_device_metadata=*/nullptr, custom_allocator);

  XlaCompiler::Options options = GenerateCompilerOptionsForPjRt(
      *device_setup_.flr(), device, platform_info,
      /*pjrt_device_compiler=*/nullptr);

  options.graph_def_version = function_library.graph_def_version();
  options.allow_cpu_custom_calls =
      (platform_info.platform_id() == se::host::kHostPlatformId);
  options.device_allocator = GetAllocator(device, stream, platform_info);
  if (platform_info.xla_device_metadata()) {
    options.shape_determination_fns =
        platform_info.xla_device_metadata()->default_shape_determination_fns();
  }

// Builds a DeviceCompiler that uses xla::LocalClient using `platform_info` and
// `compilation_device_type` (in non-TPU case) and sets *xla_device_compiler to
// point to it. Uses flags from `MarkForCompilationPassFlags` for configuring
// the persistor used in the DeviceCompiler. The platform ID from
// `platform_info` must not be null in CPU case.
Status BuildXlaDeviceCompiler(

  TF_ASSIGN_OR_RETURN(ResourceMgr * rm, GetResourceMgrForDeviceCompiler(
                                            ctx, platform_info.device_type()));
  return CompileToPjRtLoadedExecutable(
      ctx.device(), platform_info, function, args, compile_mode, has_ref_vars,
      may_alias_resource_update, ctx.function_library(), rm, compilation_result,
      client, executable);
}