return %cluster : tensor<i32>
    }
    ```

    will be transformed into:

    ```mlir
    func @computation(%arg0: tensor<i32>) -> tensor<i32> {
      %cluster = "tf_device.cluster_func"(%arg0) {func = @_func} : (tensor<i32>) -> tensor<i32>
      return %cluster : tensor<i32>
    }

    func @_func(%arg0: tensor<i32>) -> tensor<i32> {
      %identity = "tf.Identity"(%arg0) : (tensor<i32>) -> tensor<i32>

	pcdata = cursym.Func().SpillRegisterArgs(spill, newprog)

	call := obj.Appendp(pcdata, newprog)
	call.Pos = cursym.Func().Text.Pos
	call.As = obj.ACALL
	call.To.Type = obj.TYPE_BRANCH
	call.To.Name = obj.NAME_EXTERN
	morestack := "runtime.morestack"
	switch {
	case cursym.CFunc():
		morestack = "runtime.morestackc"
	case !cursym.Func().Text.From.Sym.NeedCtxt():

		// Run multiple times to shake out data races
		t.Run("goroutine launches", testLaunches)
	}
}

func BenchmarkGoroutine(b *testing.B) {
	withIdle := func(n int, fn func(b *testing.B)) func(b *testing.B) {
		return func(b *testing.B) {
			c := make(chan int)
			var ready, done sync.WaitGroup
			defer func() {
				close(c)
				done.Wait()
			}()

			for i := 0; i < n; i++ {
				ready.Add(1)

			}

			if autosize == 0 && c.cursym.Func().Text.Mark&LEAF == 0 {
				if c.cursym.Func().Text.From.Sym.NoSplit() {
					if ctxt.Debugvlog {
						ctxt.Logf("save suppressed in: %s\n", c.cursym.Name)
					}

					c.cursym.Func().Text.Mark |= LEAF
				}
			}

			p.To.Offset = int64(autosize) - ctxt.Arch.FixedFrameSize

			if c.cursym.Func().Text.Mark&LEAF != 0 {
				c.cursym.Set(obj.AttrLeaf, true)

        'a( a b c )'                                       | 'a'     | [tokens('a b c')]
        'a( a b, c d + f )'                                | 'a'     | [tokens('a b'), tokens('c d + f')]
        '_func(1+2)'                                       | '_func' | [tokens('1+2')]
        'a(A, B(1+2))'                                     | 'a'     | [token('A'), expression('B(1+2)')]

	}
	if s.CFunc() {
		fmt.Fprintf(ctxt.Bso, "cfunc ")
	}
	if s.NoSplit() {
		fmt.Fprintf(ctxt.Bso, "nosplit ")
	}
	if s.Func() != nil && s.Func().FuncFlag&abi.FuncFlagTopFrame != 0 {
		fmt.Fprintf(ctxt.Bso, "topframe ")
	}
	if s.Func() != nil && s.Func().FuncFlag&abi.FuncFlagAsm != 0 {
		fmt.Fprintf(ctxt.Bso, "asm ")
	}
	fmt.Fprintf(ctxt.Bso, "size=%d", s.Size)
	if s.Type == objabi.STEXT {

const (
	FOLL  = 1 << 0
	LABEL = 1 << 1
	LEAF  = 1 << 2
)

func preprocess(ctxt *obj.Link, cursym *obj.LSym, newprog obj.ProgAlloc) {
	autosize := int32(0)

	if cursym.Func().Text == nil || cursym.Func().Text.Link == nil {
		return
	}

	c := ctxt5{ctxt: ctxt, cursym: cursym, newprog: newprog}

	p := c.cursym.Func().Text
	autoffset := int32(p.To.Offset)
	if autoffset == -4 {

		fmt.Printf(" *\t%s\n", strings.Join(args, " "))
		fmt.Printf(" */\n")
	}
}

func (d *dstate) BeginPod(p pods.Pod) {
	d.mm = make(map[pkfunc]decodecounter.FuncPayload)
}

func (d *dstate) EndPod(p pods.Pod) {
	if d.cmd == debugDumpMode {
		d.cm.ResetModeAndGranularity()
	}
}

func (d *dstate) BeginCounterDataFile(cdf string, cdr *decodecounter.CounterDataReader, dirIdx int) {

    func.return %3 : tensor<?xi32>
  }

  // CHECK-LABEL: func @_func
  // CHECK-NOT: mhlo.is_same_data_across_replicas = true
  func.func @_func(%arg0: tensor<?xi32>, %arg1: tensor<?xi32>) -> tensor<?xi32> {
    %0 = "tf._D"(%arg0, %arg1) : (tensor<?xi32>, tensor<?xi32>) -> tensor<?xi32>
    func.return %0 : tensor<?xi32>
  }

  return %0 : tensor<1x3xf32>
}

// CHECK-LABEL: @_func
func.func private @_func(%arg0: tensor<1x3xf32>) -> (tensor<1x3xf32>) {
  return %arg0 : tensor<1x3xf32>
}
}


// -----
// SPMD: one input and no return
//
// CHECK-LABEL: func.func @serving_default(%arg0: tensor<1x3xf32>) {
// CHECK-NEXT:  "tf.IfrtCall"(%arg0)