Skip to content

Commit d0ab56c

Browse files
vwbakertensorflower-gardener
authored andcommitted
[xla:gpu:gemm_fusion] Fix crash in HoistBitcast by copying users vector
Copy operand->users() to a local vector before iterating over it, as the loop body calls RemoveInstructionAndUnusedOperands which modifies the users list, leading to iterator invalidation and crash. PiperOrigin-RevId: 933680106
1 parent aa01fe9 commit d0ab56c

2 files changed

Lines changed: 25 additions & 2 deletions

File tree

third_party/xla/xla/backends/gpu/transforms/gemm_fusion.cc

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -918,7 +918,10 @@ absl::StatusOr<bool> FusionSearchSpace::HoistBitcast(HloInstruction* instr) {
918918
CopyElementType(operand->shape(), &new_shape);
919919

920920
*operand->mutable_shape() = new_shape;
921-
for (HloInstruction* user : operand->users()) {
921+
// Copy users to a new vector to avoid invalidating the iterator.
922+
std::vector<HloInstruction*> users(operand->users().begin(),
923+
operand->users().end());
924+
for (HloInstruction* user : users) {
922925
RETURN_IF_ERROR(user->ReplaceAllUsesWith(operand));
923926
RETURN_IF_ERROR(
924927
user->parent()->RemoveInstructionAndUnusedOperands(user));

third_party/xla/xla/backends/gpu/transforms/gemm_fusion_test.cc

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -476,7 +476,7 @@ ENTRY e {
476476
)");
477477
}
478478

479-
TEST_P(GemmFusionTestV2, DoNotHoistBitcastOverParameterWithMultipleUsers) {
479+
TEST_P(GemmFusionTestV2, DoNotHoistBitcastOverParameterWithNonBitcastUsers) {
480480
ASSERT_OK_AND_ASSIGN(auto module, ParseAndReturnVerifiedModule(R"(
481481
ENTRY e {
482482
p0 = bf16[32,12] parameter(0)
@@ -514,6 +514,26 @@ ENTRY e {
514514
GmockMatch(m::Fusion(m::Parameter(), m::Bitcast(m::Parameter()))));
515515
}
516516

517+
TEST_P(GemmFusionTestV2, HoistBitcastOverParameterWithMultipleBitcastUsers) {
518+
// Regression test for b/524943134.
519+
ASSERT_OK_AND_ASSIGN(auto module, ParseAndReturnVerifiedModule(R"(
520+
HloModule m
521+
522+
ENTRY e {
523+
p0 = f32[32,3] parameter(0)
524+
p1 = f32[8,28] parameter(1)
525+
b1 = f32[32,7] bitcast(p1)
526+
b2 = f32[32,7] bitcast(p1)
527+
add = f32[32,7] add(b1, b2)
528+
ROOT d = f32[3,7] dot(p0, add),
529+
lhs_contracting_dims={0}, rhs_contracting_dims={0}
530+
})"));
531+
ASSERT_THAT(GemmFusion(gpu_version_).Run(module.get()), IsOkAndHolds(true));
532+
EXPECT_THAT(
533+
module->entry_computation()->root_instruction(),
534+
GmockMatch(m::Fusion(m::Parameter(), m::Bitcast(m::Parameter()))));
535+
}
536+
517537
TEST_P(GemmFusionTestV2, HoistBitcastAcrossTransposeWithTrivialDimension) {
518538
ASSERT_OK_AND_ASSIGN(auto module, ParseAndReturnVerifiedModule(R"(
519539
ENTRY e {

0 commit comments

Comments
 (0)