[XLA] Rename HloTestBase::CreateNewModule -> CreateNewUnverifiedModule. (Reland)
Preparation for merging HloTestBase and HloVerifiedTestBase. PiperOrigin-RevId: 221018020
This commit is contained in:
parent
748e703d9d
commit
a87acdb86d
@ -2906,7 +2906,7 @@ TEST_F(AlgebraicSimplifierTest, ConvertConvToMatmul) {
|
||||
/*feature_group_count=*/1, window, dnums, DefaultPrecisionConfig(2)));
|
||||
|
||||
// TODO(b/80488902): verify this module.
|
||||
auto module = HloTestBase::CreateNewModule();
|
||||
auto module = HloTestBase::CreateNewUnverifiedModule();
|
||||
auto* computation = module->AddEntryComputation(b.Build());
|
||||
|
||||
AlgebraicSimplifier simplifier(/*is_layout_sensitive=*/true,
|
||||
@ -3084,7 +3084,7 @@ TEST_F(AlgebraicSimplifierTest, ScalarBroadcastToTransposeReshape) {
|
||||
// Test that ReduceWindow(Pad(op, x), y) can simplify to ReduceWindow(op, x).
|
||||
TEST_F(AlgebraicSimplifierTest, FoldPadIntoReduceWindow) {
|
||||
// TODO(b/80488902): verify this module.
|
||||
auto module = HloTestBase::CreateNewModule();
|
||||
auto module = HloTestBase::CreateNewUnverifiedModule();
|
||||
HloComputation::Builder builder(TestName());
|
||||
|
||||
// Create operand to the pad.
|
||||
@ -3166,7 +3166,7 @@ TEST_F(AlgebraicSimplifierTest, FoldPadIntoReduceWindow) {
|
||||
// ReduceWindow(Convert(op), x).
|
||||
TEST_F(AlgebraicSimplifierTest, FoldConvertedPadIntoReduceWindow) {
|
||||
// TODO(b/80488902): verify this module.
|
||||
auto module = HloTestBase::CreateNewModule();
|
||||
auto module = HloTestBase::CreateNewUnverifiedModule();
|
||||
HloComputation::Builder builder(TestName());
|
||||
|
||||
// Create operand to the pad.
|
||||
|
||||
@ -117,7 +117,7 @@ TEST_F(BufferLivenessTest, ElementwiseChain) {
|
||||
auto log = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(vec_, HloOpcode::kLog, exp));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto liveness =
|
||||
@ -164,7 +164,7 @@ TEST_F(BufferLivenessTest, MultipleEntryParameters_Sequential) {
|
||||
auto add = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(vec_, HloOpcode::kAdd, negate, exp));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* entry = module->AddEntryComputation(builder.Build());
|
||||
|
||||
HloSchedule schedule(module.get());
|
||||
@ -213,7 +213,7 @@ TEST_F(BufferLivenessTest, NonElementwiseOperand) {
|
||||
auto reverse =
|
||||
builder.AddInstruction(HloInstruction::CreateReverse(vec_, negate, {0}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto liveness =
|
||||
@ -247,7 +247,7 @@ TEST_F(BufferLivenessTest, OverlappedBuffers) {
|
||||
auto add = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(vec_, HloOpcode::kAdd, negate, exp));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto liveness =
|
||||
@ -289,7 +289,7 @@ TEST_F(BufferLivenessTest, OverlappedBuffersSequentialOrder) {
|
||||
auto add = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(vec_, HloOpcode::kAdd, negate, exp));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
HloSchedule schedule(module.get());
|
||||
@ -336,7 +336,7 @@ TEST_F(BufferLivenessTest, RootInstructionIsNotLastInSequentialOrder) {
|
||||
HloInstruction::CreateSend(recv_done, token, /*channel_id=*/1));
|
||||
auto send_done = builder.AddInstruction(HloInstruction::CreateSendDone(send));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build(add));
|
||||
|
||||
HloSchedule schedule(module.get());
|
||||
@ -373,7 +373,7 @@ TEST_F(BufferLivenessTest, TupleLiveOut) {
|
||||
auto outer_tuple =
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({inner_tuple, exp}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto liveness =
|
||||
@ -393,7 +393,7 @@ TEST_F(BufferLivenessTest, TupleLiveOut) {
|
||||
|
||||
TEST_F(BufferLivenessTest, EmbeddedComputation) {
|
||||
// Test MaybeLiveOut and MayInterfere for embedded computation.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
|
||||
auto embedded_builder = HloComputation::Builder(TestName() + "_embedded");
|
||||
auto embedded_param = embedded_builder.AddInstruction(
|
||||
@ -450,7 +450,7 @@ TEST_F(BufferLivenessTest, TupleConstantLiveOut) {
|
||||
builder.AddInstruction(HloInstruction::CreateGetTupleElement(
|
||||
inner_tuple0.shape(), tuple_constant, 0));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto liveness =
|
||||
@ -514,7 +514,7 @@ TEST_F(BufferLivenessTest, IndependentTupleElements) {
|
||||
auto tuple_root =
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({add0, add1}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(BuildDummyComputation());
|
||||
module->AddEmbeddedComputation(builder.Build());
|
||||
|
||||
@ -576,7 +576,7 @@ TEST_F(BufferLivenessTest, DependentTupleElements) {
|
||||
auto tuple_root =
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({add0, add1}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(BuildDummyComputation());
|
||||
module->AddEmbeddedComputation(builder.Build());
|
||||
|
||||
@ -646,7 +646,7 @@ class FusedDynamicUpdateSliceLivenessTest : public BufferLivenessTest {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateTuple({gte0, dynamic_update_slice}));
|
||||
// Build module and get reference to entry computation.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
auto* computation = module->entry_computation();
|
||||
// Create fusion instruction based on number of tuple element 1 users.
|
||||
@ -802,7 +802,7 @@ class DynamicUpdateSliceLivenessTest : public BufferLivenessTest {
|
||||
auto tuple_root = builder.AddInstruction(
|
||||
HloInstruction::CreateTuple({gte0, dynamic_update_slice}));
|
||||
// Build module and get reference to entry computation.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(BuildDummyComputation());
|
||||
module->AddEmbeddedComputation(builder.Build());
|
||||
// Run BufferLiveness on 'module'.
|
||||
|
||||
@ -94,7 +94,7 @@ TEST_F(CopyInsertionTest, SingleParameter) {
|
||||
|
||||
EXPECT_THAT(x->users(), UnorderedElementsAre(tuple));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
InsertCopies(module.get());
|
||||
@ -114,7 +114,7 @@ TEST_F(CopyInsertionTest, SingleConstant) {
|
||||
|
||||
EXPECT_THAT(constant->users(), UnorderedElementsAre(tuple));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
InsertCopies(module.get());
|
||||
@ -127,7 +127,7 @@ TEST_F(CopyInsertionTest, SingleConstant) {
|
||||
TEST_F(CopyInsertionTest, ExistingCopiesNotRemoved) {
|
||||
// Verify that kCopy instructions which change layout and exist before
|
||||
// copy-insertion remain in the graph after copy-insertion.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
HloInstruction* constant =
|
||||
@ -181,7 +181,7 @@ TEST_F(CopyInsertionTest, MultipleConstantsAndParameters) {
|
||||
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({constant2, x, add}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
InsertCopies(module.get());
|
||||
@ -217,7 +217,7 @@ TEST_F(CopyInsertionTest, AmbiguousPointsToSet) {
|
||||
EXPECT_THAT(constant2->users(), UnorderedElementsAre(tuple1, tuple2));
|
||||
EXPECT_THAT(constant3->users(), UnorderedElementsAre(tuple2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
HloInstruction* old_root = module->entry_computation()->root_instruction();
|
||||
@ -238,7 +238,7 @@ TEST_F(CopyInsertionTest, BitcastParameter) {
|
||||
HloInstruction* bitcast = builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
ShapeUtil::MakeShape(F32, {2, 2}), HloOpcode::kBitcast, x));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_THAT(x->users(), UnorderedElementsAre(bitcast));
|
||||
@ -261,7 +261,7 @@ TEST_F(CopyInsertionTest, BitcastConstant) {
|
||||
HloInstruction* bitcast = builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
ShapeUtil::MakeShape(F32, {2, 2}), HloOpcode::kBitcast, constant));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_THAT(constant->users(), UnorderedElementsAre(bitcast));
|
||||
@ -283,7 +283,7 @@ TEST_F(CopyInsertionTest, BitcastTupleElementParameter) {
|
||||
ShapeUtil::MakeShape(F32, {2, 2}), HloOpcode::kBitcast, x));
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({bitcast}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_THAT(x->users(), UnorderedElementsAre(bitcast));
|
||||
@ -310,7 +310,7 @@ TEST_F(CopyInsertionTest, NestedTupleParameter) {
|
||||
ShapeUtil::MakeShape(F32, {42})}),
|
||||
"param0"));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_EQ(HloOpcode::kParameter,
|
||||
@ -351,7 +351,7 @@ TEST_F(CopyInsertionTest, ElementOfNestedTupleParameter) {
|
||||
auto gte = builder.AddInstruction(HloInstruction::CreateGetTupleElement(
|
||||
ShapeUtil::GetSubshape(param->shape(), {0}), param, 0));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_EQ(gte, module->entry_computation()->root_instruction());
|
||||
@ -388,7 +388,7 @@ TEST_F(CopyInsertionTest, AmbiguousTopLevelRoot) {
|
||||
builder.AddInstruction(HloInstruction::CreateGetTupleElement(
|
||||
ShapeUtil::GetSubshape(select->shape(), {0}), select, 0));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_EQ(gte, module->entry_computation()->root_instruction());
|
||||
@ -403,7 +403,7 @@ TEST_F(CopyInsertionTest, AmbiguousTopLevelRoot) {
|
||||
|
||||
class WhileCopyInsertionTest : public CopyInsertionTest {
|
||||
protected:
|
||||
WhileCopyInsertionTest() : module_(CreateNewModule()) {}
|
||||
WhileCopyInsertionTest() : module_(CreateNewUnverifiedModule()) {}
|
||||
|
||||
// Builds a While condition computation which reads the induction variable
|
||||
// from the tuple parameter, and returns a predicate indicating whether this
|
||||
@ -1295,7 +1295,7 @@ TEST_F(WhileCopyInsertionTest, InitPointsToNonDistinctUsedByTwoWhileLoops) {
|
||||
TEST_F(CopyInsertionTest, SwizzlingWhile) {
|
||||
// Test a while instruction with a body which permutes its tuple parameter
|
||||
// elements.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape loop_state_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1362,7 +1362,7 @@ TEST_F(CopyInsertionTest, CrossingParameters) {
|
||||
// | / \ |
|
||||
// | / \|
|
||||
// (p1 , p0)
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape tuple_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1395,7 +1395,7 @@ TEST_F(CopyInsertionTest, ParametersAliasing) {
|
||||
// | |
|
||||
// | |
|
||||
// (p0 , p1)
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape tuple_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1428,7 +1428,7 @@ TEST_F(CopyInsertionTest, ParameterWithNoAliasing) {
|
||||
// | |
|
||||
// | |
|
||||
// (p0 , p1)
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape tuple_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1461,7 +1461,7 @@ TEST_F(CopyInsertionTest, ParameterWithPartialAliasing) {
|
||||
// | |
|
||||
// | |
|
||||
// (p0 , p1)
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape tuple_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1496,7 +1496,7 @@ TEST_F(CopyInsertionTest, ParameterAndParallelOpsWithPartialAliasing) {
|
||||
// | | |
|
||||
// | | |
|
||||
// +-- (p0 , p1)
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape tuple_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1534,7 +1534,7 @@ TEST_F(CopyInsertionTest, ParameterAndOpsWithPartialAliasing) {
|
||||
// | Add----+
|
||||
// | | |
|
||||
// +-- (p0 , p1)
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape tuple_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1569,7 +1569,7 @@ TEST_F(CopyInsertionTest, SwizzlingWhileWithOneOp) {
|
||||
// the operation (instruction) on the element makes the live range of the
|
||||
// respective input and output elements different than if the instruction were
|
||||
// not there (as in the SwizzlingWhile test above).
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape loop_state_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1632,7 +1632,7 @@ TEST_F(CopyInsertionTest, SwizzlingWhileSharedInput) {
|
||||
// the while body is a single constant (both loop state elements are the same
|
||||
// constant). This means no copies are necessary because both loop state
|
||||
// elements are the same so interchanging them is a no-op.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape loop_state_shape =
|
||||
ShapeUtil::MakeTupleShape({scalar_shape_, scalar_shape_});
|
||||
|
||||
@ -1693,7 +1693,7 @@ TEST_F(CopyInsertionTest, SequentialWhiles) {
|
||||
const Shape loop_state_shape = ShapeUtil::MakeTupleShape(
|
||||
{element_shape, element_shape, element_shape, element_shape});
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto param_0 = builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(0, element_shape, "param_0"));
|
||||
@ -1783,7 +1783,7 @@ TEST_F(CopyInsertionTest, SequentialWhiles) {
|
||||
TEST_F(CopyInsertionTest, WhileBodyWithConstantRoot) {
|
||||
// Test a while body and condition which are each simply a constant (root of
|
||||
// computation is a constant). The body constant should be copied.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto param_0 = builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(0, scalar_shape_, "param_0"));
|
||||
|
||||
@ -60,7 +60,7 @@ TEST_F(CpuHloSupportCheckerTest, SparseUnimplemented) {
|
||||
// Since verifier is reporting sparse layouts as errors, we should
|
||||
// use a regular HloModule instead of VerifiedHloModule to avoid
|
||||
// verifier errors being triggered in the destructor.
|
||||
auto module = HloTestBase::CreateNewModule();
|
||||
auto module = HloTestBase::CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
Status status = checker().Run(module.get()).status();
|
||||
|
||||
@ -58,7 +58,7 @@ TEST_F(InstructionFusionTest, DotOperationFusion_Basic_0) {
|
||||
HloInstruction* dot = builder.AddInstruction(
|
||||
MakeDot(ShapeUtil::MakeShape(F32, {1024, 1}), exp0, arg1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(dot, computation->root_instruction());
|
||||
EXPECT_TRUE(CpuInstructionFusion().Run(module.get()).ValueOrDie());
|
||||
@ -77,7 +77,7 @@ TEST_F(InstructionFusionTest, DotOperationFusion_Basic_1) {
|
||||
HloInstruction* dot = builder.AddInstruction(
|
||||
MakeDot(ShapeUtil::MakeShape(F32, {1, 1024}), arg0, exp1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(dot, computation->root_instruction());
|
||||
EXPECT_TRUE(CpuInstructionFusion().Run(module.get()).ValueOrDie());
|
||||
@ -98,7 +98,7 @@ TEST_F(InstructionFusionTest, DotOperationNoFusion_Bitcast) {
|
||||
HloInstruction* dot = builder.AddInstruction(
|
||||
MakeDot(ShapeUtil::MakeShape(F32, {1024, 1}), bitcast0, arg1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(dot, computation->root_instruction());
|
||||
EXPECT_FALSE(CpuInstructionFusion().Run(module.get()).ValueOrDie());
|
||||
@ -119,7 +119,7 @@ TEST_F(InstructionFusionTest, DotOperationFusion_Reshape) {
|
||||
HloInstruction* dot = builder.AddInstruction(
|
||||
MakeDot(ShapeUtil::MakeShape(F32, {1024, 1}), reshape0, arg1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(dot, computation->root_instruction());
|
||||
EXPECT_TRUE(CpuInstructionFusion().Run(module.get()).ValueOrDie());
|
||||
@ -138,7 +138,7 @@ TEST_F(InstructionFusionTest, DotOperationFusion_TooLarge) {
|
||||
HloInstruction* dot = builder.AddInstruction(
|
||||
MakeDot(ShapeUtil::MakeShape(F32, {1, 32 * 1024}), arg0, exp1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(dot, computation->root_instruction());
|
||||
EXPECT_FALSE(CpuInstructionFusion().Run(module.get()).ValueOrDie());
|
||||
@ -157,7 +157,7 @@ TEST_F(InstructionFusionTest, DotOperationFusion_ElementReuse) {
|
||||
HloInstruction* dot = builder.AddInstruction(
|
||||
MakeDot(ShapeUtil::MakeShape(F32, {2, 1024}), arg0, exp1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(dot, computation->root_instruction());
|
||||
EXPECT_FALSE(CpuInstructionFusion().Run(module.get()).ValueOrDie());
|
||||
@ -321,7 +321,7 @@ TEST_F(OpcodeFusionTest, Exponential_Reshape_Negate) {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(result_shape, HloOpcode::kNegate, reshape2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -350,7 +350,7 @@ TEST_F(OpcodeFusionTest, Broadcast_Reshape_DynamicSlice_Tanh) {
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
dynamic_slice_shape, HloOpcode::kTanh, dynamic_slice4));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -370,7 +370,7 @@ TEST_F(OpcodeFusionTest, Broadcast_Negate) {
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
result_shape, HloOpcode::kNegate, broadcast1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -392,7 +392,7 @@ TEST_F(OpcodeFusionTest, DynamicSlice_Negate) {
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
result_shape, HloOpcode::kNegate, dynamic_slice2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -410,7 +410,7 @@ TEST_F(OpcodeFusionTest, Exponential_Negate) {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(param_shape, HloOpcode::kNegate, exp1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -429,7 +429,7 @@ TEST_F(OpcodeFusionTest, Reshape_Negate) {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(result_shape, HloOpcode::kNegate, reshape1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -447,7 +447,7 @@ TEST_F(OpcodeFusionTest, Reverse_Negate) {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(param_shape, HloOpcode::kNegate, reverse1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -466,7 +466,7 @@ TEST_F(OpcodeFusionTest, Slice_Negate) {
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
ShapeUtil::MakeShape(S32, {2}), HloOpcode::kNegate, slice1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -489,7 +489,7 @@ TEST_F(OpcodeFusionTest, Exponential_Transpose_Negate) {
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
result_shape, HloOpcode::kNegate, transpose2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunFusionAndCheckOpcodesWereFused(
|
||||
@ -498,7 +498,7 @@ TEST_F(OpcodeFusionTest, Exponential_Transpose_Negate) {
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, UnaryMapOfExp) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
|
||||
HloComputation::Builder builder(TestName());
|
||||
Shape shape = ShapeUtil::MakeShape(F32, {3, 4});
|
||||
@ -517,7 +517,7 @@ TEST_F(OpcodeFusionTest, UnaryMapOfExp) {
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, BinaryMapOfExps) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
|
||||
HloComputation::Builder builder(TestName());
|
||||
Shape shape = ShapeUtil::MakeShape(F32, {3, 4});
|
||||
@ -542,7 +542,7 @@ TEST_F(OpcodeFusionTest, BinaryMapOfExps) {
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, DynamicSliceWithDynamicUpdateSlice) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
|
||||
HloComputation::Builder builder(TestName());
|
||||
Shape full_shape = ShapeUtil::MakeShape(F32, {10, 100, 1000});
|
||||
@ -573,7 +573,7 @@ TEST_F(OpcodeFusionTest, DynamicSliceWithDynamicUpdateSlice) {
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, MessOfFusibleNodes) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation::Builder builder(TestName());
|
||||
|
||||
Shape full_shape = ShapeUtil::MakeShape(F32, {4, 100, 10, 100, 50});
|
||||
@ -641,7 +641,7 @@ TEST_F(OpcodeFusionTest, ReuseViaImplicitBroadcastUnary) {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(large_shape, HloOpcode::kExp, small_exp));
|
||||
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto did_fusion = CpuInstructionFusion().Run(module.get());
|
||||
@ -670,7 +670,7 @@ TEST_F(OpcodeFusionTest, ReuseViaImplicitBroadcastBinary) {
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
large_shape, HloOpcode::kAdd, small_exp, large_param));
|
||||
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto did_fusion = CpuInstructionFusion().Run(module.get());
|
||||
@ -712,7 +712,7 @@ void CreateComputationForDotAddOutputFusionTest(const string& test_name,
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, DotAddOutputFusion_1x50x19) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
CreateComputationForDotAddOutputFusionTest(TestName(), module.get(), /*m=*/1,
|
||||
/*k=*/50, /*n=*/19,
|
||||
/*add_extra_use_for_dot=*/false);
|
||||
@ -725,7 +725,7 @@ TEST_F(OpcodeFusionTest, DotAddOutputFusion_1x50x19) {
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, DotAddOutputFusion_19x50x1) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
CreateComputationForDotAddOutputFusionTest(TestName(), module.get(), /*m=*/19,
|
||||
/*k=*/50, /*n=*/1,
|
||||
/*add_extra_use_for_dot=*/false);
|
||||
@ -738,7 +738,7 @@ TEST_F(OpcodeFusionTest, DotAddOutputFusion_19x50x1) {
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, DotAddOutputFusion_19x50x19) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
CreateComputationForDotAddOutputFusionTest(TestName(), module.get(), /*m=*/19,
|
||||
/*k=*/50, /*n=*/19,
|
||||
/*add_extra_use_for_dot=*/false);
|
||||
@ -751,7 +751,7 @@ TEST_F(OpcodeFusionTest, DotAddOutputFusion_19x50x19) {
|
||||
}
|
||||
|
||||
TEST_F(OpcodeFusionTest, DotAddOutputFusion_19x50x1_multi_use) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
CreateComputationForDotAddOutputFusionTest(TestName(), module.get(), /*m=*/19,
|
||||
/*k=*/50, /*n=*/1,
|
||||
/*add_extra_use_for_dot=*/true);
|
||||
|
||||
@ -73,7 +73,7 @@ TEST_F(CpuLayoutAssignmentTest, DotWithConstantRhsTensor) {
|
||||
auto result = builder.AddInstruction(
|
||||
CreateCanonicalDot(result_shape, dot_lhs, dot_rhs));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
ComputationLayout computation_layout(computation->ComputeProgramShape());
|
||||
@ -114,7 +114,7 @@ TEST_F(CpuLayoutAssignmentTest, MultipleDotsWithSameConstantRhsTensor0) {
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
result_shape, HloOpcode::kAdd, dot_a_result, dot_b_result));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
ComputationLayout computation_layout(computation->ComputeProgramShape());
|
||||
@ -158,7 +158,7 @@ TEST_F(CpuLayoutAssignmentTest, MultipleDotsWithSameConstantRhsTensor1) {
|
||||
auto tuple_result = builder.AddInstruction(
|
||||
HloInstruction::CreateTuple({dot_a_result, dot_b_result}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
ComputationLayout computation_layout(computation->ComputeProgramShape());
|
||||
@ -192,7 +192,7 @@ TEST_F(CpuLayoutAssignmentTest, DotWithConstantLhsTensor) {
|
||||
auto dot_result = builder.AddInstruction(
|
||||
CreateCanonicalDot(result_shape, dot_lhs, dot_rhs));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
ComputationLayout computation_layout(computation->ComputeProgramShape());
|
||||
@ -232,7 +232,7 @@ TEST_F(CpuLayoutAssignmentTest, DotWithConstantRhsTensorThroughGTE) {
|
||||
auto dot_result = builder.AddInstruction(
|
||||
CreateCanonicalDot(result_shape, dot_lhs, dot_rhs));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
ComputationLayout computation_layout(computation->ComputeProgramShape());
|
||||
@ -353,7 +353,7 @@ static void AssertCorrectLayoutForDotOutputFusion(
|
||||
}
|
||||
|
||||
TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_1x50x19_dot_idx_0) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
TF_ASSERT_OK_AND_ASSIGN(
|
||||
DotOutputFusionLayoutAssignmentResult layout_assignment_result,
|
||||
RunDotOutputFusion(module.get(), TestName(), /*m=*/1, /*k=*/50, /*n=*/19,
|
||||
@ -365,7 +365,7 @@ TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_1x50x19_dot_idx_0) {
|
||||
}
|
||||
|
||||
TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_1x50x19_dot_idx_1) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
TF_ASSERT_OK_AND_ASSIGN(
|
||||
DotOutputFusionLayoutAssignmentResult layout_assignment_result,
|
||||
RunDotOutputFusion(module.get(), TestName(), /*m=*/1, /*k=*/50, /*n=*/19,
|
||||
@ -377,7 +377,7 @@ TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_1x50x19_dot_idx_1) {
|
||||
}
|
||||
|
||||
TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_19x50x1_dot_idx_0) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
TF_ASSERT_OK_AND_ASSIGN(
|
||||
DotOutputFusionLayoutAssignmentResult layout_assignment_result,
|
||||
RunDotOutputFusion(module.get(), TestName(), /*m=*/19, /*k=*/50, /*n=*/1,
|
||||
@ -389,7 +389,7 @@ TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_19x50x1_dot_idx_0) {
|
||||
}
|
||||
|
||||
TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_19x50x1_dot_idx_1) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
TF_ASSERT_OK_AND_ASSIGN(
|
||||
DotOutputFusionLayoutAssignmentResult layout_assignment_result,
|
||||
RunDotOutputFusion(module.get(), TestName(), /*m=*/19, /*k=*/50, /*n=*/1,
|
||||
@ -401,7 +401,7 @@ TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_19x50x1_dot_idx_1) {
|
||||
}
|
||||
|
||||
TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_19x50x19_dot_idx_0) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
TF_ASSERT_OK_AND_ASSIGN(
|
||||
DotOutputFusionLayoutAssignmentResult layout_assignment_result,
|
||||
RunDotOutputFusion(module.get(), TestName(), /*m=*/19, /*k=*/50, /*n=*/19,
|
||||
@ -413,7 +413,7 @@ TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_19x50x19_dot_idx_0) {
|
||||
}
|
||||
|
||||
TEST_F(CpuLayoutAssignmentTest, DotOutputFusion_19x50x19_dot_idx_1) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
TF_ASSERT_OK_AND_ASSIGN(
|
||||
DotOutputFusionLayoutAssignmentResult layout_assignment_result,
|
||||
RunDotOutputFusion(module.get(), TestName(), /*m=*/19, /*k=*/50, /*n=*/19,
|
||||
|
||||
@ -50,7 +50,7 @@ class CpuEigenDotOperationTest
|
||||
/*entry_point_name=*/"entry",
|
||||
/*relocation_model=*/CpuAotCompilationOptions::RelocationModel::Static};
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(std::move(entry_computation));
|
||||
|
||||
CompileAheadOfTimeAndVerifyIr(std::move(hlo_module), options,
|
||||
|
||||
@ -46,7 +46,7 @@ class CpuExternalConstantsTest : public CpuCodegenTest {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, param, constant));
|
||||
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
CompileAndVerifyIr(std::move(module), filecheck_pattern,
|
||||
|
||||
@ -91,7 +91,7 @@ TEST_P(CpuUnaryIntrinsicTest, DoIt) {
|
||||
/*entry_point_name=*/"entry",
|
||||
/*relocation_model=*/CpuAotCompilationOptions::RelocationModel::Static};
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(std::move(computation));
|
||||
|
||||
string check_lines{spec.check_lines.data(), spec.check_lines.size()};
|
||||
|
||||
@ -56,7 +56,7 @@ TEST_F(CpuNoAliasTest, Concat) {
|
||||
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(std::move(computation));
|
||||
|
||||
// Now that we have an HLO module, build an llvm_ir::AliasAnalysis for it.
|
||||
|
||||
@ -44,7 +44,7 @@ class GpuHloScheduleTest : public HloVerifiedTestBase {
|
||||
.ConsumeValueOrDie();
|
||||
}
|
||||
|
||||
std::unique_ptr<HloModule> CreateNewModule() {
|
||||
std::unique_ptr<HloModule> CreateNewUnverifiedModule() {
|
||||
HloModuleConfig config;
|
||||
auto debug_options = GetDebugOptionsForTest();
|
||||
debug_options.set_xla_gpu_disable_multi_streaming(false);
|
||||
@ -79,7 +79,7 @@ TEST_F(GpuHloScheduleTest, SequentialMatMul) {
|
||||
HloInstruction* dot2 =
|
||||
builder.AddInstruction(CreateCanonicalDot(f32_2x2_, dot1, z));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build(dot2));
|
||||
|
||||
std::unique_ptr<StreamAssignment> streams = AssignStreams(*module);
|
||||
@ -139,7 +139,7 @@ TEST_F(GpuHloScheduleTest, SequentialAdd) {
|
||||
HloInstruction* add3 = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kAdd, add1, add2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build(add3));
|
||||
|
||||
std::unique_ptr<StreamAssignment> streams = AssignStreams(*module);
|
||||
@ -209,7 +209,7 @@ TEST_F(GpuHloScheduleTest, ConcurrentMatMul) {
|
||||
HloInstruction* add =
|
||||
builder.AddInstruction(CreateCanonicalDot(f32_2x2_, dot1, dot2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build(add));
|
||||
|
||||
std::unique_ptr<StreamAssignment> streams = AssignStreams(*module);
|
||||
@ -288,7 +288,7 @@ TEST_F(GpuHloScheduleTest, LatticeMatMul) {
|
||||
HloInstruction* d40 =
|
||||
builder.AddInstruction(CreateCanonicalDot(f32_2x2_, d30, d31));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build(d40));
|
||||
|
||||
std::unique_ptr<StreamAssignment> streams = AssignStreams(*module);
|
||||
|
||||
@ -60,7 +60,7 @@ TEST_F(GpuHloSupportCheckerTest, SparseUnimplemented) {
|
||||
// Since verifier is reporting sparse layouts as errors, we should
|
||||
// use a regular HloModule instead of VerifiedHloModule to avoid
|
||||
// verifier errors being triggered in the destructor.
|
||||
auto module = HloTestBase::CreateNewModule();
|
||||
auto module = HloTestBase::CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
Status status = checker().Run(module.get()).status();
|
||||
|
||||
@ -61,7 +61,7 @@ TEST_F(LayoutAssignmentTest, Elementwise) {
|
||||
HloInstruction::CreateParameter(1, ashape, "y"));
|
||||
auto add = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(ashape, HloOpcode::kAdd, x, y));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation =
|
||||
module->AddEntryComputation(builder.Build(add));
|
||||
|
||||
@ -148,7 +148,7 @@ TEST_F(LayoutAssignmentTest, BatchNormInference) {
|
||||
{operand, scale, offset, mean, variance, epsilon, feature_index},
|
||||
kCudnnBatchNormForwardInferenceCallTarget));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation =
|
||||
module->AddEntryComputation(builder.Build(batchnorm));
|
||||
|
||||
@ -217,7 +217,7 @@ TEST_F(LayoutAssignmentTest, BatchNormTraining) {
|
||||
batchnorm_shape, {operand, scale, offset, epsilon, feature_index},
|
||||
kCudnnBatchNormForwardTrainingCallTarget));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation =
|
||||
module->AddEntryComputation(builder.Build(batchnorm));
|
||||
|
||||
@ -298,7 +298,7 @@ TEST_F(LayoutAssignmentTest, BatchNormGrad) {
|
||||
feature_index},
|
||||
kCudnnBatchNormBackwardCallTarget));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation =
|
||||
module->AddEntryComputation(builder.Build(batchnorm));
|
||||
|
||||
|
||||
@ -41,7 +41,7 @@ TEST_F(InstructionFusionTest,
|
||||
builder.AddInstruction(HloInstruction::CreateBroadcast(
|
||||
ShapeUtil::MakeShape(S32, {1}), exp1, {0}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(broadcast2, computation->root_instruction());
|
||||
EXPECT_FALSE(GpuInstructionFusion(/*may_duplicate=*/true)
|
||||
@ -61,7 +61,7 @@ TEST_F(InstructionFusionTest,
|
||||
builder.AddInstruction(HloInstruction::CreateBroadcast(
|
||||
ShapeUtil::MakeShape(S32, {1}), negate1, {0}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(broadcast2, computation->root_instruction());
|
||||
EXPECT_TRUE(GpuInstructionFusion(/*may_duplicate=*/true)
|
||||
@ -80,7 +80,7 @@ TEST_F(InstructionFusionTest,
|
||||
HloInstruction* reshape2 = builder.AddInstruction(
|
||||
HloInstruction::CreateReshape(ShapeUtil::MakeShape(S32, {}), exp1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(reshape2, computation->root_instruction());
|
||||
EXPECT_TRUE(GpuInstructionFusion(/*may_duplicate=*/true)
|
||||
@ -99,7 +99,7 @@ TEST_F(InstructionFusionTest,
|
||||
HloInstruction* transpose2 = builder.AddInstruction(
|
||||
HloInstruction::CreateTranspose(ShapeUtil::MakeShape(S32, {}), exp1, {}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(transpose2, computation->root_instruction());
|
||||
EXPECT_TRUE(GpuInstructionFusion(/*may_duplicate=*/true)
|
||||
@ -117,7 +117,7 @@ TEST_F(InstructionFusionTest, PotentialBitcastReshapeOfDotUnfused) {
|
||||
auto reshape2 = builder.AddInstruction(HloInstruction::CreateReshape(
|
||||
ShapeUtil::MakeShape(S32, {1, 1, 1}), dot1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(reshape2, computation->root_instruction());
|
||||
EXPECT_FALSE(GpuInstructionFusion(/*may_duplicate=*/true)
|
||||
@ -134,7 +134,7 @@ TEST_F(InstructionFusionTest, PotentialBitcastTransposeOfDotUnfused) {
|
||||
auto transpose2 = builder.AddInstruction(HloInstruction::CreateTranspose(
|
||||
ShapeUtil::MakeShape(S32, {1, 1}), dot1, {0, 1}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(transpose2, computation->root_instruction());
|
||||
EXPECT_FALSE(GpuInstructionFusion(/*may_duplicate=*/true)
|
||||
@ -723,7 +723,7 @@ TEST_F(InstructionFusionTest, AvoidsLargeFusion) {
|
||||
sum = b.AddInstruction(
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, sum, param));
|
||||
}
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(b.Build());
|
||||
EXPECT_TRUE(GpuInstructionFusion(/*may_duplicate=*/true)
|
||||
.Run(module.get())
|
||||
|
||||
@ -580,7 +580,7 @@ TEST_F(MultiOutputFusionTest, AvoidsLargeFusion) {
|
||||
// ...
|
||||
// where each of the (pi * pj)'s is represented as a fusion node so that
|
||||
// multi-output fusion will pay attention to it.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation::Builder b(TestName());
|
||||
Shape shape = ShapeUtil::MakeShape(F32, {10, 100});
|
||||
|
||||
|
||||
@ -30,7 +30,7 @@ namespace gpu {
|
||||
|
||||
class StreamAssignmentTest : public HloVerifiedTestBase {
|
||||
protected:
|
||||
std::unique_ptr<HloModule> CreateNewModule() {
|
||||
std::unique_ptr<HloModule> CreateNewUnverifiedModule() {
|
||||
HloModuleConfig config;
|
||||
auto debug_options = GetDebugOptionsForTest();
|
||||
debug_options.set_xla_gpu_disable_multi_streaming(false);
|
||||
@ -55,7 +55,7 @@ TEST_F(StreamAssignmentTest, SequentialMatMul) {
|
||||
HloInstruction* dot2 =
|
||||
builder.AddInstruction(CreateCanonicalDot(f32_2x2_, dot1, z));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build(dot2));
|
||||
|
||||
std::unique_ptr<StreamAssignment> assignment = AssignStreams(*module);
|
||||
@ -76,7 +76,7 @@ TEST_F(StreamAssignmentTest, ConcurrentMatMul) {
|
||||
HloInstruction* add = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kAdd, dot1, dot2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build(add));
|
||||
|
||||
std::unique_ptr<StreamAssignment> assignment = AssignStreams(*module);
|
||||
@ -120,7 +120,7 @@ TEST_F(StreamAssignmentTest, LatticeMatMul) {
|
||||
HloInstruction* d40 =
|
||||
builder.AddInstruction(CreateCanonicalDot(f32_2x2_, d30, d31));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build(d40));
|
||||
|
||||
std::unique_ptr<StreamAssignment> assignment = AssignStreams(*module);
|
||||
|
||||
@ -23,7 +23,8 @@ limitations under the License.
|
||||
namespace xla {
|
||||
namespace gpu {
|
||||
|
||||
std::unique_ptr<HloModule> GpuCodegenTest::CreateNewModuleWithFTZ(bool ftz) {
|
||||
std::unique_ptr<HloModule> GpuCodegenTest::CreateNewUnverifiedModuleWithFTZ(
|
||||
bool ftz) {
|
||||
HloModuleConfig config;
|
||||
auto debug_options = GetDebugOptionsFromFlags();
|
||||
debug_options.set_xla_gpu_ftz(ftz);
|
||||
|
||||
@ -26,9 +26,9 @@ namespace gpu {
|
||||
// Tests that verify IR or PTX emitted by the GPU backend is as expected.
|
||||
class GpuCodegenTest : public LlvmIrGenTestBase {
|
||||
protected:
|
||||
// Like HloTestBase::CreateNewModule(), with a flag for configuring the ftz
|
||||
// option.
|
||||
std::unique_ptr<HloModule> CreateNewModuleWithFTZ(bool ftz);
|
||||
// Like HloTestBase::CreateNewUnverifiedModule(), with a flag for configuring
|
||||
// the ftz option.
|
||||
std::unique_ptr<HloModule> CreateNewUnverifiedModuleWithFTZ(bool ftz);
|
||||
|
||||
// Compiles the given HLO module to PTX and verifies the PTX matches the given
|
||||
// FileCheck pattern. (See http://llvm.org/docs/CommandGuide/FileCheck.html).
|
||||
|
||||
@ -46,7 +46,7 @@ TEST_F(GpuCopyTest, UseMemcpy) {
|
||||
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(std::move(computation));
|
||||
|
||||
// There should not be any kernel prefixed "copy".
|
||||
|
||||
@ -39,7 +39,7 @@ class GpuFtzTest : public GpuCodegenTest {
|
||||
/* parameter_number=*/1, param_shape, "y"));
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(param_shape, op, x, y));
|
||||
|
||||
auto hlo_module = CreateNewModuleWithFTZ(ftz_);
|
||||
auto hlo_module = CreateNewUnverifiedModuleWithFTZ(ftz_);
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
return hlo_module;
|
||||
}
|
||||
@ -54,7 +54,7 @@ class GpuFtzTest : public GpuCodegenTest {
|
||||
/* parameter_number=*/0, param_shape, "x"));
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(param_shape, op, x));
|
||||
|
||||
auto hlo_module = CreateNewModuleWithFTZ(ftz_);
|
||||
auto hlo_module = CreateNewUnverifiedModuleWithFTZ(ftz_);
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
return hlo_module;
|
||||
}
|
||||
|
||||
@ -51,7 +51,7 @@ TEST_F(GpuIndexTest, CompatibleUseLinearIndex) {
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
ShapeUtil::MakeShape(PRED, {5, 7, 2}), HloOpcode::kGe, param_x, param_y));
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Check the optimized IR as the unoptimized IR contains dead udiv and urem.
|
||||
|
||||
@ -48,7 +48,7 @@ TEST_F(GpuLdgTest, LdgForParamRead) {
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, param, param));
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(std::move(computation));
|
||||
|
||||
CompileAndVerifyPtx(std::move(hlo_module), R"(
|
||||
@ -73,7 +73,7 @@ TEST_F(GpuLdgTest, LdgForNonParamRead) {
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({add, square}));
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(std::move(computation));
|
||||
|
||||
CompileAndVerifyPtx(std::move(hlo_module), R"(
|
||||
@ -95,7 +95,7 @@ TEST_F(GpuLdgTest, LdgForNonParamRead) {
|
||||
// reduce in the foreseeable future. But if that turns out to be wrong, I give
|
||||
// you, future reader, permission to delete this test.
|
||||
TEST_F(GpuLdgTest, NoLdgWhenSharingBuffer) {
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
HloComputation::Builder builder(TestName());
|
||||
|
||||
HloComputation* reduce_computation;
|
||||
|
||||
@ -47,7 +47,7 @@ TEST_F(GpuNoAliasTest, Concat) {
|
||||
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(std::move(computation));
|
||||
|
||||
CompileAndVerifyIr(std::move(hlo_module),
|
||||
|
||||
@ -29,7 +29,7 @@ namespace {
|
||||
class WhileTransformerTest : public HloTestBase {
|
||||
protected:
|
||||
WhileTransformerTest()
|
||||
: module_(CreateNewModule()),
|
||||
: module_(CreateNewUnverifiedModule()),
|
||||
induction_variable_shape_(ShapeUtil::MakeShape(S32, {})),
|
||||
data_shape_(ShapeUtil::MakeShape(F32, {8})),
|
||||
condition_result_shape_(ShapeUtil::MakeShape(PRED, {})) {}
|
||||
|
||||
@ -65,7 +65,7 @@ class HloComputationTest : public HloTestBase {
|
||||
};
|
||||
|
||||
TEST_F(HloComputationTest, GetEmbeddedComputationsEmpty) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto negate_computation =
|
||||
module->AddEntryComputation(CreateNegateComputation());
|
||||
EXPECT_TRUE(negate_computation->MakeEmbeddedComputationsList().empty());
|
||||
@ -73,7 +73,7 @@ TEST_F(HloComputationTest, GetEmbeddedComputationsEmpty) {
|
||||
|
||||
TEST_F(HloComputationTest, GetEmbeddedComputationsOneComputation) {
|
||||
// Create computation which calls one other computation.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto negate_computation =
|
||||
module->AddEmbeddedComputation(CreateNegateComputation());
|
||||
auto map_computation =
|
||||
@ -85,7 +85,7 @@ TEST_F(HloComputationTest, GetEmbeddedComputationsOneComputation) {
|
||||
|
||||
TEST_F(HloComputationTest, GetEmbeddedComputationsDiamond) {
|
||||
// Create computations with a diamond-shaped callgraph.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto negate_computation =
|
||||
module->AddEmbeddedComputation(CreateNegateComputation());
|
||||
auto map1_computation =
|
||||
@ -119,7 +119,7 @@ TEST_F(HloComputationTest, PostOrderSingleton) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto constant = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(42.0f)));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_THAT(computation->MakeInstructionPostOrder(), ElementsAre(constant));
|
||||
}
|
||||
@ -134,7 +134,7 @@ TEST_F(HloComputationTest, PostOrderSimple) {
|
||||
HloInstruction::CreateUnary(r0f32_, HloOpcode::kNegate, constant));
|
||||
auto negate2 = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(r0f32_, HloOpcode::kNegate, negate1));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_THAT(computation->MakeInstructionPostOrder(),
|
||||
ElementsAre(constant, negate1, negate2));
|
||||
@ -151,7 +151,7 @@ TEST_F(HloComputationTest, PostOrderTrace) {
|
||||
builder.AddInstruction(HloInstruction::CreateTrace("foobar", negate1));
|
||||
auto negate2 = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(r0f32_, HloOpcode::kNegate, negate1));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
// Trace instructions should be at the end of the sort.
|
||||
EXPECT_THAT(computation->MakeInstructionPostOrder(),
|
||||
@ -170,7 +170,7 @@ TEST_F(HloComputationTest, PostOrderDisconnectedInstructions) {
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(42.0f)));
|
||||
auto constant4 = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(42.0f)));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_THAT(computation->MakeInstructionPostOrder(),
|
||||
UnorderedElementsAre(constant1, constant2, constant3, constant4));
|
||||
@ -192,7 +192,7 @@ TEST_F(HloComputationTest, PostOrderWithMultipleRoots) {
|
||||
r0f32_, HloOpcode::kAdd, constant2, constant3));
|
||||
auto add3 = builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
r0f32_, HloOpcode::kAdd, constant1, constant3));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
auto post_order = computation->MakeInstructionPostOrder();
|
||||
EXPECT_EQ(6, post_order.size());
|
||||
@ -217,7 +217,7 @@ TEST_F(HloComputationTest, VisitWithMultipleRoots) {
|
||||
constant2, constant3));
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(r0f32_, HloOpcode::kAdd,
|
||||
constant1, constant3));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
// Visitor which keeps track of which instructions have been visited.
|
||||
class TestVisitor : public DfsHloVisitorWithDefault {
|
||||
@ -257,7 +257,7 @@ TEST_F(HloComputationTest, DeepCopyArray) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto constant = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR1<float>({1.0, 2.0, 3.0})));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
auto copy = computation->DeepCopyInstruction(constant).ValueOrDie();
|
||||
|
||||
@ -274,7 +274,7 @@ TEST_F(HloComputationTest, DeepCopyTuple) {
|
||||
auto tuple = builder.AddInstruction(
|
||||
HloInstruction::CreateTuple({constant1, constant2}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
auto tuple_copy = computation->DeepCopyInstruction(tuple).ValueOrDie();
|
||||
|
||||
@ -376,7 +376,7 @@ TEST_F(HloComputationTest, DeepCopyToken) {
|
||||
// copied.
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto token = builder.AddInstruction(HloInstruction::CreateToken());
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
auto copy = computation->DeepCopyInstruction(token).ValueOrDie();
|
||||
|
||||
@ -393,7 +393,7 @@ TEST_F(HloComputationTest, DeepCopyTokenTuple) {
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(42.0)));
|
||||
auto tuple =
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({token, constant}));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
auto copy = computation->DeepCopyInstruction(tuple).ValueOrDie();
|
||||
|
||||
@ -412,7 +412,7 @@ TEST_F(HloComputationTest, CycleDetection) {
|
||||
HloInstruction::CreateUnary(r0f32_, HloOpcode::kNegate, constant));
|
||||
auto add = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(r0f32_, HloOpcode::kAdd, negate, negate));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
// Add a control dependency to create a cycle.
|
||||
ASSERT_IS_OK(add->AddControlDependencyTo(negate));
|
||||
@ -440,7 +440,7 @@ TEST_F(HloComputationTest, RemoveInstructionWithDuplicateOperand) {
|
||||
r0f32_, HloOpcode::kAdd, dead_negate, dead_negate));
|
||||
auto negate = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(r0f32_, HloOpcode::kNegate, constant));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(4, computation->instruction_count());
|
||||
EXPECT_THAT(computation->root_instruction(), op::Negate(constant));
|
||||
@ -466,7 +466,7 @@ TEST_F(HloComputationTest, CloneWithControlDependency) {
|
||||
HloInstruction::CreateParameter(0, r0f32_, "param0"));
|
||||
auto negate = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(r0f32_, HloOpcode::kNegate, param));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation =
|
||||
module->AddEntryComputation(builder.Build(/*root_instruction=*/add));
|
||||
|
||||
@ -505,7 +505,7 @@ TEST_F(HloComputationTest, Stringification) {
|
||||
2, PrecisionConfig::DEFAULT);
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateDot(sout, x, reshape, dot_dnums, precision_config));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto options = HloPrintOptions().set_print_metadata(false);
|
||||
@ -540,7 +540,7 @@ TEST_F(HloComputationTest, StringificationIndent) {
|
||||
2, PrecisionConfig::DEFAULT);
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateDot(sout, x, reshape, dot_dnums, precision_config));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto options =
|
||||
@ -576,7 +576,7 @@ TEST_F(HloComputationTest, StringificationCanonical) {
|
||||
2, PrecisionConfig::DEFAULT);
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateDot(sout, x, reshape, dot_dnums, precision_config));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto options = HloPrintOptions().set_print_metadata(false);
|
||||
|
||||
@ -387,7 +387,7 @@ TEST_F(FusionCostAnalysis, LoopFusion) {
|
||||
HloInstruction::CreateBinary(r2f32, HloOpcode::kSubtract, mul, clamp));
|
||||
auto tuple = HloInstruction::CreateTuple({sub, sub, mul, c1});
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto* computation = module->AddEntryComputation(builder.Build());
|
||||
auto* fusion = computation->CreateFusionInstruction(
|
||||
{sub, mul, exp, clamp, add}, HloInstruction::FusionKind::kLoop);
|
||||
@ -429,7 +429,7 @@ TEST_F(FusionCostAnalysis, NoLayout) {
|
||||
auto add = builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
shape_with_layout, HloOpcode::kAdd, c1, broadcast));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto* computation = module->AddEntryComputation(builder.Build());
|
||||
auto* fusion = computation->CreateFusionInstruction(
|
||||
{add, broadcast}, HloInstruction::FusionKind::kLoop);
|
||||
@ -472,7 +472,7 @@ TEST_F(DomainCostAnalysis, DomainCost) {
|
||||
auto domain = builder.AddInstruction(
|
||||
HloInstruction::CreateDomain(tuple->shape(), tuple, nullptr, nullptr));
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_EQ(hlo_module->entry_computation()->root_instruction(), domain);
|
||||
|
||||
@ -43,7 +43,7 @@ using ::testing::UnorderedElementsAre;
|
||||
class HloDataflowAnalysisTest : public HloTestBase,
|
||||
public ::testing::WithParamInterface<bool> {
|
||||
protected:
|
||||
HloDataflowAnalysisTest() : module_(CreateNewModule()) {}
|
||||
HloDataflowAnalysisTest() : module_(CreateNewUnverifiedModule()) {}
|
||||
|
||||
// Run dataflow analysis on the member module. For convenience returns a
|
||||
// reference to the generated analysis stored in analysis_.
|
||||
@ -1884,7 +1884,7 @@ INSTANTIATE_TEST_CASE_P(HloDataflowAnalysisInstantiation,
|
||||
class HloDataflowAnalysisTestBase : public HloTestBase {
|
||||
protected:
|
||||
void BuildModule(std::unique_ptr<HloComputation> computation) {
|
||||
module_ = CreateNewModule();
|
||||
module_ = CreateNewUnverifiedModule();
|
||||
computation_ = module_->AddEntryComputation(std::move(computation));
|
||||
}
|
||||
|
||||
@ -2476,7 +2476,7 @@ TEST_F(CanShareOperandBufferWithUserTest, WhileCanShare) {
|
||||
return builder.Build();
|
||||
};
|
||||
|
||||
module_ = CreateNewModule();
|
||||
module_ = CreateNewUnverifiedModule();
|
||||
HloComputation* cond_computation =
|
||||
module_->AddEmbeddedComputation(make_cond());
|
||||
HloComputation* body_computation =
|
||||
@ -2511,7 +2511,7 @@ TEST_F(CanShareOperandBufferWithUserTest, CallToComputationWithFusionRoot) {
|
||||
auto add = sub_builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, sub_param, ones));
|
||||
|
||||
module_ = CreateNewModule();
|
||||
module_ = CreateNewUnverifiedModule();
|
||||
auto sub_computation = module_->AddEmbeddedComputation(sub_builder.Build());
|
||||
sub_computation->CreateFusionInstruction({add, ones},
|
||||
HloInstruction::FusionKind::kLoop);
|
||||
|
||||
@ -59,7 +59,7 @@ TEST_F(HloDceTest, NoDeadCode) {
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
constant1->shape(), HloOpcode::kAdd, constant1, constant2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_EQ(3, computation->instruction_count());
|
||||
@ -80,7 +80,7 @@ TEST_F(HloDceTest, InstructionsWithSideEffect) {
|
||||
HloInstruction::CreateSend(constant, token, /*channel_id=*/0));
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_EQ(4, computation->instruction_count());
|
||||
@ -110,7 +110,7 @@ TEST_F(HloDceTest, DeadParameters) {
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
live_param->shape(), HloOpcode::kNegate, live_param));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_EQ(5, computation->instruction_count());
|
||||
@ -150,7 +150,7 @@ TEST_F(HloDceTest, ControlDependencies) {
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
constant1->shape(), HloOpcode::kAdd, constant1, constant2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Add a control dependency between two instructions.
|
||||
@ -175,7 +175,7 @@ TEST_F(HloDceTest, ControlDependencies) {
|
||||
|
||||
// Tests that a dead call instruction is removed.
|
||||
TEST_F(HloDceTest, DeadInstructionWithCalledComputation) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
Shape shape = ShapeUtil::MakeShape(F32, {});
|
||||
|
||||
// Called computation for the call instruction.
|
||||
@ -215,7 +215,7 @@ TEST_F(HloDceTest, DeadInstructionWithCalledComputation) {
|
||||
// Tests that a while instruction with an infeed (effectul instruction) in its
|
||||
// body is not removed, even its user count is 0.
|
||||
TEST_F(HloDceTest, CalledComputationWithSideEffect) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
Shape shape = ShapeUtil::MakeShape(F32, {});
|
||||
|
||||
// Condition computation of a while instruction.
|
||||
@ -270,7 +270,7 @@ TEST_F(HloDceTest, CalledComputationWithSideEffect) {
|
||||
|
||||
// Tests that a nested call instruction with a side effect is not removed.
|
||||
TEST_F(HloDceTest, CalledComputationWithNestedSideEffect) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
Shape shape = ShapeUtil::MakeShape(F32, {});
|
||||
|
||||
// Nested called computation with a side effect.
|
||||
@ -323,7 +323,7 @@ TEST_F(HloDceTest, CalledComputationWithNestedSideEffect) {
|
||||
}
|
||||
|
||||
TEST_F(HloDceTest, RemoveDeadSubcomputation) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation::Builder builder(TestName());
|
||||
|
||||
HloComputation::Builder subcomp_builder("reduction_subcomp");
|
||||
@ -364,7 +364,7 @@ TEST_F(HloDceTest, RemoveDeadSubcomputation) {
|
||||
}
|
||||
|
||||
TEST_F(HloDceTest, KeepUsedSubcomputation) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation::Builder builder(TestName());
|
||||
|
||||
HloComputation::Builder subcomp_builder("reduction_subcomp");
|
||||
|
||||
@ -65,7 +65,7 @@ TEST_F(HloSchedulingTest, LastUseScheduledFirst) {
|
||||
auto sub = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(vec, HloOpcode::kSubtract, add, negate));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
HloMemoryScheduler scheduler([](const BufferValue& buffer) {
|
||||
@ -172,7 +172,7 @@ TEST_F(HloSchedulingTest, TuplesAreAccountedCorrectly) {
|
||||
builder.AddInstruction(HloInstruction::CreateBinary(r1f32, HloOpcode::kAdd,
|
||||
tuple_elm, abs_abs2));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
TF_ASSERT_OK_AND_ASSIGN(HloSchedule schedule,
|
||||
ScheduleModule(*module,
|
||||
@ -218,7 +218,7 @@ TEST_F(HloSchedulingTest, MultiOutputFusionAccountedCorrectly) {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(r1f32, HloOpcode::kAdd, tuple_elm, exp));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto fusion = computation->CreateFusionInstruction(
|
||||
@ -242,7 +242,7 @@ TEST_F(HloSchedulingTest, MultiOutputFusionAccountedCorrectly) {
|
||||
}
|
||||
|
||||
TEST_F(HloSchedulingTest, HeapSimulatorAccountsForSubcomputations) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape r1f32 = ShapeUtil::MakeShape(F32, {4});
|
||||
|
||||
// param != 0
|
||||
|
||||
@ -63,7 +63,7 @@ class HloModuleTest : public HloTestBase {
|
||||
|
||||
TEST_F(HloModuleTest, OneComputationPostOrder) {
|
||||
// Create a module with a single computation.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(CreateConstantComputation());
|
||||
|
||||
EXPECT_THAT(module->MakeComputationPostOrder(),
|
||||
@ -72,7 +72,7 @@ TEST_F(HloModuleTest, OneComputationPostOrder) {
|
||||
|
||||
TEST_F(HloModuleTest, TwoComputationsPostOrder) {
|
||||
// Create a module with two unconnected computations.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation1 = module->AddEntryComputation(CreateConstantComputation());
|
||||
auto computation2 =
|
||||
module->AddEmbeddedComputation(CreateConstantComputation());
|
||||
@ -88,7 +88,7 @@ TEST_F(HloModuleTest, TwoComputationsPostOrder) {
|
||||
|
||||
TEST_F(HloModuleTest, CloneTest) {
|
||||
// Create and copy a module with a diamond call graph of computations.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation1 =
|
||||
module->AddEmbeddedComputation(CreateConstantComputation());
|
||||
auto computation2 =
|
||||
@ -111,7 +111,7 @@ TEST_F(HloModuleTest, CloneTest) {
|
||||
}
|
||||
|
||||
TEST_F(HloModuleTest, CloneHasFusion) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
|
||||
// Create the fused computation.
|
||||
HloComputation* fused_computation;
|
||||
@ -154,7 +154,7 @@ TEST_F(HloModuleTest, CloneHasFusion) {
|
||||
|
||||
TEST_F(HloModuleTest, DiamondComputationsPostOrder) {
|
||||
// Create a module with a diamond call graph of computations.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation1 =
|
||||
module->AddEmbeddedComputation(CreateConstantComputation());
|
||||
auto computation2 =
|
||||
@ -174,7 +174,7 @@ TEST_F(HloModuleTest, DiamondComputationsPostOrder) {
|
||||
|
||||
TEST_F(HloModuleTest, LargeConstantToString) {
|
||||
// Create a module with a single computation.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder("Constant");
|
||||
std::vector<float> values(16, 42.0);
|
||||
builder.AddInstruction(
|
||||
@ -194,8 +194,8 @@ TEST_F(HloModuleTest, LargeConstantToString) {
|
||||
}
|
||||
|
||||
TEST_F(HloModuleTest, UniqueModuleId) {
|
||||
auto module_a = CreateNewModule();
|
||||
auto module_b = CreateNewModule();
|
||||
auto module_a = CreateNewUnverifiedModule();
|
||||
auto module_b = CreateNewUnverifiedModule();
|
||||
EXPECT_NE(module_a->unique_id(), module_b->unique_id());
|
||||
}
|
||||
|
||||
|
||||
@ -53,7 +53,7 @@ TEST_F(HloOrderingTest, InstructionsInDifferentComputations) {
|
||||
// %c = Constant(42.0f)
|
||||
//
|
||||
// This results in a diamond-shaped callgraph.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape scalar_shape = ShapeUtil::MakeShape(xla::F32, {});
|
||||
|
||||
auto builder_c = HloComputation::Builder("C");
|
||||
@ -126,7 +126,7 @@ TEST_F(HloOrderingTest, InstructionsInWhileComputations) {
|
||||
// %constant = Constant(1.0)
|
||||
// return While(%constant, body, condition)
|
||||
//
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape scalar_shape = ShapeUtil::MakeShape(xla::F32, {});
|
||||
|
||||
auto body_builder = HloComputation::Builder("body");
|
||||
@ -176,7 +176,7 @@ TEST_F(HloOrderingTest, InstructionsInWhileComputations) {
|
||||
|
||||
TEST_F(HloOrderingTest, ParametersDefinedBeforeOthers) {
|
||||
// Entry parameter should always be defined before other instruction.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape scalar_shape = ShapeUtil::MakeShape(xla::F32, {});
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto constant = builder.AddInstruction(
|
||||
@ -209,7 +209,7 @@ TEST_F(HloOrderingTest, ValuesInWhileComputations) {
|
||||
// %while = While(%constant, body, condition)
|
||||
// %add = Add(%constant, %while)
|
||||
//
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape scalar_shape = ShapeUtil::MakeShape(xla::F32, {});
|
||||
|
||||
auto body_builder = HloComputation::Builder("body");
|
||||
@ -407,7 +407,7 @@ TEST_F(HloOrderingTest,
|
||||
// %dead = Constant(123.0)
|
||||
//
|
||||
// %root should interfere with %dead.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape scalar_shape = ShapeUtil::MakeShape(xla::F32, {});
|
||||
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
@ -455,7 +455,7 @@ TEST_F(HloOrderingTest,
|
||||
// ROOT %call = call({%c}), subcomputation
|
||||
//
|
||||
// %root should interfere with %dead.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
const Shape scalar_shape = ShapeUtil::MakeShape(xla::F32, {});
|
||||
|
||||
auto subbuilder = HloComputation::Builder(TestName() + ".sub");
|
||||
|
||||
@ -66,7 +66,7 @@ class HloSubcomputationUnificationTest : public HloTestBase {
|
||||
};
|
||||
|
||||
TEST_F(HloSubcomputationUnificationTest, UnifyIdentities) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
|
||||
auto callee1 =
|
||||
@ -103,7 +103,7 @@ TEST_F(HloSubcomputationUnificationTest, UnifyIdentities) {
|
||||
}
|
||||
|
||||
TEST_F(HloSubcomputationUnificationTest, UnifyAdditions) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
|
||||
auto callee1 =
|
||||
@ -143,7 +143,7 @@ TEST_F(HloSubcomputationUnificationTest, UnifyAdditions) {
|
||||
|
||||
// Do not unify subcomputations with different parameter shapes.
|
||||
TEST_F(HloSubcomputationUnificationTest, DifferentParameterShapes) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
|
||||
auto callee1 =
|
||||
@ -184,7 +184,7 @@ TEST_F(HloSubcomputationUnificationTest, DifferentParameterShapes) {
|
||||
// Regression test for b/31466798. Checks that entry_computation is still valid
|
||||
// after unification.
|
||||
TEST_F(HloSubcomputationUnificationTest, TwoIdenticalComputations) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
HloComputation::Builder builder("pow");
|
||||
auto x =
|
||||
|
||||
@ -66,7 +66,7 @@ TEST_F(HloVerifierTest, NullInstructionParent) {
|
||||
HloInstruction::CreateParameter(0, scalar_shape, "param"));
|
||||
HloInstruction* negate = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(scalar_shape, HloOpcode::kNegate, param));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
TF_ASSERT_OK(verifier().Run(module.get()).status());
|
||||
@ -85,7 +85,7 @@ TEST_F(HloVerifierTest, NullComputationParent) {
|
||||
HloInstruction::CreateParameter(0, scalar_shape, "param"));
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(scalar_shape, HloOpcode::kNegate, param));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
HloComputation* computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
TF_ASSERT_OK(verifier().Run(module.get()).status());
|
||||
@ -104,7 +104,7 @@ TEST_F(HloVerifierTest, DifferentOperandParents) {
|
||||
HloInstruction::CreateParameter(0, scalar_shape, "param"));
|
||||
HloInstruction* negate = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(scalar_shape, HloOpcode::kNegate, param));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
HloComputation::Builder emb_builder(TestName());
|
||||
@ -138,7 +138,7 @@ TEST_F(HloVerifierTest, ResetsShapeVerifierState) {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(s2, HloOpcode::kMultiply, add, add));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Run the verifier twice. It should fail both times, because it shouldn't
|
||||
@ -303,7 +303,7 @@ TEST_F(HloVerifierTest, NegativeInteriorPaddingNotAllowed) {
|
||||
HloInstruction::CreateConstant(LiteralUtil::Zero(F32))),
|
||||
padding_config));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto status = verifier().Run(module.get()).status();
|
||||
@ -327,7 +327,7 @@ TEST_F(HloVerifierTest, PadNegativeInteriorDilationNotAllowed) {
|
||||
HloInstruction::CreateConstant(LiteralUtil::Zero(F32).Clone())),
|
||||
padding_config));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
EXPECT_THAT(verifier().Run(module.get()).status().error_message(),
|
||||
|
||||
@ -117,7 +117,7 @@ TEST_F(InstructionFusionTest, PotentialBitcastReshapeOfParameterUnfused) {
|
||||
auto reshape1 = builder.AddInstruction(
|
||||
HloInstruction::CreateReshape(ShapeUtil::MakeShape(S32, {1, 1}), param0));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(reshape1, computation->root_instruction());
|
||||
EXPECT_FALSE(
|
||||
@ -133,7 +133,7 @@ TEST_F(InstructionFusionTest, PotentialBitcastSimpleReshapeOfParameterUnfused) {
|
||||
auto reshape1 = builder.AddInstruction(
|
||||
HloInstruction::CreateReshape(ShapeUtil::MakeShape(S32, {1, 1}), param0));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(reshape1, computation->root_instruction());
|
||||
EXPECT_FALSE(
|
||||
@ -149,7 +149,7 @@ TEST_F(InstructionFusionTest, PotentialBitcastTransposeOfParameterUnfused) {
|
||||
auto transpose1 = builder.AddInstruction(HloInstruction::CreateTranspose(
|
||||
ShapeUtil::MakeShape(S32, {}), param0, {}));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(transpose1, computation->root_instruction());
|
||||
EXPECT_FALSE(
|
||||
@ -172,7 +172,7 @@ TEST_F(InstructionFusionTest, AvoidDuplicationIfNotAllFusible) {
|
||||
HloInstruction* unary = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kAbs, binary1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(unary, computation->root_instruction());
|
||||
EXPECT_FALSE(
|
||||
@ -361,7 +361,7 @@ TEST_F(InstructionFusionTest, AllowUnaryDuplication) {
|
||||
HloInstruction* unary2 = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kAbs, unary1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(unary2, computation->root_instruction());
|
||||
EXPECT_TRUE(
|
||||
@ -385,7 +385,7 @@ TEST_F(InstructionFusionTest, AllowEffectiveUnaryDuplication) {
|
||||
HloInstruction* unary = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kAbs, binary1));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
EXPECT_EQ(unary, computation->root_instruction());
|
||||
EXPECT_TRUE(
|
||||
|
||||
@ -54,7 +54,7 @@ TEST_F(ReducePrecisionInsertionTest, BeforeUnaryInstruction) {
|
||||
HloInstruction* b = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, a));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected state before adding ops.
|
||||
@ -84,7 +84,7 @@ TEST_F(ReducePrecisionInsertionTest, BeforeBinaryInstruction) {
|
||||
HloInstruction* c = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, a, b));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected state before adding ops.
|
||||
@ -113,7 +113,7 @@ TEST_F(ReducePrecisionInsertionTest, BeforeZeroInputInstruction) {
|
||||
HloInstruction* b = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, a));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected state before adding ops.
|
||||
@ -146,7 +146,7 @@ TEST_F(ReducePrecisionInsertionTest, AvoidAddingDuplicateInstructions) {
|
||||
HloInstruction* d = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, b, c));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected state before adding ops.
|
||||
@ -178,7 +178,7 @@ TEST_F(ReducePrecisionInsertionTest, AfterRootInstruction) {
|
||||
HloInstruction* b = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, a));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected state before adding ops.
|
||||
@ -215,7 +215,7 @@ TEST_F(ReducePrecisionInsertionTest, AfterNonRootInstruction) {
|
||||
HloInstruction* c = builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, a_cos, b_cos));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected graph before adding ops.
|
||||
@ -242,7 +242,7 @@ TEST_F(ReducePrecisionInsertionTest, OutputIsNotFloat) {
|
||||
HloInstruction* y = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, x));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected graph before adding ops.
|
||||
@ -268,7 +268,7 @@ TEST_F(ReducePrecisionInsertionTest, ShouldReduceOutputPrecisionIsFalse) {
|
||||
HloInstruction* y = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, x));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected graph before adding ops.
|
||||
@ -294,7 +294,7 @@ TEST_F(ReducePrecisionInsertionTest, InsertionIsNotRecursive) {
|
||||
HloInstruction* b = builder.AddInstruction(
|
||||
HloInstruction::CreateReducePrecision(shape, a, 8, 23));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected state before adding ops.
|
||||
@ -321,7 +321,7 @@ TEST_F(ReducePrecisionInsertionTest, SkipRedundantReducePrecisionAfter) {
|
||||
HloInstruction* y = builder.AddInstruction(
|
||||
HloInstruction::CreateReducePrecision(shape, x, 5, 10));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected graph before adding ops.
|
||||
@ -349,7 +349,7 @@ TEST_F(ReducePrecisionInsertionTest, AddNonRedundantReducePrecision) {
|
||||
HloInstruction* y = builder.AddInstruction(
|
||||
HloInstruction::CreateReducePrecision(shape, x, 8, 23));
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Confirm expected graph before adding ops.
|
||||
@ -375,7 +375,7 @@ TEST_F(ReducePrecisionInsertionTest, IgnoreOpsInsideFusionNode) {
|
||||
builder.AddInstruction(HloInstruction::CreateParameter(0, shape, "x"));
|
||||
HloInstruction* y = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, x));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Manually fuse the kCos operation into a fusion operation.
|
||||
@ -411,7 +411,7 @@ TEST_F(ReducePrecisionInsertionTest, OpGetsInsertedInHeadOfFusionNode) {
|
||||
builder.AddInstruction(HloInstruction::CreateParameter(0, shape, "x"));
|
||||
HloInstruction* y = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, x));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Manually fuse the kCos operation into a fusion operation.
|
||||
@ -458,7 +458,7 @@ TEST_F(ReducePrecisionInsertionTest, OpGetsInsertedInTailOfFusionNode) {
|
||||
builder.AddInstruction(HloInstruction::CreateParameter(0, shape, "x"));
|
||||
HloInstruction* y = builder.AddInstruction(
|
||||
HloInstruction::CreateUnary(shape, HloOpcode::kCos, x));
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto computation = module->AddEntryComputation(builder.Build());
|
||||
|
||||
// Manually fuse the kCos operation into a fusion operation.
|
||||
|
||||
@ -172,7 +172,7 @@ TEST_F(TransposeFoldingTest, FuseDotWithConstantOperands) {
|
||||
HloInstruction* mul = builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
add->shape(), HloOpcode::kMultiply, add, sub));
|
||||
|
||||
auto module = CreateNewModule("fuse_with_constant_operands");
|
||||
auto module = CreateNewUnverifiedModule("fuse_with_constant_operands");
|
||||
HloComputation* entry_computation =
|
||||
module->AddEntryComputation(builder.Build(mul));
|
||||
HloInstruction* call = module->OutlineExpressionFromComputation(
|
||||
@ -247,7 +247,7 @@ TEST_F(TransposeFoldingTest, FoldConvDimSwapTransposeRhs) {
|
||||
conv_shape.ValueOrDie(), x, transpose_y,
|
||||
/*feature_group_count=*/1, window, dnums, DefaultPrecisionConfig(2)));
|
||||
|
||||
auto module = CreateNewModule("test_module");
|
||||
auto module = CreateNewUnverifiedModule("test_module");
|
||||
HloComputation* entry_computation =
|
||||
module->AddEntryComputation(builder.Build(conv));
|
||||
FoldTranspose(module.get());
|
||||
@ -302,7 +302,7 @@ TEST_F(TransposeFoldingTest, FoldConvComplexTransposeRhs) {
|
||||
conv_shape.ValueOrDie(), x, transpose_y,
|
||||
/*feature_group_count=*/1, window, dnums, DefaultPrecisionConfig(2)));
|
||||
|
||||
auto module = CreateNewModule("test_module");
|
||||
auto module = CreateNewUnverifiedModule("test_module");
|
||||
HloComputation* entry_computation =
|
||||
module->AddEntryComputation(builder.Build(conv));
|
||||
FoldTranspose(module.get());
|
||||
@ -362,7 +362,7 @@ TEST_F(TransposeFoldingTest, FoldConvTransposeLhs) {
|
||||
conv_shape.ValueOrDie(), transpose_x, y,
|
||||
/*feature_group_count=*/1, window, dnums, DefaultPrecisionConfig(2)));
|
||||
|
||||
auto module = CreateNewModule("test_module");
|
||||
auto module = CreateNewUnverifiedModule("test_module");
|
||||
HloComputation* entry_computation =
|
||||
module->AddEntryComputation(builder.Build(conv));
|
||||
FoldTranspose(module.get());
|
||||
@ -428,7 +428,7 @@ TEST_F(TransposeFoldingTest, FoldConvComplexTransposeLhs) {
|
||||
conv_shape.ValueOrDie(), transpose_x, y,
|
||||
/*feature_group_count=*/1, window, dnums, DefaultPrecisionConfig(2)));
|
||||
|
||||
auto module = CreateNewModule("test_module");
|
||||
auto module = CreateNewUnverifiedModule("test_module");
|
||||
HloComputation* entry_computation =
|
||||
module->AddEntryComputation(builder.Build(conv));
|
||||
FoldTranspose(module.get());
|
||||
|
||||
@ -48,7 +48,7 @@ class TuplePointsToAnalysisTest : public HloTestBase {
|
||||
}
|
||||
|
||||
void BuildModule(std::unique_ptr<HloComputation> computation) {
|
||||
module_ = CreateNewModule();
|
||||
module_ = CreateNewUnverifiedModule();
|
||||
module_->AddEntryComputation(std::move(computation));
|
||||
}
|
||||
|
||||
@ -809,7 +809,7 @@ TEST_F(FusionPointsToAnalysisTest, FusionParam0TwoUsers) {
|
||||
class PointsToAnalysisTestBase : public HloTestBase {
|
||||
protected:
|
||||
void BuildModule(std::unique_ptr<HloComputation> computation) {
|
||||
module_ = CreateNewModule();
|
||||
module_ = CreateNewUnverifiedModule();
|
||||
computation_ = module_->AddEntryComputation(std::move(computation));
|
||||
}
|
||||
|
||||
@ -1176,7 +1176,7 @@ TEST_F(CanShareOperandBufferWithUserTest, WhileCanShare) {
|
||||
return builder.Build();
|
||||
};
|
||||
|
||||
module_ = CreateNewModule();
|
||||
module_ = CreateNewUnverifiedModule();
|
||||
HloComputation* cond_computation =
|
||||
module_->AddEmbeddedComputation(make_cond());
|
||||
HloComputation* body_computation =
|
||||
@ -1211,7 +1211,7 @@ TEST_F(CanShareOperandBufferWithUserTest, CallToComputationWithFusionRoot) {
|
||||
auto add = sub_builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(shape, HloOpcode::kAdd, sub_param, ones));
|
||||
|
||||
module_ = CreateNewModule();
|
||||
module_ = CreateNewUnverifiedModule();
|
||||
auto sub_computation = module_->AddEmbeddedComputation(sub_builder.Build());
|
||||
sub_computation->CreateFusionInstruction({add, ones},
|
||||
HloInstruction::FusionKind::kLoop);
|
||||
|
||||
@ -45,7 +45,8 @@ class ZeroSizedHloEliminationTest : public HloTestBase {
|
||||
0, ShapeUtil::MakeShape(F32, {3, 0}), "zero sized param"))) {}
|
||||
|
||||
StatusOr<bool> RunZeroSizedElimination() {
|
||||
auto module = CreateNewModule("zero_sized_elimination_test_module");
|
||||
auto module =
|
||||
CreateNewUnverifiedModule("zero_sized_elimination_test_module");
|
||||
module->AddEntryComputation(builder_.Build());
|
||||
return ZeroSizedHloElimination{}.Run(module.get());
|
||||
}
|
||||
|
||||
@ -42,7 +42,7 @@ XLA_TEST_F(BroadcastTest, BroadcastScalarToScalar) {
|
||||
ShapeUtil::MakeShape(F32, {}), input, {}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -58,7 +58,7 @@ XLA_TEST_F(BroadcastTest, BroadcastScalarTo2D) {
|
||||
ShapeUtil::MakeShape(F32, {2, 2}), input, {}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -81,7 +81,7 @@ XLA_TEST_F(BroadcastTest, BroadcastVectorTo2D) {
|
||||
builder.AddInstruction(HloInstruction::CreateTuple({element1, element2}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -102,7 +102,7 @@ XLA_TEST_F(BroadcastTest, Broadcast2DTo2D) {
|
||||
ShapeUtil::MakeShape(F32, {2, 2}), input, {0, 1}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -121,7 +121,7 @@ XLA_TEST_F(BroadcastTest, Broadcast2DTo2DTranspose) {
|
||||
ShapeUtil::MakeShape(F32, {2, 2}), input, {1, 0}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -138,7 +138,7 @@ XLA_TEST_F(BroadcastTest, Broadcast2DTo3D) {
|
||||
ShapeUtil::MakeShape(F32, {2, 3, 2}), input, {0, 2}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -158,7 +158,7 @@ TEST_F(BroadcastTest, Broadcast_R1_2_To_R4_2x2x3x3) {
|
||||
ShapeUtil::MakeShape(F32, {2, 2, 3, 3}), input, {1}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -183,7 +183,7 @@ TEST_F(BroadcastTest, Broadcast_R1_1025_To_R4_3x3x3x1025) {
|
||||
ShapeUtil::MakeShape(F32, {3, 3, 3, r1_size}), input, {3}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -214,7 +214,7 @@ XLA_TEST_F(BroadcastTest, Broadcast_R1_64_To_R4_32x64x7x7) {
|
||||
ShapeUtil::MakeShape(F32, {32, 64, 7, 7}), input, {1}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -230,7 +230,7 @@ TEST_F(BroadcastTest, Broadcast_R0_to_R4_64x64x3x3) {
|
||||
ShapeUtil::MakeShape(F32, {64, 64, 3, 3}), input, {}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
LOG(INFO) << hlo_module->ToString();
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
@ -253,7 +253,7 @@ TEST_F(BroadcastTest, Broadcast_R2_2x2_To_R4_3x3x2x2) {
|
||||
ShapeUtil::MakeShape(F32, {3, 3, 2, 2}), input, {2, 3}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
@ -287,7 +287,7 @@ TEST_F(BroadcastTest, Broadcast_R3_2x3x4_to_R4_2x3x4x5) {
|
||||
ShapeUtil::MakeShape(F32, {2, 3, 4, 5}), input, {0, 1, 2}));
|
||||
|
||||
// Create HLO module, compile, and execute.
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
auto result = ExecuteAndTransfer(std::move(hlo_module), {});
|
||||
|
||||
|
||||
@ -45,7 +45,7 @@ class CopyOpTest : public HloTestBase {
|
||||
builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
constant->shape(), HloOpcode::kCopy, constant));
|
||||
auto computation = builder.Build();
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(std::move(computation));
|
||||
|
||||
Literal result = ExecuteAndTransfer(std::move(module), {});
|
||||
@ -98,7 +98,7 @@ XLA_TEST_F(CopyOpTest, CopyParameterScalar) {
|
||||
|
||||
auto computation = builder.Build();
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(std::move(computation));
|
||||
|
||||
Literal result = ExecuteAndTransfer(std::move(module), {&literal});
|
||||
@ -119,7 +119,7 @@ XLA_TEST_F(CopyOpTest, CopyConstantR2Twice) {
|
||||
|
||||
auto computation = builder.Build();
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(std::move(computation));
|
||||
Literal result = ExecuteAndTransfer(std::move(module), {});
|
||||
LiteralTestUtil::ExpectR2Near<float>({{1.0, 2.0}, {3.0, 4.0}}, result,
|
||||
@ -143,7 +143,7 @@ XLA_TEST_F(CopyOpTest, CopyConstantR2DifferentLayouts) {
|
||||
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(std::move(computation));
|
||||
Literal result = ExecuteAndTransfer(std::move(module), {});
|
||||
|
||||
@ -175,7 +175,7 @@ void CopyOpTest::TestCopyConstantLayout021(size_t n1, size_t n2, size_t n3) {
|
||||
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(std::move(computation));
|
||||
ForceResultLayout(module.get(), LayoutUtil::MakeLayout({1, 2, 0}));
|
||||
Literal result = ExecuteAndTransfer(std::move(module), {});
|
||||
@ -209,7 +209,7 @@ void CopyOpTest::TestCopyConstantLayoutR4(size_t n1, size_t n2, size_t n3,
|
||||
|
||||
std::unique_ptr<HloComputation> computation = builder.Build();
|
||||
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
module->AddEntryComputation(std::move(computation));
|
||||
ForceResultLayout(module.get(), LayoutUtil::MakeLayout(permutation));
|
||||
Literal result = ExecuteAndTransfer(std::move(module), {});
|
||||
|
||||
@ -70,7 +70,7 @@ class CustomCallTest : public HloTestBase {
|
||||
};
|
||||
|
||||
XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(CustomCallR0F32Add2)) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
|
||||
auto constant = builder.AddInstruction(
|
||||
@ -85,7 +85,7 @@ XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(CustomCallR0F32Add2)) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(CustomCallR2F32Reduce)) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
|
||||
Array2D<float> array(2, 2);
|
||||
@ -106,7 +106,7 @@ XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(CustomCallR2F32Reduce)) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(UsedInOtherComputations)) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto b = HloComputation::Builder(TestName());
|
||||
|
||||
auto input = b.AddInstruction(
|
||||
@ -130,7 +130,7 @@ XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(UsedInOtherComputations)) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(InputAndOutputLayoutDiffer)) {
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto b = HloComputation::Builder(TestName());
|
||||
|
||||
auto input =
|
||||
@ -155,7 +155,7 @@ XLA_TEST_F(CustomCallTest, DISABLED_ON_GPU(LayoutConstrained)) {
|
||||
// The argument and result of the computation are set to different layouts,
|
||||
// but the custom call is layout constrained to a fixed operand and result
|
||||
// layout, so the correct result should be produced.
|
||||
auto module = CreateNewModule();
|
||||
auto module = CreateNewUnverifiedModule();
|
||||
auto b = HloComputation::Builder(TestName());
|
||||
|
||||
auto input =
|
||||
|
||||
@ -81,7 +81,7 @@ class FusionTest : public HloTestBase {
|
||||
}
|
||||
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
|
||||
auto prim_type = primitive_util::NativeToPrimitiveType<T>();
|
||||
|
||||
@ -183,7 +183,7 @@ XLA_TEST_F(FusionTest, Test) {
|
||||
// (-{{1.0, 1.0, 1.0}, {0.0, 0.0, 0.0}}),
|
||||
// {{0.5, 0.5, 0.5}, {0.5, 0.5, 0.5}})) = {{0.5}, {2.72}}
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<float>({{1.0}, {2.0}, {3.0}})));
|
||||
auto const1 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
@ -231,7 +231,7 @@ XLA_TEST_F(FusionTest, Parameter) {
|
||||
// Build a computation and fuse part of it so the fusion instruction has an
|
||||
// operand parameter.
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<float>({{1.0, 2.0, 3.0}})));
|
||||
auto copy1 = builder.AddInstruction(HloInstruction::CreateUnary(
|
||||
@ -266,7 +266,7 @@ XLA_TEST_F(FusionTest, RandomizedParallelPartition) {
|
||||
ShapeUtil::MakeShapeWithLayout(F32, {rand_dim0_size, dim1_size}, {1, 0});
|
||||
// Build simple fusion computation: y = x^2 (elementwise).
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
|
||||
auto two = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(2.0)));
|
||||
@ -290,7 +290,7 @@ XLA_TEST_F(FusionTest, RandomizedParallelPartition) {
|
||||
|
||||
XLA_TEST_F(FusionTest, BroadcastIntoBinaryOp) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const_vector = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR1<float>({1.0, 2.0, 3.0})));
|
||||
auto const_array = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
@ -314,7 +314,7 @@ XLA_TEST_F(FusionTest, BroadcastIntoBinaryOp) {
|
||||
|
||||
XLA_TEST_F(FusionTest, ReshapeToScalar) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto single_element_array = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR2<int32>({{5}})));
|
||||
auto reshape = builder.AddInstruction(HloInstruction::CreateReshape(
|
||||
@ -329,7 +329,7 @@ XLA_TEST_F(FusionTest, ReshapeToScalar) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Reshape_3by2_1by2by3) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<int32>({{1, 2}, {3, 4}, {5, 6}})));
|
||||
auto reshape1 = builder.AddInstruction(HloInstruction::CreateReshape(
|
||||
@ -344,7 +344,7 @@ XLA_TEST_F(FusionTest, Reshape_3by2_1by2by3) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Reshape_1by2by3_3by2) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR3<int32>({{{1, 2, 3}, {4, 5, 6}}})));
|
||||
auto reshape1 = builder.AddInstruction(
|
||||
@ -359,7 +359,7 @@ XLA_TEST_F(FusionTest, Reshape_1by2by3_3by2) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Reshape_1by1by1_) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR3<int32>({{{7}}})));
|
||||
auto reshape1 = builder.AddInstruction(
|
||||
@ -374,7 +374,7 @@ XLA_TEST_F(FusionTest, Reshape_1by1by1_) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Reshape__1by1by1) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<int32>(7)));
|
||||
auto reshape1 = builder.AddInstruction(HloInstruction::CreateReshape(
|
||||
@ -389,7 +389,7 @@ XLA_TEST_F(FusionTest, Reshape__1by1by1) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Reshape__) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<int32>(7)));
|
||||
auto reshape1 = builder.AddInstruction(
|
||||
@ -404,7 +404,7 @@ XLA_TEST_F(FusionTest, Reshape__) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Reshape_3by3_3by3) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<int32>({{1, 2, 3}, {4, 5, 6}, {7, 8, 9}})));
|
||||
auto reshape1 = builder.AddInstruction(
|
||||
@ -419,7 +419,7 @@ XLA_TEST_F(FusionTest, Reshape_3by3_3by3) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Transpose_2by3) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<int32>({{1, 2, 3}, {4, 5, 6}})));
|
||||
auto reshape1 = builder.AddInstruction(HloInstruction::CreateTranspose(
|
||||
@ -434,7 +434,7 @@ XLA_TEST_F(FusionTest, Transpose_2by3) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Transpose_3by3) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<int32>({{1, 2, 3}, {4, 5, 6}, {7, 8, 9}})));
|
||||
auto reshape1 = builder.AddInstruction(HloInstruction::CreateTranspose(
|
||||
@ -449,7 +449,7 @@ XLA_TEST_F(FusionTest, Transpose_3by3) {
|
||||
|
||||
XLA_TEST_F(FusionTest, Reverse) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR1<int32>({1, 2, 3})));
|
||||
auto reverse1 = builder.AddInstruction(HloInstruction::CreateReverse(
|
||||
@ -465,7 +465,7 @@ XLA_TEST_F(FusionTest, Reverse) {
|
||||
|
||||
XLA_TEST_F(FusionTest, ReverseNegate) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR1<int32>({1, 2, 3})));
|
||||
auto reverse1 = builder.AddInstruction(HloInstruction::CreateReverse(
|
||||
@ -483,7 +483,7 @@ XLA_TEST_F(FusionTest, ReverseNegate) {
|
||||
|
||||
XLA_TEST_F(FusionTest, BroadcastNegate) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<int32>(1)));
|
||||
auto broadcast1 = builder.AddInstruction(HloInstruction::CreateBroadcast(
|
||||
@ -501,7 +501,7 @@ XLA_TEST_F(FusionTest, BroadcastNegate) {
|
||||
|
||||
XLA_TEST_F(FusionTest, SliceNegate) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR1<int32>({1, 2, 3, 4})));
|
||||
auto slice1 = builder.AddInstruction(HloInstruction::CreateSlice(
|
||||
@ -519,7 +519,7 @@ XLA_TEST_F(FusionTest, SliceNegate) {
|
||||
|
||||
XLA_TEST_F(FusionTest, DynamicSliceNegate) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR1<int32>({1, 2, 3, 4})));
|
||||
auto const1 = builder.AddInstruction(
|
||||
@ -541,7 +541,7 @@ XLA_TEST_F(FusionTest, DynamicSliceNegate) {
|
||||
|
||||
XLA_TEST_F(FusionTest, ReshapeNegate) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR1<int32>({1, 2, 3, 4})));
|
||||
auto reshape1 = builder.AddInstruction(
|
||||
@ -559,7 +559,7 @@ XLA_TEST_F(FusionTest, ReshapeNegate) {
|
||||
|
||||
XLA_TEST_F(FusionTest, TransposeNegate) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<int32>({{1, 2}, {3, 4}})));
|
||||
auto transpose1 = builder.AddInstruction(HloInstruction::CreateTranspose(
|
||||
@ -587,7 +587,7 @@ std::unique_ptr<HloComputation> MakeReduceTestComputation() {
|
||||
}
|
||||
|
||||
XLA_TEST_F(FusionTest, DISABLED_ON_CPU(Reduce)) {
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
@ -607,7 +607,7 @@ XLA_TEST_F(FusionTest, DISABLED_ON_CPU(Reduce)) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(FusionTest, DISABLED_ON_CPU(ReduceImplicitBroadcast)) {
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
@ -630,7 +630,7 @@ XLA_TEST_F(FusionTest, DISABLED_ON_CPU(ReduceImplicitBroadcast)) {
|
||||
|
||||
XLA_TEST_F(FusionTest, DISABLED_ON_CPU(ReduceWindow)) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
auto const0 = builder.AddInstruction(HloInstruction::CreateConstant(
|
||||
LiteralUtil::CreateR2<int32>({{2, 3, 5}, {7, 11, 13}, {17, 19, 23}})));
|
||||
auto const1 = builder.AddInstruction(
|
||||
@ -682,7 +682,7 @@ XLA_TEST_F(FusionTest, DISABLED_ON_CPU(ReduceWindow)) {
|
||||
// into a fusion, it should remain shared, rather than being duplicated
|
||||
// within the fusion.
|
||||
XLA_TEST_F(FusionTest, SharedConstant) {
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto const0 = builder.AddInstruction(
|
||||
|
||||
@ -107,7 +107,8 @@ HloTestBase::HloTestBase(se::Platform* test_platform,
|
||||
instruction_can_change_layout_func);
|
||||
}
|
||||
|
||||
std::unique_ptr<HloModule> HloTestBase::CreateNewModule(const string& name) {
|
||||
std::unique_ptr<HloModule> HloTestBase::CreateNewUnverifiedModule(
|
||||
const string& name) {
|
||||
return absl::make_unique<HloModule>(name, GetModuleConfigForTest());
|
||||
}
|
||||
|
||||
|
||||
@ -72,7 +72,12 @@ class HloTestBase : public ::testing::Test {
|
||||
// options from command-line flags. If you want a fresh HloModule object and
|
||||
// then add HloComputations to it, it's recommended to use this method in your
|
||||
// tests.
|
||||
std::unique_ptr<HloModule> CreateNewModule(const string& name = TestName());
|
||||
//
|
||||
// This returns a vanilla HloModule that doesn't run the HLO verifier on
|
||||
// destruction. If you want to run the verifier, you want
|
||||
// HloVerifiedTestBase::CreateNewVerifiedModule.
|
||||
std::unique_ptr<HloModule> CreateNewUnverifiedModule(
|
||||
const string& name = TestName());
|
||||
|
||||
// Runs the hlo_pass with the provided module and returns the result. This
|
||||
// function also verifies that the module remains unchanged when hlo_pass
|
||||
|
||||
@ -69,12 +69,12 @@ class HloVerifiedTestBase : public HloTestBase {
|
||||
std::unique_ptr<VerifiedHloModule> CreateNewVerifiedModule(
|
||||
const string& name = TestName());
|
||||
|
||||
// CreateNewModule creates an *unverified* module, which presumably isn't what
|
||||
// you want if you're using HloVerifiedTestBase, so we delete this function to
|
||||
// keep you from accidentally calling it. If you really want it, you can get
|
||||
// it by calling HloTestBase::CreateNewModule().
|
||||
std::unique_ptr<HloModule> CreateNewModule(const string& name = TestName()) =
|
||||
delete;
|
||||
// CreateNewUnverifiedModule creates an *unverified* module, which presumably
|
||||
// isn't what you want if you're using HloVerifiedTestBase, so we delete this
|
||||
// function to keep you from accidentally calling it. If you really want it,
|
||||
// you can get it by calling HloTestBase::CreateNewUnverifiedModule().
|
||||
std::unique_ptr<HloModule> CreateNewUnverifiedModule(
|
||||
const string& name = TestName()) = delete;
|
||||
|
||||
private:
|
||||
bool verifier_layout_sensitive_;
|
||||
|
||||
@ -38,8 +38,8 @@ XLA_TEST_F(HloVerifiedTestBaseTest, NoModule) {
|
||||
// Test shouldn't fail if no module is created at all.
|
||||
}
|
||||
|
||||
XLA_TEST_F(HloVerifiedTestBaseTest, GoodCreateNewModule) {
|
||||
// Call CreateNewModule and build up a valid module.
|
||||
XLA_TEST_F(HloVerifiedTestBaseTest, GoodCreateNewUnverifiedModule) {
|
||||
// Call CreateNewUnverifiedModule and build up a valid module.
|
||||
auto module = CreateNewVerifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto input = builder.AddInstruction(
|
||||
@ -50,8 +50,8 @@ XLA_TEST_F(HloVerifiedTestBaseTest, GoodCreateNewModule) {
|
||||
}
|
||||
|
||||
// This test is expected to fail. See test class comment.
|
||||
XLA_TEST_F(HloVerifiedTestBaseTest, DISABLED_BadCreateNewModule) {
|
||||
// Call CreateNewModule and build up a invalid module.
|
||||
XLA_TEST_F(HloVerifiedTestBaseTest, DISABLED_BadCreateNewUnverifiedModule) {
|
||||
// Call CreateNewUnverifiedModule and build up a invalid module.
|
||||
auto module = CreateNewVerifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto input = builder.AddInstruction(
|
||||
|
||||
@ -68,7 +68,7 @@ class LLVMCompilerTest : public ::testing::Test {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(42.0)));
|
||||
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
|
||||
compiler->SetPreOptimizationHook(pre_opt_hook);
|
||||
@ -90,7 +90,7 @@ class LLVMCompilerTest : public ::testing::Test {
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(42.0)));
|
||||
|
||||
std::unique_ptr<HloModule> hlo_module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> hlo_module = CreateNewUnverifiedModule();
|
||||
hlo_module->AddEntryComputation(builder.Build());
|
||||
|
||||
auto module_group = absl::make_unique<HloModuleGroup>("test_module_group");
|
||||
@ -124,7 +124,7 @@ class LLVMCompilerTest : public ::testing::Test {
|
||||
return ::testing::UnitTest::GetInstance()->current_test_info()->name();
|
||||
}
|
||||
|
||||
static std::unique_ptr<HloModule> CreateNewModule() {
|
||||
static std::unique_ptr<HloModule> CreateNewUnverifiedModule() {
|
||||
HloModuleConfig config;
|
||||
config.set_debug_options(GetDebugOptionsFromFlags());
|
||||
return absl::make_unique<HloModule>(TestName(), config);
|
||||
|
||||
@ -62,7 +62,7 @@ class MultiOutputFusionTest : public HloTestBase {
|
||||
|
||||
void RunTest2D(bool manual_fusion, int64 size) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
|
||||
const Shape elem_shape0 = ShapeUtil::MakeShapeWithLayout(F32, {}, {});
|
||||
const Shape elem_shape2 =
|
||||
@ -122,7 +122,7 @@ class MultiOutputFusionTest : public HloTestBase {
|
||||
|
||||
void RunTest1D(bool manual_fusion, int size) {
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto hlo_module = CreateNewModule();
|
||||
auto hlo_module = CreateNewUnverifiedModule();
|
||||
|
||||
const Shape elem_shape_F32 =
|
||||
ShapeUtil::MakeShapeWithDescendingLayout(F32, {size});
|
||||
|
||||
@ -28,7 +28,7 @@ namespace {
|
||||
class TokenHloTest : public HloTestBase {};
|
||||
|
||||
XLA_TEST_F(TokenHloTest, SingleTokenInstruction) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
builder.AddInstruction(HloInstruction::CreateToken());
|
||||
|
||||
@ -39,7 +39,7 @@ XLA_TEST_F(TokenHloTest, SingleTokenInstruction) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(TokenHloTest, TokenTree) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto token0 = builder.AddInstruction(HloInstruction::CreateToken());
|
||||
auto token1 = builder.AddInstruction(HloInstruction::CreateToken());
|
||||
@ -54,7 +54,7 @@ XLA_TEST_F(TokenHloTest, TokenTree) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(TokenHloTest, InvalidTokenShapedEntryParameter) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(0, ShapeUtil::MakeShape(F32, {}), "p0"));
|
||||
@ -75,7 +75,7 @@ XLA_TEST_F(TokenHloTest, InvalidTokenShapedEntryParameter) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(TokenHloTest, InvalidTupleTokenShapedEntryParameter) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
builder.AddInstruction(HloInstruction::CreateParameter(
|
||||
0,
|
||||
@ -95,7 +95,7 @@ XLA_TEST_F(TokenHloTest, InvalidTupleTokenShapedEntryParameter) {
|
||||
}
|
||||
|
||||
XLA_TEST_F(TokenHloTest, InvalidOperandToTokenInstruction) {
|
||||
std::unique_ptr<HloModule> module = CreateNewModule();
|
||||
std::unique_ptr<HloModule> module = CreateNewUnverifiedModule();
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
auto param = builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(0, ShapeUtil::MakeShape(F32, {}), "p0"));
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user