Skip to content

Commit a493a0d

Browse files
seherellistensorflower-gardener
authored andcommitted
[XLA:SchedulingGroups][2/N] Do not combine all-reduces that have scheduling annotations.
PiperOrigin-RevId: 931253508
1 parent 3380ca0 commit a493a0d

3 files changed

Lines changed: 35 additions & 1 deletion

File tree

third_party/xla/xla/hlo/transforms/collectives/BUILD

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -370,6 +370,7 @@ cc_library(
370370
hdrs = ["all_reduce_combiner.h"],
371371
deps = [
372372
"//xla:shape_util",
373+
"//xla:side_effect_util",
373374
"//xla:status_macros",
374375
"//xla:xla_data_proto_cc",
375376
"//xla/hlo/ir:hlo",

third_party/xla/xla/hlo/transforms/collectives/all_reduce_combiner.cc

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,7 @@ limitations under the License.
4141
#include "xla/service/hlo_domain_map.h"
4242
#include "xla/shape.h"
4343
#include "xla/shape_util.h"
44+
#include "xla/side_effect_util.h"
4445
#include "xla/status_macros.h"
4546
#include "xla/xla_data.pb.h"
4647
#include "tsl/platform/errors.h"
@@ -169,7 +170,10 @@ absl::StatusOr<bool> AllReduceCombiner::RunWithKeyCombiner(
169170

170171
auto key_fn = [&domain_map, &combine_key](const HloInstruction* instruction)
171172
-> std::optional<AllReduceCombiner::GroupKey> {
172-
if (instruction->opcode() != HloOpcode::kAllReduce) {
173+
if (instruction->opcode() != HloOpcode::kAllReduce ||
174+
(instruction->has_frontend_attributes() &&
175+
instruction->frontend_attributes().map().contains(
176+
kXlaSchedulingGroupIdAttr))) {
173177
return std::nullopt;
174178
}
175179
return combine_key(instruction, *domain_map);

third_party/xla/xla/hlo/transforms/collectives/all_reduce_combiner_test.cc

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -542,5 +542,34 @@ TEST_F(AllReduceCombinerTest, PreservesMetadata) {
542542
op::GetTupleElement(combined_all_reduce, 1)));
543543
}
544544

545+
TEST_F(AllReduceCombinerTest, DoNotCombineOpWithSchedulingAnnotation) {
546+
absl::string_view hlo_text = R"(
547+
HloModule Module
548+
549+
%add (x: f32[], y: f32[]) -> f32[] {
550+
%x = f32[] parameter(0)
551+
%y = f32[] parameter(1)
552+
ROOT %add = f32[] add(f32[] %x, f32[] %y)
553+
}
554+
555+
ENTRY entry {
556+
%param.0 = f32[32] parameter(0)
557+
%param.1 = f32[32] parameter(1)
558+
%all-reduce.0 = f32[32] all-reduce(%param.0), replica_groups={},
559+
to_apply=%add, frontend_attributes={_scheduling_group_id="0"}
560+
%all-reduce.1 = f32[32] all-reduce(%param.1), replica_groups={},
561+
to_apply=%add
562+
ROOT tuple = (f32[32], f32[32]) tuple(%all-reduce.0, %all-reduce.1)
563+
}
564+
)";
565+
TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr<HloModule> module,
566+
ParseAndReturnVerifiedModule(hlo_text));
567+
AllReduceCombiner combine(1024 * 1024, kMaxCombineCount);
568+
EXPECT_THAT(combine.Run(module.get()), absl_testing::IsOkAndHolds(false));
569+
EXPECT_THAT(module->entry_computation()->root_instruction(),
570+
op::Tuple(op::AllReduce(op::Parameter(0)),
571+
op::AllReduce(op::Parameter(1))));
572+
}
573+
545574
} // namespace
546575
} // namespace xla

0 commit comments

Comments
 (0)