// 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,

	{ET_EXEC, "ET_EXEC"},
	{EM_860, "EM_860"},
	{SHN_LOPROC, "SHN_LOPROC"},
	{SHT_PROGBITS, "SHT_PROGBITS"},
	{SHF_MERGE + SHF_TLS, "SHF_MERGE+SHF_TLS"},
	{PT_LOAD, "PT_LOAD"},
	{PF_W + PF_R + 0x50, "PF_W+PF_R+0x50"},
	{DT_SYMBOLIC, "DT_SYMBOLIC"},
	{DF_BIND_NOW, "DF_BIND_NOW"},
	{DF_1_PIE, "DF_1_PIE"},
	{NT_FPREGSET, "NT_FPREGSET"},
	{STB_GLOBAL, "STB_GLOBAL"},
	{STT_COMMON, "STT_COMMON"},

class XlaKernelCreator : public CustomKernelCreator {
 public:
  // Given a NodeDef 'node_def' and the function library runtime 'flr', returns
  // true if 'node_def' is a call to a compilable function defined in 'flr',
  // with the kXlaCompileAttr set.
  bool CanCreateKernel(
      const FunctionLibraryRuntime& flr,
      const std::shared_ptr<const NodeProperties>& props) const override;

    OptimizerOptions opts;
    device_mgr_ = std::make_unique<StaticDeviceMgr>(std::move(devices));
    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");
  }

  FunctionLibraryRuntime* flr_;

    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");
}

 * Called from standard ppc64 C ABI, where r2, r14-r31, f14-f31 are
 * callee-save, so they must be saved explicitly.
 */
.globl crosscall_ppc64
crosscall_ppc64:
	// Start with standard C stack frame layout and linkage
	mflr	%r0
	std	%r0, 16(%r1)	// Save LR in caller's frame
	mfcr	%r0
	std	%r0, 8(%r1)	// Save CR in caller's frame
	stdu	%r1, -FRAME_SIZE(%r1)
	std	%r2, 24(%r1)

	FOR_EACH_GPR std
	FOR_EACH_FPR stfd
	FOR_EACH_VR stvx