ReplicaGroup group;
  group.add_replica_ids(0);
  group.add_replica_ids(1);
  XlaOp reduce_scatter =
      ReduceScatter(x, to_apply, /*scatter_dimension=*/1, /*shard_count=*/2,
                    /*replica_groups=*/{group});

  std::vector<XlaOp> tuple_values;
  tuple_values.push_back(reduce_scatter);
  xla::Tuple(&builder, tuple_values);

  TF_ASSERT_OK_AND_ASSIGN(XlaComputation computation, builder.Build());

      BuildReduceBody<AddOp>(element_type, &reduce_scatter.getComputation(),
                             &rewriter);
    } else if (reduce_op == "Mul") {
      BuildReduceBody<MulOp>(element_type, &reduce_scatter.getComputation(),
                             &rewriter);
    } else if (reduce_op == "Min") {
      BuildReduceBody<MinOp>(element_type, &reduce_scatter.getComputation(),
                             &rewriter);

      %cst = "tf.Const"() {value = dense<0> : tensor<i32>} : () -> tensor<i32>
      %cst_0 = "tf.Const"() {value = dense<[[0, 4], [1, 5], [2, 6], [3, 7]]> : tensor<4x2xi32>} : () -> tensor<4x2xi32>
      // CHECK:          "mhlo.reduce_scatter"(%arg0)
      // CHECK{LITERAL}: replica_groups = dense<[[0, 4], [1, 5], [2, 6], [3, 7]]>
      // CHECK-SAME:     scatter_dimension = 0
      //

    %cst = "tf.Const"() {value = dense<0> : tensor<i32>} : () -> tensor<i32>
    %cst_0 = "tf.Const"() {value = dense<[[0, 4], [1, 5], [2, 6], [3, 7]]> : tensor<4x2xi32>} : () -> tensor<4x2xi32>
    // CHECK:          "mhlo.reduce_scatter"(%arg0)
    // CHECK{LITERAL}: replica_groups = dense<[[0, 4], [1, 5], [2, 6], [3, 7]]>
    // CHECK-SAME:     scatter_dimension = 0
    //

}

def TF_XlaReduceScatterOp : TF_Op<"XlaReduceScatter", [Pure]> {
  let summary = "Wraps the XLA ReduceScatter operator";

  let description = [{
documented at https://www.tensorflow.org/xla/operation_semantics#reducescatter.
  }];

  let arguments = (ins

*   `tf.test`:

    *   Added `tf.test.experimental.sync_devices`, which is useful for accurately measuring performance in benchmarks.

*   `tf.experimental.dtensor`:

    *   Added experimental support to ReduceScatter fuse on GPU (NCCL).

### Bug Fixes and Other Changes

*   `tf.SavedModel`: