return success();
}

// Forms clusters with ops of the same `_replication_info` attribute under a
// block.
//
// For a given block, clusters are formed via grouping ops by
// `_replication_info` attributes. For every cluster formed:
//   1. Find associated TPUReplicateMetadata attributes with the same
//      `_replication_info` attribute.
//   2. Find users not in cluster that are interleaved between cluster ops.

    %control = tf_executor.island wraps "tf.TPUReplicateMetadata"() {_tpu_replicate = "cluster", allow_soft_placement = false, computation_shape...

        // already have a cluster assigned to them (via `_replication_info`
        // attribute value).
        // `tf.Identity` op is also treated as special tpu ops since it can play
        // a role as connection between `tf.TPUReplicatedInput` or
        // `tf.TPUPartitionedInput`. For example, we have the follow pseudocode:
        // %0 = tf_executor.island wraps "tf.OpA" (){_replication_info = 'c'}

    For example, a non replicated `tf_device.cluster_func`:

    ```mlir
    func @tf_tpu_rewrite(%arg0: tensor<i8>) {
      %0 = "tf_device.cluster_func"(%arg0) {_xla_compile_device_type = "TPU", _replication_info = "cluster0", func = @func} : (tensor<i8>) -> tensor<i8>
      return
    }
    ```

    will be rewritten as:

    ```mlir
    func @tf_tpu_rewrite(%arg0: tensor<i8>) {

  func.func @main() {

  // CHECK-LABEL: @converts_cluster
  func.func @converts_cluster() {
    // CHECK: "tf.XlaLaunch"()

  func.func @main(%arg0: tensor<*xf32> {tf.device = "/job:localhost/replica:0/task:0/device:CPU:0"}) {

```

wll be replaced by `_replication_info="cluster"` and  `_xla_compile_device_type="TPU"`.

```mlir

// Verifies that all MLIR Ops have the expected attributes.
std::unique_ptr<mlir::OperationPass<mlir::func::FuncOp>>
CreateVerifyClusteringPass();

// Creates a pass that forms clusters from operations of the same
// `_replication_info` attribute.
std::unique_ptr<mlir::OperationPass<mlir::ModuleOp>>
CreateTPUClusterFormationPass(bool strict_clusters = false);

// Creates a pass that extracts outside compilation (Host ops inside device

            tf_device.return %4#0, %4#1 : tensor<2048xi32>, tensor<2048xi32>
            }) {device = "TPU_REPLICATED_HOST_0"} : () -> (tensor<2048xi32>, tensor<2048xi32>)
      %1 = "tf_device.cluster_func"(%0#0, %0#1) {_replication_info = "cluster_test_fn", func = @tpu_func} : (tensor<2048xi32>, tensor<2048xi32>) -> tensor<2048xi32>
      return %1: tensor<2048xi32>
   }
   // CHECK-LABEL: func @tpu_func
   // CHECK: mhlo.type_extensions