tensorflow/compiler/xla/service/spmd/dot_handler.cc
// Return if it would be better to match the LHS operand or RHS operand
// of a dot for non-contracting partitioning.
Return if it would be better to match the LHS operand or RHS operand
of a dot for non-contracting partitioning.
Return if it would be better to match the LHS operand or RHS operand of a dot for non - contracting partitioning .
bool LhsIsBestMatchForNonContractingPartitioning(
    const DotConvDimsMapping& dims_mapping, const PartitionedHlo& lhs,
    const PartitionedHlo& rhs, const Shape& output_base_shape,
    const HloSharding& output_sharding, const SpmdPartitionerOptions& options,
    int64_t num_partitions, int64_t lhs_non_contracting_partitions,
    int64_t rhs_non_contracting_partitions, int64_t lhs_matching_partitions,
    int64_t rhs_matching_partitions, int64_t lhs_contracting_partitions,
    int64_t rhs_contracting_partitions,
    int64_t output_lhs_non_contracting_partitions,
    int64_t output_rhs_non_contracting_partitions, int64_t lhs_batch_partitions,
    int64_t rhs_batch_partitions, SpmdBuilder* b, const Window& conv_window,
    const std::function<StatusOr<HloInstruction*>(
        HloInstruction*, HloInstruction*, SpmdBuilder*,
        const Window& conv_window)>& create_sharded_dot,
    SpmdPartitioningVisitor* visitor) {
  const bool may_group_on_lhs_non_contracting =
      lhs_non_contracting_partitions == output_lhs_non_contracting_partitions &&
      lhs_non_contracting_partitions > 1;
  const bool may_group_on_rhs_non_contracting =
      rhs_non_contracting_partitions == output_rhs_non_contracting_partitions &&
      rhs_non_contracting_partitions > 1;
  bool lhs_matching = may_group_on_lhs_non_contracting &&
                      (!may_group_on_rhs_non_contracting ||
                       lhs_non_contracting_partitions *
                               ShapeUtil::ByteSizeOf(rhs.hlo()->shape()) <
                           rhs_non_contracting_partitions *
                               ShapeUtil::ByteSizeOf(lhs.hlo()->shape()));
  if (may_group_on_lhs_non_contracting && may_group_on_rhs_non_contracting &&
      options.choose_faster_windowed_einsum_over_mem) {
    const DotDimensionIndexMapping indices_map = ComputeDimensionIndexMapping(
        dims_mapping, lhs.base_shape().rank(), rhs.base_shape().rank(),
        output_base_shape.rank());
    absl::optional<int64_t> lhs_matching_iterations;
    absl::optional<int64_t> rhs_matching_iterations;
    std::tie(lhs_matching_iterations, rhs_matching_iterations) =
        EstimateWindowedEinsumIterationsForNonContractingPartitioning(
            dims_mapping, lhs, rhs, output_base_shape, output_sharding, options,
            num_partitions, lhs_non_contracting_partitions,
            rhs_non_contracting_partitions, lhs_matching_partitions,
            rhs_matching_partitions, lhs_contracting_partitions,
            rhs_contracting_partitions, output_lhs_non_contracting_partitions,
            output_rhs_non_contracting_partitions, lhs_batch_partitions,
            rhs_batch_partitions, conv_window);
    if (lhs_matching_iterations && rhs_matching_iterations) {
      const int64_t lhs_all_gather_bytes =
          ShapeUtil::ByteSizeOf(lhs.hlo()->shape()) *
          rhs_non_contracting_partitions;
      const int64_t rhs_all_gather_bytes =
          ShapeUtil::ByteSizeOf(rhs.hlo()->shape()) *
          lhs_non_contracting_partitions;
      auto lhs_grouped =
          GetNonContractingPartitionGroupedShardingForMatchedOperand(
              true, lhs.sharding(), output_sharding,
              dims_mapping.lhs_non_contracting_dims);
      auto lhs_all_gather_subgroups = lhs_grouped.device_groups;
      auto rhs_grouped =
          GetNonContractingPartitionGroupedShardingForMatchedOperand(
              false, rhs.sharding(), output_sharding,
              dims_mapping.rhs_non_contracting_dims);
      auto rhs_all_gather_subgroups = rhs_grouped.device_groups;
      const double lhs_all_gather_time_in_ms =
          visitor->GetCommunicationTimeInMilliSec(
              lhs_all_gather_bytes,
              visitor->CreateReplicaGroups(lhs_all_gather_subgroups));
      const double rhs_all_gather_time_in_ms =
          visitor->GetCommunicationTimeInMilliSec(
              rhs_all_gather_bytes,
              visitor->CreateReplicaGroups(rhs_all_gather_subgroups));
      HloInstruction* compute_lhs = lhs.hlo();
      Shape lhs_original_shape = compute_lhs->shape();
      *compute_lhs->mutable_shape() =
          GetPerGroupBaseShape(lhs_grouped, lhs.base_shape());
      HloInstruction* compute_rhs = rhs.hlo();
      Shape rhs_original_shape = compute_rhs->shape();
      *compute_rhs->mutable_shape() =
          GetPerGroupBaseShape(rhs_grouped, rhs.base_shape());
      HloInstruction* dot =
          create_sharded_dot(compute_lhs, compute_rhs, b, conv_window)
              .ValueOrDie();
      const double computation_time_in_ms =
          visitor->GetComputationTimeInMilliSec(dot);
      *compute_lhs->mutable_shape() = lhs_original_shape;
      *compute_rhs->mutable_shape() = rhs_original_shape;
      VLOG(2) << "lhs: " << lhs.hlo()->ToString() << "\n"
              << "rhs: " << rhs.hlo()->ToString() << "\n"
              << "lhs_non_contracting_partitions: "
              << lhs_non_contracting_partitions
              << " rhs_non_contracting_partitions: "
              << rhs_non_contracting_partitions << "\n"
              << "lhs_matching_iterations: " << *lhs_matching_iterations
              << " rhs_matching_iterations: " << *rhs_matching_iterations
              << "\n"
              << "lhs_all_gather_bytes: " << lhs_all_gather_bytes
              << " rhs_all_gather_bytes: " << rhs_all_gather_bytes << "\n"
              << "lhs_all_gather_time_in_ms: " << lhs_all_gather_time_in_ms
              << " rhs_all_gather_time_in_ms: " << rhs_all_gather_time_in_ms
              << "\n"
              << "dot: " << dot->ToString() << "\n"
              << "computation_time_in_ms: " << computation_time_in_ms;
      if (computation_time_in_ms == 0.0 || lhs_all_gather_time_in_ms == 0.0 ||
          rhs_all_gather_time_in_ms == 0.0) {
        lhs_matching = *lhs_matching_iterations < *rhs_matching_iterations;
      } else if ((computation_time_in_ms <= lhs_all_gather_time_in_ms) &&
                 (computation_time_in_ms <= rhs_all_gather_time_in_ms)) {
        lhs_matching = lhs_all_gather_bytes / rhs_non_contracting_partitions >
                       rhs_all_gather_bytes / lhs_non_contracting_partitions;
      } else {
        lhs_matching = lhs_all_gather_time_in_ms > rhs_all_gather_time_in_ms;
      }
    }
  }
  return lhs_matching;
}
bool LhsIsBestMatchForNonContractingPartitioning ( const DotConvDimsMapping & dims_mapping , const PartitionedHlo & lhs , const PartitionedHlo & rhs , const Shape & output_base_shape , const HloSharding & output_sharding , const SpmdPartitionerOptions & options , int64_t num_partitions , int64_t lhs_non_contracting_partitions , int64_t rhs_non_contracting_partitions , int64_t lhs_matching_partitions , int64_t rhs_matching_partitions , int64_t lhs_contracting_partitions , int64_t rhs_contracting_partitions , int64_t output_lhs_non_contracting_partitions , int64_t output_rhs_non_contracting_partitions , int64_t lhs_batch_partitions , int64_t rhs_batch_partitions , SpmdBuilder * b , const Window & conv_window , const std :: function < StatusOr < HloInstruction * > ( HloInstruction * , HloInstruction * , SpmdBuilder * , const Window & conv_window ) > & create_sharded_dot , SpmdPartitioningVisitor * visitor ) { const bool may_group_on_lhs_non_contracting = lhs_non_contracting_partitions == output_lhs_non_contracting_partitions && lhs_non_contracting_partitions > 1 ; const bool may_group_on_rhs_non_contracting = rhs_non_contracting_partitions == output_rhs_non_contracting_partitions && rhs_non_contracting_partitions > 1 ; bool lhs_matching = may_group_on_lhs_non_contracting && ( ! may_group_on_rhs_non_contracting || lhs_non_contracting_partitions * ShapeUtil :: ByteSizeOf ( rhs . hlo ( ) -> shape ( ) ) < rhs_non_contracting_partitions * ShapeUtil :: ByteSizeOf ( lhs . hlo ( ) -> shape ( ) ) ) ; if ( may_group_on_lhs_non_contracting && may_group_on_rhs_non_contracting && options . choose_faster_windowed_einsum_over_mem ) { const DotDimensionIndexMapping indices_map = ComputeDimensionIndexMapping ( dims_mapping , lhs . base_shape ( ) . rank ( ) , rhs . base_shape ( ) . rank ( ) , output_base_shape . rank ( ) ) ; absl :: optional < int64_t > lhs_matching_iterations ; absl :: optional < int64_t > rhs_matching_iterations ; std :: tie ( lhs_matching_iterations , rhs_matching_iterations ) = EstimateWindowedEinsumIterationsForNonContractingPartitioning ( dims_mapping , lhs , rhs , output_base_shape , output_sharding , options , num_partitions , lhs_non_contracting_partitions , rhs_non_contracting_partitions , lhs_matching_partitions , rhs_matching_partitions , lhs_contracting_partitions , rhs_contracting_partitions , output_lhs_non_contracting_partitions , output_rhs_non_contracting_partitions , lhs_batch_partitions , rhs_batch_partitions , conv_window ) ; if ( lhs_matching_iterations && rhs_matching_iterations ) { const int64_t lhs_all_gather_bytes = ShapeUtil :: ByteSizeOf ( lhs . hlo ( ) -> shape ( ) ) * rhs_non_contracting_partitions ; const int64_t rhs_all_gather_bytes = ShapeUtil :: ByteSizeOf ( rhs . hlo ( ) -> shape ( ) ) * lhs_non_contracting_partitions ; auto lhs_grouped = GetNonContractingPartitionGroupedShardingForMatchedOperand ( true , lhs . sharding ( ) , output_sharding , dims_mapping . lhs_non_contracting_dims ) ; auto lhs_all_gather_subgroups = lhs_grouped . device_groups ; auto rhs_grouped = GetNonContractingPartitionGroupedShardingForMatchedOperand ( false , rhs . sharding ( ) , output_sharding , dims_mapping . rhs_non_contracting_dims ) ; auto rhs_all_gather_subgroups = rhs_grouped . device_groups ; const double lhs_all_gather_time_in_ms = visitor -> GetCommunicationTimeInMilliSec ( lhs_all_gather_bytes , visitor -> CreateReplicaGroups ( lhs_all_gather_subgroups ) ) ; const double rhs_all_gather_time_in_ms = visitor -> GetCommunicationTimeInMilliSec ( rhs_all_gather_bytes , visitor -> CreateReplicaGroups ( rhs_all_gather_subgroups ) ) ; HloInstruction * compute_lhs = lhs . hlo ( ) ; Shape lhs_original_shape = compute_lhs -> shape ( ) ; * compute_lhs -> mutable_shape ( ) = GetPerGroupBaseShape ( lhs_grouped , lhs . base_shape ( ) ) ; HloInstruction * compute_rhs = rhs . hlo ( ) ; Shape rhs_original_shape = compute_rhs -> shape ( ) ; * compute_rhs -> mutable_shape ( ) = GetPerGroupBaseShape ( rhs_grouped , rhs . base_shape ( ) ) ; HloInstruction * dot = create_sharded_dot ( compute_lhs , compute_rhs , b , conv_window ) . ValueOrDie ( ) ; const double computation_time_in_ms = visitor -> GetComputationTimeInMilliSec ( dot ) ; * compute_lhs -> mutable_shape ( ) = lhs_original_shape ; * compute_rhs -> mutable_shape ( ) = rhs_original_shape ; VLOG ( 2 ) << " " << lhs . hlo ( ) -> ToString ( ) << " \n " << " " << rhs . hlo ( ) -> ToString ( ) << " \n " << " " << lhs_non_contracting_partitions << " " << rhs_non_contracting_partitions << " \n " << " " << * lhs_matching_iterations << " " << * rhs_matching_iterations << " \n " << " " << lhs_all_gather_bytes << " " << rhs_all_gather_bytes << " \n " << " " << lhs_all_gather_time_in_ms << " " << rhs_all_gather_time_in_ms << " \n " << " " << dot -> ToString ( ) << " \n " << " " << computation_time_in_ms ; if ( computation_time_in_ms == 0.0 || lhs_all_gather_time_in_ms == 0.0 || rhs_all_gather_time_in_ms == 0.0 ) { lhs_matching = * lhs_matching_iterations < * rhs_matching_iterations ; } else if ( ( computation_time_in_ms <= lhs_all_gather_time_in_ms ) && ( computation_time_in_ms <= rhs_all_gather_time_in_ms ) ) { lhs_matching = lhs_all_gather_bytes / rhs_non_contracting_partitions > rhs_all_gather_bytes / lhs_non_contracting_partitions ; } else { lhs_matching = lhs_all_gather_time_in_ms > rhs_all_gather_time_in_ms ; } } } return lhs_matching ; }
Return if it would be better to match the LHS operand or RHS operand
of a dot for non-contracting partitioning.
Return if it would be better to match the LHS operand or RHS operand of a dot for non - contracting partitioning .
// If both match output non-contracting dimensions, choose the one which // will result in smaller replication of the other operand. // If both grouping are available and the option to choose faster windowed // einsums vs saving memory is enabled then try to determine which of the // operands will have more overlapping benefits for the windowed einsum // when matched (if a windowed einsum is gonna be generated at all). // 1) When computation is shorter than both all_gathers, we choose to overlap // with the smaller all_gather as it has potentially smaller extra // collective-permute overhead outside of the while loop; 2) Otherwise, we // choose the all_gather with longer runtime to overlap with. /*lhs_matching=*/ /*lhs_matching=*/