"tf.Yield"(%0) : (tensor<512x256xf32>) -> ()
    }, {
      "tf.Yield"(%arg2) : (tensor<512x256xf32>) -> ()
    }) { is_stateless = true}: (tensor<i1>) -> tensor<512x256xf32>
    "tf.Yield"(%1) : (tensor<512x256xf32>) -> ()
  }, {
    "tf.Yield"(%arg2) : (tensor<512x256xf32>) -> ()
  }) { is_stateless = true}: (tensor<i1>) -> tensor<512x256xf32>
  func.return %2 : tensor<512x256xf32>

    %0:2 = tf_executor.island wraps "tf.RecvTPUEmbeddingActivations"() {config = "test_config_recv_embedding"} : () -> tensor<512x256xf32> loc("RecvTPUEmbedding")
    %1:1 = tf_executor.island wraps "tf.SendTPUEmbeddingGradients"(%0) {N = 1 : i64, NN = 0 : i64, config = "test_config_send_embedding", operandSegmentSizes = array<i32: 1, 0>} : (tensor<512x256xf32>) -> () loc("SendTPUEmbedding")
    tf_executor.fetch
  }
  func.return
}

    "tf_device.cluster"() ({
      %0 = "tf.RecvTPUEmbeddingActivations"() {config = "test_config_recv_embedding"} : () -> tensor<512x256xf32>
      "tf.SendTPUEmbeddingGradients"(%0) {N = 1 : i64, NN = 0 : i64, config = "test_config_send_embedding", operandSegmentSizes = array<i32: 1, 0>} : (tensor<512x256xf32>) -> ()
      "tf.A"() {_xla_outside_compilation = "cluster1"} : () -> ()
      tf_device.return

module {
  func.func @main(%arg0: tensor<*x!tf_type.resource>, %arg1: tensor<*x!tf_type.resource>, %arg2: tensor<*x!tf_type.resource<tensor<512x256xf32>>>) {
    %cst = "tf.Const"() {value = dense<1> : tensor<i1>} : () -> tensor<i1>
    %0 = "tf.While"(%cst) {body = @while_body, cond = @while_cond, is_stateless = false} : (tensor<i1>) -> (tensor<i1>)
    return
  }

module {
  func.func @main(%arg0: tensor<*x!tf_type.resource>, %arg1: tensor<*x!tf_type.resource>, %arg2: tensor<*x!tf_type.resource<tensor<512x256xf32>>>) {
    %cst = "tf.Const"() {value = dense<1> : tensor<i1>} : () -> tensor<i1>
    %0 = "tf.While"(%cst) {body = @while_body, cond = @while_cond, is_stateless = false} : (tensor<i1>) -> (tensor<i1>)
    return
  }

}

func.func @einsum_ellipsis(%arg0: tensor<1x512x128xf32>, %arg1: tensor<128x256xf32>) -> tensor<1x512x256xf32> {
  %0 = "tf.Einsum"(%arg0, %arg1) {device = "", equation = "...x,xy->...y"} : (tensor<1x512x128xf32>, tensor<128x256xf32>) -> tensor<1x512x256xf32>
  func.return %0 : tensor<1x512x256xf32>
// CHECK-LABEL: einsum_ellipsis

  func.return
}

// -----

func.func @testSendTPUEmbeddingGradients(%x: tensor<512x256xf32>) {
  "tf.SendTPUEmbeddingGradients"(%x) {N = 1 : i64, NN = 0 : i64, config = "", operandSegmentSizes = array<i32: 1, 0>} : (tensor<512x256xf32>) -> ()
  func.return
}

// -----

//===--------------------------------------------------------------------===//

    %0 = "tfl.batch_matmul"(%arg0, %cst_0) {adj_x = false, adj_y = false} : (tensor<1x256xf32>, tensor<256x256xi8>) -> tensor<1x256xf32>
    %1 = "tfl.batch_matmul"(%0, %cst_1) {adj_x = false, adj_y = false} : (tensor<1x256xf32>, tensor<256x256x!quant.uniform<i8:f32, 1.000000e+00>>) -> tensor<1x256xf32>
    %2:2 = "tfl.while"(%cst_2, %1) ({
    ^bb0(%arg1: tensor<i32>,  %arg2: tensor<1x256xf32>):
      %cst_3 = arith.constant dense<10> : tensor<i32>

    }) {dimensions = dense<1> : tensor<1xi64>} : (tensor<1x256xf32>, tensor<f32>) -> tensor<1xf32>
    %5  = mhlo.add %3, %arg6 : tensor<1xf32>
    "mhlo.return"(%1, %arg3, %arg4, %arg5, %5) : (tensor<i32>, tensor<i32>, tensor<i32>, tensor<1x256xf32>, tensor<1xf32>) -> ()
  }) : (tensor<i32>, tensor<i32>, tensor<i32>, tensor<1x256xf32>, tensor<1xf32>) -> (tensor<i32>, tensor<i32>, tensor<i32>, tensor<1x256xf32>, tensor<1xf32>)

  %w = arith.constant dense<127.0> : tensor<512x2xf32>
  %mm = "tfl.batch_matmul"(%0, %w) {adj_x = false, adj_y = false} : (tensor<1x3x3x512xf32>, tensor<512x2xf32>) -> tensor<1x3x3x2xf32>
  %mm_s = "quantfork.stats"(%mm) {layerStats = dense<[0.000000e+00, 1.000000e+01]> : tensor<2xf32>} : (tensor<1x3x3x2xf32>) -> tensor<1x3x3x2xf32>