message: [XLA] Update documentation on all-to-all instruction.
  Previously this discussed an old, non-tuple-shaped version of the instruction.
  No functional change.
diff: @@ -500,19 +500,34 @@ class HloInstruction {
  const std::vector<ReplicaGroup>& replica_groups,
  absl::string_view barrier, const absl::optional<int64>& all_reduce_id);
  - // This op handles the communication of an Alltoall operation. On each core,
  - // the operands are N ops in the same shape, where N is the number of cores
  - // participating the Alltoall. Then the N operands are scattered to N cores,
  - // e.g., the ith operand is sent to the ith core. Then each core gathers the
  - // received data into a tuple.
  + // An all-to-all op takes N array operands of the same shape and scatters them
  + // to N replicas. Each replica gathers the results into a tuple.
  //
  - // - `replica_groups`: each ReplicaGroup contains a list of replica id. If
  - // empty, all replicas belong to one group in the order of 0 - (n-1). Alltoall
  - // will be applied within subgroups in the specified order. For example,
  - // replica groups = {{1,2,3},{4,5,0}} means, an Alltoall will be applied
  - // within replica 1, 2, 3, and in the gather phase, the received blocks will
  - // be concatenated in the order of 1, 2, 3; another Alltoall will be applied
  - // within replica 4, 5, 0, and the concatenation order is 4, 5, 0.
  + // For example, suppose we have 3 replicas, with replica i passing inputs
  + // [a_i, b_i, c_i] to its all-to-all op. Then the resulting tuples are
  + //
  + // replica 0: (a_0, a_1, a_2)
  + // replica 1: (b_0, b_1, b_2)
  + // replica 2: (c_0, c_1, c_2).
  + //
  + // If replica_groups is set, the op is sharded and the replicas are permuted.
  + // To explain by way of example, suppose we have replica_groups={{1,2},{3,0}}.
  + // Then each replica passes two operands, say [a_i, b_i], and the result is
  + //
  + // replica 0: (b_3, b_0)
  + // replica 1: (a_1, a_2)
  + // replica 2: (b_1, b_2)
  + // replica 3: (a_3, a_0).
  + //
  + // All replica groups must have the same number of elements, and the number of
  + // operands must be equal to the size of one replica group. Each replica must
  + // appear in exactly one group.
  + //
  + // Note that this instruction is different than the all-to-all op in
  + // xla_builder.h. The version in XlaBuilder takes one input and slices it,
  + // and then concatenates the results into a single array. This instruction
  + // takes multiple inputs and returns a tuple; it doesn't slice or concatenate.
  + // It is used to implement the higher-level instruction in XlaBuilder.
  static std::unique_ptr<HloInstruction> CreateAllToAll(
  const Shape& shape, absl::Span<HloInstruction* const> operands,
  const std::vector<ReplicaGroup>& replica_groups);
  