// using library runtime `runtime` on a device `dev` with given
// `inputs_arg_shape_and_dtype` and `input_handles`.
absl::StatusOr<std::string> GetCompilerIr(
    IrExportStage stage, ProcessFunctionLibraryRuntime* pflr,
    absl::string_view func_name, Device* dev, EagerContext* context,
    absl::Span<const ArgShapeAndDType> input_arg_shape_and_dtype,
    absl::Span<const TensorHandle* const> input_handles,

  Graph* graph = options.graph->get();

  OptimizerOptions opts;
  auto pflr = std::make_unique<ProcessFunctionLibraryRuntime>(
      nullptr, options.session_options->env, /*config=*/nullptr,
      TF_GRAPH_DEF_VERSION, options.flib_def, opts);
  FunctionLibraryRuntime* flr =
      pflr->GetFLR(ProcessFunctionLibraryRuntime::kDefaultFLRDevice);

  for (Node* node : graph->nodes()) {

  Device* GetDevice(const string& device_name);
  FunctionLibraryRuntime* flr() { return flr_; }

 private:
  FunctionLibraryRuntime* flr_;
  std::unique_ptr<DeviceMgr> device_mgr_;
  std::unique_ptr<FunctionLibraryDefinition> lib_def_;
  std::unique_ptr<ProcessFunctionLibraryRuntime> pflr_;
};

}  // namespace tensorflow

namespace tensorflow {

bool XlaKernelCreator::CanCreateKernel(
    const FunctionLibraryRuntime& flr,
    const std::shared_ptr<const NodeProperties>& props) const {
  return CanCreateXlaKernel(props->node_def) &&
         !XlaOpRegistry::IsCompilationDevice(flr.device()->device_type());
}

static Status CreateXlaKernel(FunctionLibraryRuntime* flr,
                              const NodeDef& node_def,

    bool may_alias_resource_update, FunctionLibraryRuntime* flr,
    ResourceMgr* rm, const XlaCompiler::CompilationResult** compilation_result,
    xla::PjRtClient** client, xla::PjRtLoadedExecutable** executable) {
  PjRtDeviceCompiler* pjrt_device_compiler;
  DeviceCompilationProfiler* profiler;
  TF_RETURN_IF_ERROR(GetOrCreatePjRtDeviceCompilerAndProfiler(
      platform_info, rm, flr, &pjrt_device_compiler, &profiler));

  if (fdef.has_value()) {
    TF_CHECK_OK(lib_def_->AddFunctionDef(*fdef));
  }
  pflr_ = std::make_unique<ProcessFunctionLibraryRuntime>(
      device_mgr_.get(), Env::Default(), /*config=*/nullptr,
      TF_GRAPH_DEF_VERSION, lib_def_.get(), opts,
      /*default_thread_pool=*/nullptr, /*cluster_flr=*/nullptr);
  flr_ = pflr_->GetFLR("/job:localhost/replica:0/task:0/cpu:0");
}

// A function which expects R2 to hold TOC, and has a distinct local entry.
	.globl	toc_func
	.type	toc_func, @function
toc_func:
	addis	2,12,.TOC.-toc_func@ha
	addi	2,2,.TOC.-toc_func@l
	.localentry toc_func, .-toc_func
	mflr	0
	std	0,16(1)
	stdu	1,-32(1)

	// Call a NOTOC function which clobbers R2.
	bl	notoc_nor2_func
	nop

	// Call libc random. This should generate a TOC relative plt stub.
	bl	random
	nop

// computation.
// - `flr`: the FunctionLibraryRuntime for the `function`.
Status CompileToPjRtLoadedExecutable(
    const DeviceBase* device, const XlaPlatformInfo& platform_info,
    const NameAttrList& function,
    const std::vector<XlaCompiler::Argument>& args,
    DeviceCompileMode compile_mode, bool has_ref_vars,
    bool may_alias_resource_update, FunctionLibraryRuntime* flr,

  TF_EXPECT_OK(CompileToPjRtLoadedExecutable(
      device, platform_info, function, SampleArgsForAddXY(),
      DeviceCompileMode::kStrict, /*has_ref_vars=*/true,
      /*may_alias_resource_update=*/true, device_setup.flr(), &resource_mgr,
      &compilation_result, &pjrt_client, &pjrt_executable));

  EXPECT_TRUE(compilation_result != nullptr);
  EXPECT_TRUE(pjrt_executable != nullptr);
  EXPECT_TRUE(pjrt_client != nullptr);
}

// the persistor used in the DeviceCompiler. The platform ID from
// `platform_info` must not be null in CPU case.
Status BuildXlaDeviceCompiler(
    DeviceBase* dev, FunctionLibraryRuntime* flr,
    const XlaPlatformInfo& platform_info, DeviceType compilation_device_type,
    DeviceCompiler<xla::LocalExecutable, xla::LocalClient>**
        xla_device_compiler);