experiments/STT-tensorflow
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Remove TuplePointsToAnalysis::CanShareOperandBufferWithUser.

Browse Source 
The split-brain state of this complicated function and its counter
part in DataflowAnalysis, after a few syncs, got eventually out of
sync. This has become the source of many nasty bugs. This CL removes
the one in tuple points to analysis let the one in data flow analysis
become the source of truth.

PiperOrigin-RevId: 257455657
This commit is contained in:
Yunxing Dai 2019-07-10 12:14:17 -07:00 committed by TensorFlower Gardener
parent a7ecc1b1a0
commit 36e969ee0b
3 changed files with 1 additions and 530 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

143
tensorflow/compiler/xla/service/tuple_points_to_analysis.cc
View File

@ -708,147 +708,4 @@ bool TuplePointsToAnalysis::HasUniqueFusedUseOfOperandAt(
fused_param_uses[0].first == fusion->fused_expression_root() &&
fused_param_uses[0].second == use_operand_index;
}
// User and operand can share buffers iff both instructions emit the same shape
// and layout, and 'user' meets one of the following qualifications:
//
// (1) Is element-wise. Or...
// (2) Is a loop fusion instruction where the only use of 'operand' at 'index'
// in the set 'user.fused_instructions' is a DynamicUpdateSlice fused root
// at operand 0. Or...
// (3) Is a kDot -> kAdd output fusion instruction where the only use of
// 'operand' at 'index' in the set 'user.fused_instructions' is a kAdd fused
// root at operand 0 or 1. Or...
// (4) The 'user' of 'operand' is DynamicUpdateSlice or While at operand index
// 0.
// (5) The 'user' of 'operand' is Sort, and it is the only user.
// (6) The 'user' of 'operand' is TriangularSolve, it is the second operand,
// and it is the only user.
//
// (2) and (3) can only be determined if points-to analysis is available.
bool TuplePointsToAnalysis::CanShareOperandBufferWithUser(
HloInstruction* operand, const ShapeIndex& operand_index,
HloInstruction* user, const ShapeIndex& user_index) const {
CHECK(user->IsUserOf(operand))
<< "user: " << user->ToString() << " operand: " << operand->ToString();
const Shape& operand_subshape =
ShapeUtil::GetSubshape(operand->shape(), operand_index);
const Shape& user_subshape =
ShapeUtil::GetSubshape(user->shape(), user_index);
// Check that operand and user emit the same shape and layout.
if (!ShapeUtil::Equal(operand_subshape, user_subshape)) {
return false;
}
if (user->opcode() == HloOpcode::kFusion) {
if (user->IsLoopFusion() || user->IsInputFusion()) {
if (user->fused_expression_root()->opcode() ==
HloOpcode::kDynamicUpdateSlice) {
// Loop fusion with kDynamicUpdateSlice fused root.
//
// Returns true iff there is exactly one use of 'operand' at shape index
// 'operand_index', and this singleton use is the fused root at operand
// index 0.
return HasUniqueFusedUseOfOperandAt(operand, operand_index, user, 0);
} else {
HloInstruction* fusion_param =
user->fused_parameter(user->operand_index(operand));
return HloDataflowAnalysis::AreTransitiveUsesElementwiseOrTuple(
fusion_param);
}
} else if (user->IsOutputFusion() &&
user->fused_expression_root()->opcode() == HloOpcode::kAdd) {
// Output fusion with kAdd fused root.
// Check if one operand of kAdd fused root is kDot or kConvolution.
auto* add = user->fused_expression_root();
auto add_operand_it =
absl::c_find_if(add->operands(), [&](HloInstruction* operand) {
return operand->opcode() == HloOpcode::kConvolution ||
operand->opcode() == HloOpcode::kDot;
});
if (add_operand_it == add->operands().end()) {
return false;
}
auto* matched_add_operand = *add_operand_it;
// Calculate operand index of 'add' operand which was not matched above.
const int64 other_add_operand_index =
matched_add_operand == add->operand(0) ? 1 : 0;
// Returns true iff there is exactly one use of 'operand' at shape index
// 'operand_index', and this singleton use is the fused root (at operand
// index 'other_add_operand_index').
return HasUniqueFusedUseOfOperandAt(operand, operand_index, user,
other_add_operand_index);
} else if (user->IsCustomFusion()) {
std::vector<int64> operand_indices = user->OperandIndices(operand);
return operand_indices.size() == 1 && operand_indices[0] == 0 &&
absl::c_any_of(
user->fused_instructions_computation()->instructions(),
[](const HloInstruction* hlo) {
return hlo->opcode() == HloOpcode::kScatter;
});
}
}
if (user->opcode() == HloOpcode::kDynamicUpdateSlice ||
user->opcode() == HloOpcode::kScatter ||
user->opcode() == HloOpcode::kWhile) {
// We eliminated other users in BufferLiveness::live_range_strictly_before,
// so here we just need to check that the use is at operand index 0.
std::vector<int64> operand_indices = user->OperandIndices(operand);
return operand_indices.size() == 1 && operand_indices[0] == 0;
}
if (user->opcode() == HloOpcode::kSort) {
// Only valid if there are no other users.
if (operand->users().size() != 1) {
return false;
}
// If we only sort keys, the output of sort is not a tuple, so we can always
// share the buffer.
if (user->operand_count() == 1) {
return true;
}
CHECK(!user_index.empty());
// Only share with the right tuple element buffer.
std::vector<int64> operand_indices = user->OperandIndices(operand);
return operand_indices.size() == 1 && user_index[0] == operand_indices[0];
}
if (user->opcode() == HloOpcode::kTriangularSolve) {
// Only valid if there are no other users.
if (operand->users().size() != 1) {
return false;
}
std::vector<int64> operand_indices = user->OperandIndices(operand);
return operand_indices.size() == 1 && operand_indices[0] == 1;
}
if (user->opcode() == HloOpcode::kCall) {
// TODO(b/62548313): Remove when buffer assignment is module scoped and
// does not assign buffers to calls.
// Find called computation parameter associated with 'operand'.
const std::vector<int64> operand_indices = user->OperandIndices(operand);
if (operand_indices.size() > 1) {
return false;
}
CHECK_EQ(1, operand_indices.size());
auto* param = user->to_apply()->parameter_instruction(operand_indices[0]);
// Get all uses of 'operand' at 'index' in called computation.
auto param_uses = GetAllUsesOfInstructionAtIndex(param, operand_index);
// Return true iff:
// *) There exists exactly one use of 'operand' in called computation.
// *) The unique use is by the root instruction of called computation.
// (Note: we check the root of the called computation, because the
// root result buffer is required to alias with the Call result buffer).
// *) The root instruction of the called computation is element-wise on
// 'operand'.
auto* callee_root = user->to_apply()->root_instruction();
return param_uses.size() == 1 && param_uses[0].first == callee_root &&
callee_root->IsElementwiseOnOperand(param_uses[0].second);
}
// Loop fusions that contain transposing copies won't reach here as they have
// different layouts, which fails the check in the beginning of this function.
//
// Multi-output fusion will fail the check here as tuples are not considered
// an elementwise operation.
return user->IsElementwiseOnOperand(user->operand_index(operand));
}
} // namespace xla

