}
		case BlockRetJmp:
			if len(b.Succs) != 0 {
				f.Fatalf("retjmp block %s len(Succs)==%d, want 0", b, len(b.Succs))
			}
			if b.NumControls() != 1 {
				f.Fatalf("retjmp block %s has nil control", b)
			}
			if !b.Controls[0].Type.IsMemory() {
				f.Fatalf("retjmp block %s has non-memory control value %s", b, b.Controls[0].LongString())
			}
		case BlockPlain:
			if len(b.Succs) != 1 {

func (b *Block) removeSucc(i int) {
	n := len(b.Succs) - 1
	if i != n {
		e := b.Succs[n]
		b.Succs[i] = e
		// Update the other end of the edge we moved.
		e.b.Preds[e.i].i = i
	}
	b.Succs[n] = Edge{}
	b.Succs = b.Succs[:n]
	b.Func.invalidateCFG()
}

func (b *Block) swapSuccessors() {
	if len(b.Succs) != 2 {
		b.Fatalf("swapSuccessors with len(Succs)=%d", len(b.Succs))
	}
	e0 := b.Succs[0]

// JumpTable   [integer Value]  [succ1,succ2,..]

var genericBlocks = []blockData{
	{name: "Plain"},                  // a single successor
	{name: "If", controls: 1},        // if Controls[0] goto Succs[0] else goto Succs[1]
	{name: "Defer", controls: 1},     // Succs[0]=defer queued, Succs[1]=defer recovered. Controls[0] is call op (of memory type)

				// A variable is live on output from this block
				// if it is live on input to some successor.
				//
				// out[b] = \bigcup_{s \in succ[b]} in[s]
				newliveout.Copy(lv.blockEffects(b.Succs[0].Block()).livein)
				for _, succ := range b.Succs[1:] {
					newliveout.Or(newliveout, lv.blockEffects(succ.Block()).livein)
				}
			}

			if !be.liveout.Eq(newliveout) {
				change = true
				be.liveout.Copy(newliveout)
			}

		if int(ID) < len(f.Cache.blocks) {
			b = &f.Cache.blocks[ID]
			b.ID = ID
		} else {
			b = &Block{ID: ID}
		}
	}
	b.Kind = kind
	b.Func = f
	b.Preds = b.predstorage[:0]
	b.Succs = b.succstorage[:0]
	b.Values = b.valstorage[:0]
	f.Blocks = append(f.Blocks, b)
	f.invalidateCFG()
	return b
}

func (f *Func) freeBlock(b *Block) {
	if b.Func == nil {

    }
  }

  for (auto succ : GetSuccessors(op)) {
    if (is_cluster_op && !IsTpuRegularOp(succ) && has_cluster_metadata) {
      if (!CheckReplicatedIOOp(succ, metadata_map[cluster], op)) return false;
    }
    if (is_cluster_op && IsTpuRegularOp(succ)) {
      if (!CheckClusterSuccessors(succ, cluster, op, metadata_map))
        return false;
    }
    if (!is_cluster_op) {

// use[n] refers to values that are used at node n.
//
// Given a node n, variables' liveliness is defined like the following:
// live_in[n] = use[n] U (live_out[n] - def[n])
// live_out[n] = U {live_in[s] | s ε succ[n]}
//
// Consider a sequence of op:
//
// ```
// node 1: %0 = stablehlo.constant
// node 2: %1 = stablehlo.constant
// node 3: %2 = stablehlo.add %0, %1
// node 4: %3 = stablehlo.multiply %2, %1