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,

		},
		[]ProgHeader{
			{PT_PHDR, PF_R + PF_X, 0x40, 0x400040, 0x400040, 0x1c0, 0x1c0, 0x8},
			{PT_INTERP, PF_R, 0x200, 0x400200, 0x400200, 0x1c, 0x1c, 1},
			{PT_LOAD, PF_R + PF_X, 0x0, 0x400000, 0x400000, 0x684, 0x684, 0x200000},
			{PT_LOAD, PF_R + PF_W, 0x688, 0x600688, 0x600688, 0x210, 0x218, 0x200000},
			{PT_DYNAMIC, PF_R + PF_W, 0x6b0, 0x6006b0, 0x6006b0, 0x1a0, 0x1a0, 0x8},
			{PT_NOTE, PF_R, 0x21c, 0x40021c, 0x40021c, 0x20, 0x20, 0x4},

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

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

  EXPECT_EQ(options.device_type, compilation_device_type);
  EXPECT_EQ(options.device_ordinal, 0);
  EXPECT_NE(options.flib_def, nullptr);

      bool *has_outside_compilation) {
    OptimizerOptions opts;
    pflr_ = std::make_unique<ProcessFunctionLibraryRuntime>(
        device_mgr_.get(), Env::Default(), /*config=*/nullptr,
        TF_GRAPH_DEF_VERSION, fld, opts,
        /*default_thread_pool=*/nullptr);
    auto flr = pflr_->GetFLR("/job:localhost/replica:0/task:0/cpu:0");
    return ExtractOutsideCompilationForFunction(

      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,
                                      &xla_device_compiler));
  core::ScopedUnref xla_device_compiler_ref(xla_device_compiler);

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