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

    const NameAttrList& func_name_attrs, const string& new_func_name,
    const string& host_graph_func_name,
    const std::map<string, int>& host_compute_core, FunctionLibraryRuntime* flr,
    FunctionLibraryDefinition* fld, std::vector<string>* shape_inference_graphs,
    bool* has_outside_compilation);

// Rewrites XLA computation in `clusters` to replace outside compilation nodes

      host_compute_core, flr, fld, shape_inference_graphs,
      &then_branch_has_outside_compilation));
  TF_RETURN_IF_ERROR(ExtractOutsideCompilationForFunction(
      xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
      else_branch, else_branch_xla_func_name, else_branch_host_func_name,
      host_compute_core, flr, fld, shape_inference_graphs,

		if tlssize != 0 {
			ph := newElfPhdr()
			ph.Type = elf.PT_TLS
			ph.Flags = elf.PF_R
			ph.Memsz = tlssize
			ph.Align = uint64(ctxt.Arch.RegSize)
		}
	}

	if ctxt.HeadType == objabi.Hlinux || ctxt.HeadType == objabi.Hfreebsd {
		ph := newElfPhdr()
		ph.Type = elf.PT_GNU_STACK
		ph.Flags = elf.PF_W + elf.PF_R
		ph.Align = uint64(ctxt.Arch.RegSize)
	} else if ctxt.HeadType == objabi.Hopenbsd {

  std::unique_ptr<tensorflow::ScopedStepContainer> step_container_;
  std::unique_ptr<tensorflow::FunctionLibraryDefinition> flib_def_;
  std::unique_ptr<tensorflow::ProcessFunctionLibraryRuntime> pflr_;
  tensorflow::OpKernelContext::Params params_;

  xla::XlaBuilder xla_builder_;
};

}  // namespace mhlo
}  // namespace mlir

	OP_LD_R12_R12    = OP_LD | 12<<21 | 12<<16    // ld r12,0(r12)
	OP_MTCTR_R12     = OP_MTCTR | 12<<21          // mtctr r12
	OP_MFLR_R12      = OP_MFLR | 12<<21           // mflr r12
	OP_MFLR_R0       = OP_MFLR | 0<<21            // mflr r0
	OP_MTLR_R0       = OP_MTLR | 0<<21            // mtlr r0

	// This is a special, preferred form of bcl to obtain the next
	// instruction address (NIA, aka PC+4) in LR.

  FunctionLibraryRuntime::Handle handle;
  TF_RETURN_IF_ERROR(
      flr->Instantiate(function.name(), AttrSlice(&function.attr()), &handle));
  *fbody = flr->GetFunctionBody(handle);
  CHECK(*fbody);  // Can't be nullptr since we just instantiated it.
  const DataTypeVector& arg_types = (*fbody)->arg_types;
  std::vector<bool> const_args(arg_types.size());

      tensorflow::OpRegistry::Global(), tensorflow::FunctionDefLibrary());
  pflr_ = std::make_unique<tensorflow::ProcessFunctionLibraryRuntime>(
      device_mgr_.get(), tensorflow::Env::Default(), /*config=*/nullptr,
      version_or.value(), flib_def_.get(), tensorflow::OptimizerOptions());
  params_.function_library = pflr_->GetFLR(device_->name());
  return success();
}

// in `fbody` as well as the indices of its constant and resource arguments.
// `fbody` is owned by `flr`.
// `constant_arg_indices` and `resource_arg_indices` should be empty vector.
// They are sorted in ascending order on this function's return.
Status GetBodyAndConstantsAndResources(FunctionLibraryRuntime* flr,
                                       const NameAttrList& function,

	"cmp",
}

var ppcNeed = []string{
	"BR main.main(SB)",
	"CALL main.Println(SB)",
	"RET",
}

var ppcPIENeed = []string{
	"BR",
	"CALL",
	"RET",
}

var ppcGnuNeed = []string{
	"mflr",
	"lbz",
	"beq",
}

func mustHaveDisasm(t *testing.T) {
	switch runtime.GOARCH {
	case "loong64":
		t.Skipf("skipping on %s", runtime.GOARCH)
	case "mips", "mipsle", "mips64", "mips64le":