namespace tensorflow {
namespace tfcompile {

using QuantizeXlaFn = std::function<Status(const tf2xla::Config& config,
                                           xla::XlaComputation* computation)>;

// Set the static quantization function to the `fn` if it hasn't been set.
// Return false if the static function has been set.
bool RegisterQuantizeFn(const QuantizeXlaFn& fn);

}  // namespace tfcompile

}

// Represents an entry in the XLA compile cache.
message XlaSerializedCacheEntry {
  // Used to uniqely identify this entry in its persisted representation.
  XlaSerializedCacheKey key = 1;

  // The computation (HLO) that compilation was done for. It is correlated to
  // the input TF graph so we can use it to fingerprint the compiled binary. We
  // serialize this rather than the input graphdef because it provides a

namespace tf2xla {
namespace internal {

// Legalize the given MLIR module to XLA HLO using a combination of the MLIR
// Bridge and XlaBuilder
absl::StatusOr<XlaCompilationResult> LegalizeTfToHlo(
    const tpu::MlirToHloArgs& computation,
    const tpu::TPUCompileMetadataProto& metadata, bool use_tuple_args,
    llvm::StringRef device_type,
    XlaShapeLayoutHelpers::ShapeDeterminationFns shape_determination_fns,

//
// Like qcasts, a dcast is allowed to have both its operand and result
// as non quantized types. This facilitates transformations and marks edges
// where the computation must be carried out in the expressed type.
//
// Especially early in transformation, it is common to have dcasts on
// all operands to ops that must operate with the expressed type (typically

    $mlir_module is a serialized MLIR module with a `main` function that contains target computation.
    $metadata is a serialized TPUCompileMetadataProto describing the shapes and types of the inputs to the computation, as well as a mapping onto the TPU pod topology.

	// Now compute mant*(2^e2)*(10^q).
	// Is it an exact computation?
	// Only small positive powers of 10 are exact (5^28 has 66 bits).
	exact := q <= 27 && q >= 0

	di, dexp2, d0 := mult64bitPow10(mant, e2, q)
	if dexp2 >= 0 {
		panic("not enough significant bits after mult64bitPow10")
	}
	// As a special case, computation might still be exact, if exponent

//
// If you modify this table, you MUST run 'go generate' to regenerate anames.go!
const (
	// Unprivileged ISA (Document Version 20190608-Base-Ratified)

	// 2.4: Integer Computational Instructions
	AADDI = obj.ABaseRISCV + obj.A_ARCHSPECIFIC + iota
	ASLTI
	ASLTIU
	AANDI
	AORI
	AXORI
	ASLLI
	ASRLI
	ASRAI
	ALUI
	AAUIPC
	AADD
	ASLT
	ASLTU
	AAND

	GoStart(2, x, y, 0, x / 2, y / 2, y, a, n);
	GoStart(3, x, y, x / 2, x, y / 2, y, a, n);
	GoWait(0);
	GoWait(1);
	GoWait(2);
	GoWait(3);
}

// The actual computation.  This is called in parallel.
void
DoStep(int xdim, int ydim, int xstart, int xend, int ystart, int yend, int *a, int *n)
{
	int x, y, c, i, j;

	for(x = xstart; x < xend; x++) {
		for(y = ystart; y < yend; y++) {

//
// Like qcasts, a dcast is allowed to have both its operand and result
// as non quantized types. This facilitates transformations and marks edges
// where the computation must be carried out in the expressed type.
//
// Especially early in transformation, it is common to have dcasts on
// all operands to ops that must operate with the expressed type (typically

                    **host_graph, fld);
  }

  return absl::OkStatus();
}

// Expand XLA computation's outside compilation host side graph into main graph.
// Add a control edge between sequencer node and the XLA computation node.
Status ExpandHostGraphIntoMainGraph(Graph* main_graph,
                                    FunctionLibraryDefinition* fld,