10
tensorflow/compiler/xla/service/tuple_points_to_analysis.h
View File

@ -17,6 +17,7 @@ limitations under the License.
#define TENSORFLOW_COMPILER_XLA_SERVICE_TUPLE_POINTS_TO_ANALYSIS_H_
#include <stddef.h>
#include <iosfwd>
#include <memory>
#include <set>
@ -265,15 +266,6 @@ class TuplePointsToAnalysis : public DfsHloVisitorWithDefault {
const ShapeIndex& index,
const HloInstruction* user) const;
// Returns true if 'user' (at 'user_index') can share a buffer with its
// operand 'operand' (at 'operand_index'). Returns false otherwise.
//
// REQUIRES: 'operand' is an operand of 'user'.
bool CanShareOperandBufferWithUser(HloInstruction* operand,
const ShapeIndex& operand_index,
HloInstruction* user,
const ShapeIndex& user_index) const;
private:
explicit TuplePointsToAnalysis(
const HloModule* module,

378
tensorflow/compiler/xla/service/tuple_points_to_analysis_test.cc
View File

@ -928,383 +928,5 @@ TEST_F(DoesNotUseOperandBufferTest, FusedDynamicUpdateSlice) {
EXPECT_FALSE(
points_to_analysis_->DoesNotUseOperandBuffer(tuple, {1}, fusion));
}
class CanShareOperandBufferWithUserTest : public PointsToAnalysisTestBase {};
TEST_F(CanShareOperandBufferWithUserTest, ElementWiseSameShape) {
auto builder = HloComputation::Builder(TestName());
Shape shape = ShapeUtil::MakeShape(F32, {8});
auto param = builder.AddInstruction(
HloInstruction::CreateParameter(0, shape, "param"));
auto exp = builder.AddInstruction(
HloInstruction::CreateUnary(shape, HloOpcode::kExp, param));
auto log = builder.AddInstruction(
HloInstruction::CreateUnary(shape, HloOpcode::kLog, exp));
BuildModuleAndRunAnalysis(builder.Build());
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(param, {}, exp, {}));
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(exp, {}, log, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, ElementWiseDifferentShape) {
auto builder = HloComputation::Builder(TestName());
Shape in_shape = ShapeUtil::MakeShape(F32, {8});
Shape out_shape = ShapeUtil::MakeShape(PRED, {8});
auto param0 = builder.AddInstruction(
HloInstruction::CreateParameter(0, in_shape, "param0"));
auto param1 = builder.AddInstruction(
HloInstruction::CreateParameter(1, in_shape, "param1"));
auto result = builder.AddInstruction(HloInstruction::CreateCompare(
out_shape, param0, param1, ComparisonDirection::kEq));
BuildModuleAndRunAnalysis(builder.Build());
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(param0, {},
result, {}));
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(param1, {},
result, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, CopyShares) {
auto builder = HloComputation::Builder(TestName());
Shape shape = ShapeUtil::MakeShape(F32, {8});
auto param = builder.AddInstruction(
HloInstruction::CreateParameter(0, shape, "param"));
auto exp = builder.AddInstruction(
HloInstruction::CreateUnary(shape, HloOpcode::kExp, param));
auto copy = builder.AddInstruction(
HloInstruction::CreateUnary(shape, HloOpcode::kCopy, exp));
BuildModuleAndRunAnalysis(builder.Build());
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(param, {}, exp, {}));
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(exp, {}, copy, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, FusedDynamicUpdateSlice) {
auto builder = HloComputation::Builder(TestName());
Shape data_shape = ShapeUtil::MakeShape(F32, {8});
auto tuple = builder.AddInstruction(HloInstruction::CreateParameter(
0, ShapeUtil::MakeTupleShape({data_shape, data_shape}), "tuple"));
auto gte0 = builder.AddInstruction(
HloInstruction::CreateGetTupleElement(data_shape, tuple, 0));
auto gte1 = builder.AddInstruction(
HloInstruction::CreateGetTupleElement(data_shape, tuple, 1));
// Create a DynamicUpdateSlice instruction of tuple element 1.
auto starts = builder.AddInstruction(
HloInstruction::CreateConstant(LiteralUtil::CreateR0<int32>(2)));
auto update = builder.AddInstruction(HloInstruction::CreateConstant(
LiteralUtil::CreateR1<float>({2.f, 2.f, 2.f})));
auto dynamic_update_slice =
builder.AddInstruction(HloInstruction::CreateDynamicUpdateSlice(
data_shape, gte1, update, {starts}));
builder.AddInstruction(
HloInstruction::CreateTuple({gte0, dynamic_update_slice}));
BuildModule(builder.Build());
auto fusion = computation_->CreateFusionInstruction(
{dynamic_update_slice, starts, update, gte1},
HloInstruction::FusionKind::kLoop);
RunAnalysis();
// The fusion instruction can share with tuple element 1.
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(tuple, {0},
fusion, {}));
EXPECT_TRUE(points_to_analysis_->CanShareOperandBufferWithUser(tuple, {1},
fusion, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, DynamicUpdateSliceCanShare) {
auto builder = HloComputation::Builder(TestName());
Shape data_shape = ShapeUtil::MakeShape(F32, {8});
Shape update_shape = ShapeUtil::MakeShape(F32, {4});
Shape starts_shape = ShapeUtil::MakeShape(S32, {});
auto data = builder.AddInstruction(
HloInstruction::CreateParameter(0, data_shape, "data"));
auto update = builder.AddInstruction(
HloInstruction::CreateParameter(1, update_shape, "update"));
auto starts = builder.AddInstruction(
HloInstruction::CreateParameter(2, starts_shape, "starts"));
auto dus = builder.AddInstruction(HloInstruction::CreateDynamicUpdateSlice(
data_shape, data, update, {starts}));
BuildModuleAndRunAnalysis(builder.Build());
// The DynamicUpdateSlice instruction can share with the data operand, but not
// with update or starts.
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(data, {}, dus, {}));
EXPECT_FALSE(
points_to_analysis_->CanShareOperandBufferWithUser(update, {}, dus, {}));
EXPECT_FALSE(
points_to_analysis_->CanShareOperandBufferWithUser(starts, {}, dus, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, ScatterCanShare) {
const char* hlo_text = R"(
HloModule TensorFlowScatterV1
update_s32 (lhs: s32[], rhs: s32[]) -> s32[] {
lhs = s32[] parameter(0)
ROOT rhs = s32[] parameter(1)
}
ENTRY main {
operand = s32[3,3] parameter(0)
indices = s32[2] parameter(1)
updates = s32[2,3] parameter(2)
ROOT scatter = s32[3,3] scatter(operand, indices, updates),
to_apply=update_s32,
update_window_dims={1},
inserted_window_dims={0},
scatter_dims_to_operand_dims={0},
index_vector_dim=1
}
)";
TF_ASSERT_OK_AND_ASSIGN(module_, ParseAndReturnUnverifiedModule(hlo_text));
computation_ = module_->entry_computation();
RunAnalysis();
HloInstruction* operand_param = computation_->parameter_instruction(0);
HloInstruction* indices_param = computation_->parameter_instruction(1);
HloInstruction* updates_param = computation_->parameter_instruction(2);
HloInstruction* scatter = computation_->root_instruction();
EXPECT_TRUE(points_to_analysis_->CanShareOperandBufferWithUser(
operand_param, {}, scatter, {}));
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(
indices_param, {}, scatter, {}));
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(
updates_param, {}, scatter, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, SortCanShare) {
auto builder = HloComputation::Builder(TestName());
module_ = CreateNewVerifiedModule();
Shape keys_shape = ShapeUtil::MakeShape(F32, {8});
auto keys = builder.AddInstruction(
HloInstruction::CreateParameter(0, keys_shape, "keys"));
TF_ASSERT_OK_AND_ASSIGN(
auto* sort, MakeSortHlo(keys_shape, {keys}, 0, /*is_stable=*/false,
&builder, module_.get()));
computation_ = module_->AddEntryComputation(builder.Build());
RunAnalysis();
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(keys, {}, sort, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, SortCanShareWithTupleUser) {
auto builder = HloComputation::Builder(TestName());
module_ = CreateNewVerifiedModule();
Shape keys_shape = ShapeUtil::MakeShape(F32, {8});
Shape values_shape = ShapeUtil::MakeShape(F32, {8});
auto keys = builder.AddInstruction(
HloInstruction::CreateParameter(0, keys_shape, "keys"));
auto values = builder.AddInstruction(
HloInstruction::CreateParameter(1, values_shape, "values"));
TF_ASSERT_OK_AND_ASSIGN(
auto* sort,
MakeSortHlo(ShapeUtil::MakeTupleShape({keys_shape, values_shape}),
{keys, values}, 0, /*is_stable=*/false, &builder,
module_.get()));
computation_ = module_->AddEntryComputation(builder.Build());
RunAnalysis();
// The buffer for the keys can be shared with the first tuple entry.
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(keys, {}, sort, {0}));
// The buffer for the values can be shared with the second tuple entry.
EXPECT_TRUE(points_to_analysis_->CanShareOperandBufferWithUser(values, {},
sort, {1}));
// Verify that the buffers are not shared with the "wrong" tuple entry.
EXPECT_FALSE(
points_to_analysis_->CanShareOperandBufferWithUser(keys, {}, sort, {1}));
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(values, {},
sort, {0}));
}
TEST_F(CanShareOperandBufferWithUserTest, FusedDotAdd) {
auto builder = HloComputation::Builder(TestName());
Shape data_shape = ShapeUtil::MakeShape(F32, {2, 2});
auto a = builder.AddInstruction(HloInstruction::CreateConstant(
LiteralUtil::CreateR2<float>({{1.0, 0.0}, {0.0, 1.0}})));
auto b = builder.AddInstruction(HloInstruction::CreateConstant(
LiteralUtil::CreateR2<float>({{2.0, 2.0}, {2.0, 2.0}})));
DotDimensionNumbers dot_dnums;
dot_dnums.add_lhs_contracting_dimensions(1);
dot_dnums.add_rhs_contracting_dimensions(0);
PrecisionConfig precision_config;
precision_config.mutable_operand_precision()->Resize(
/*new_size=*/2, PrecisionConfig::DEFAULT);
auto dot = builder.AddInstruction(
HloInstruction::CreateDot(data_shape, a, b, dot_dnums, precision_config));
auto one = builder.AddInstruction(
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(1.0)));
auto add_operand = builder.AddInstruction(
HloInstruction::CreateBroadcast(data_shape, one, {1}));
auto add = builder.AddInstruction(HloInstruction::CreateBinary(
data_shape, HloOpcode::kAdd, dot, add_operand));
BuildModule(builder.Build());
auto fusion = computation_->CreateFusionInstruction(
{add, dot}, HloInstruction::FusionKind::kOutput);
RunAnalysis();
// Output fused dot add should be able to share buffer with 'add_operand'.
EXPECT_TRUE(points_to_analysis_->CanShareOperandBufferWithUser(
add_operand, {}, fusion, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, OutputFusionCantAliasOperandBuffer) {
auto builder = HloComputation::Builder(TestName());
Shape data_shape = ShapeUtil::MakeShape(F32, {2, 2});
auto one = builder.AddInstruction(
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(1.0)));
auto operand = builder.AddInstruction(
HloInstruction::CreateBroadcast(data_shape, one, {1}));
auto reverse = builder.AddInstruction(
HloInstruction::CreateReverse(data_shape, operand, {0, 1}));
auto two = builder.AddInstruction(HloInstruction::CreateConstant(
LiteralUtil::CreateR2<float>({{2.0, 2.0}, {2.0, 2.0}})));
auto add = builder.AddInstruction(
HloInstruction::CreateBinary(data_shape, HloOpcode::kAdd, reverse, two));
BuildModule(builder.Build());
auto fusion = computation_->CreateFusionInstruction(
{add, two, reverse}, HloInstruction::FusionKind::kOutput);
RunAnalysis();
// Output fused operand->reverse->add cannot alias operand buffer 'operand'.
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(operand, {},
fusion, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, WhileCanShare) {
Shape data_shape = ShapeUtil::MakeShape(F32, {8});
auto make_cond = [&data_shape]() {
auto builder = HloComputation::Builder(TestName() + ".Cond");
auto data = builder.AddInstruction(
HloInstruction::CreateParameter(0, data_shape, "data"));
builder.AddInstruction(HloInstruction::CreateCompare(
ShapeUtil::MakeShape(PRED, {}), data, data, ComparisonDirection::kEq));
return builder.Build();
};
auto make_body = [&data_shape]() {
auto builder = HloComputation::Builder(TestName() + ".Body");
auto data = builder.AddInstruction(
HloInstruction::CreateParameter(0, data_shape, "data"));
builder.AddInstruction(
HloInstruction::CreateBinary(data_shape, HloOpcode::kAdd, data, data));
return builder.Build();
};
module_ = CreateNewUnverifiedModule();
HloComputation* cond_computation =
module_->AddEmbeddedComputation(make_cond());
HloComputation* body_computation =
module_->AddEmbeddedComputation(make_body());
auto builder = HloComputation::Builder(TestName());
auto data = builder.AddInstruction(
HloInstruction::CreateParameter(0, data_shape, "data"));
auto whil = builder.AddInstruction(HloInstruction::CreateWhile(
data_shape, cond_computation, body_computation, data));
computation_ = module_->AddEntryComputation(builder.Build());
RunAnalysis();
// The While instruction can share with the data operand.
EXPECT_TRUE(
points_to_analysis_->CanShareOperandBufferWithUser(data, {}, whil, {}));
}
// Tests that Call can alias operand buffer if the only use of the operand
// in the called computation is an elementwise instruction.
TEST_F(CanShareOperandBufferWithUserTest, CallToComputationWithFusionRoot) {
Shape shape = ShapeUtil::MakeShape(F32, {8});
// Build sub-computation with fusion root.
auto sub_builder = HloComputation::Builder(TestName() + "_sub");
auto sub_param = sub_builder.AddInstruction(
HloInstruction::CreateParameter(0, shape, "sub_param"));
auto one = sub_builder.AddInstruction(
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(1.0)));
auto ones = sub_builder.AddInstruction(
HloInstruction::CreateBroadcast(shape, one, {1}));
auto add = sub_builder.AddInstruction(
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, sub_param, ones));
module_ = CreateNewUnverifiedModule();
auto sub_computation = module_->AddEmbeddedComputation(sub_builder.Build());
sub_computation->CreateFusionInstruction({add, ones},
HloInstruction::FusionKind::kLoop);
// Build entry-computation with kCall which calls 'sub_computation'.
auto builder = HloComputation::Builder(TestName());
auto param = builder.AddInstruction(
HloInstruction::CreateParameter(0, shape, "param"));
auto reverse =
builder.AddInstruction(HloInstruction::CreateReverse(shape, param, {0}));
auto call = builder.AddInstruction(
HloInstruction::CreateCall(shape, {reverse}, sub_computation));
computation_ = module_->AddEntryComputation(builder.Build());
RunAnalysis();
EXPECT_TRUE(points_to_analysis_->CanShareOperandBufferWithUser(reverse, {},
call, {}));
}
TEST_F(CanShareOperandBufferWithUserTest, LoopFusionWithElementwiseOperand) {
Shape full_shape = ShapeUtil::MakeShape(F32, {16, 32});
Shape broadcast_shape = ShapeUtil::MakeShape(F32, {16});
auto builder = HloComputation::Builder(TestName() + "_fusion");
auto param0 = builder.AddInstruction(
HloInstruction::CreateParameter(0, full_shape, "full"));
auto param1 = builder.AddInstruction(
HloInstruction::CreateParameter(1, broadcast_shape, "small"));
auto broadcast = builder.AddInstruction(
HloInstruction::CreateBroadcast(full_shape, param1, {0}));
auto add = builder.AddInstruction(HloInstruction::CreateBinary(
full_shape, HloOpcode::kAdd, param0, broadcast));
BuildModule(builder.Build());
auto fusion = computation_->CreateFusionInstruction(
{add, broadcast}, HloInstruction::FusionKind::kLoop);
RunAnalysis();
EXPECT_TRUE(points_to_analysis_->CanShareOperandBufferWithUser(param0, {},
fusion, {}));
EXPECT_FALSE(points_to_analysis_->CanShareOperandBufferWithUser(param1, {},
fusion, {}));
}
} // namespace
} // namespace xla