@@ -7998,31 +7998,32 @@ ENTRY entry {
7998
7998
EXPECT_TRUE(repacker_ran);
7999
7999
}
8000
8000
8001
- TEST_F(MemorySpaceAssignmentTest, ReduceReservedScopedVmemIfOperandInVmem) {
8002
- // This test is designed to test UpdateReservedScopedVmemSize() in MSA, which
8003
- // will invoke reserved_scoped_memory_fn to update scoped allocation
8004
- // size. UpdateReservedScopedVmemSize() should iterate through all scheduled
8005
- // instruction and check if either their operands or outputs has been assigned
8006
- // in alternate memory. If so, corresponding operand/output will be passed to
8007
- // reserved_scoped_memory_fn. We isolate UpdateReservedScopedVmemSize() by
8008
- // constructing a dummy reserved_scoped_memory_fn that return +1 when operand
8009
- // set is empty, and return +2 when output set is empty, because if either set
8010
- // of an instruction is empty, it is gureented that some scoped allocation is
8011
- // required. We use +1/+2 to distinguish the correctness of each set.
8012
- // Therefore, after MSA pass, for each instruction, there are a few possible
8013
- // outcomes:
8001
+ TEST_F(MemorySpaceAssignmentTest,
8002
+ ReduceReservedScopedAllocationIfOperandInAlternateMemory) {
8003
+ // This test is designed to test UpdateReservedScopedAllocationSize() in MSA,
8004
+ // which will invoke reserved_scoped_memory_fn to update scoped allocation
8005
+ // size. UpdateReservedScopedAllocationSize() should iterate through all
8006
+ // scheduled instruction and check if either their operands or outputs has
8007
+ // been assigned in alternate memory. If so, corresponding operand/output will
8008
+ // be passed to reserved_scoped_memory_fn. We isolate
8009
+ // UpdateReservedScopedAllocationSize() by constructing a dummy
8010
+ // reserved_scoped_memory_fn that return +1 when operand set is empty, and
8011
+ // return +2 when output set is empty, because if either set of an instruction
8012
+ // is empty, it is gureented that some scoped allocation is required. We use
8013
+ // +1/+2 to distinguish the correctness of each set. Therefore, after MSA
8014
+ // pass, for each instruction, there are a few possible outcomes:
8014
8015
// 1. If both operand set and output set are not empty, scoped allocation
8015
8016
// size should be 0, since reserved_scoped_memory_fn will return 0.
8016
8017
// 2. If only operand set is empty, scoped allocation size should be 2, since
8017
8018
// reserved_scoped_memory_fn will return 2.
8018
8019
// 3. If only output set is empty, scoped allocation size should be 1, since
8019
8020
// reserved_scoped_memory_fn will return 1.
8020
8021
// 4. If both sets are empty, scoped allocation size should be 3.
8021
- // Initially, UpdateReservedScopedVmemSize () will only be invoked after each
8022
- // MSA repacking, we use a similar test HLO module as used in "Repack" test.
8023
- // This test is capable of testing if UpdateReservedScopedVmemSize() can
8024
- // correctly pass operand/output set of all instructions to
8025
- // reserved_scoped_memory_fn.
8022
+ // Initially, UpdateReservedScopedAllocationSize () will only be invoked after
8023
+ // each MSA repacking, we use a similar test HLO module as used in "Repack"
8024
+ // test. This test is capable of testing if
8025
+ // UpdateReservedScopedAllocationSize() can correctly pass operand/output set
8026
+ // of all instructions to reserved_scoped_memory_fn.
8026
8027
absl::string_view hlo_string = R"(
8027
8028
HloModule bug, is_scheduled=true
8028
8029
@@ -8059,11 +8060,13 @@ TEST_F(MemorySpaceAssignmentTest, ReduceReservedScopedVmemIfOperandInVmem) {
8059
8060
int64_t scoped_memory_size = 0;
8060
8061
if (operands_in_alternate_memory.empty()) {
8061
8062
scoped_memory_size += 1;
8062
- LOG(INFO) << instruction->name() << " has no operand in vmem";
8063
+ LOG(INFO) << instruction->name()
8064
+ << " has no operand in alternate memory";
8063
8065
}
8064
8066
if (outputs_in_alternate_memory.empty()) {
8065
8067
scoped_memory_size += 2;
8066
- LOG(INFO) << instruction->name() << " has no output in vmem";
8068
+ LOG(INFO) << instruction->name()
8069
+ << " has no output in alternate memory";
8067
8070
}
8068
8071
return scoped_memory_size;
8069
8072
};
@@ -8196,7 +8199,7 @@ TEST_F(MemorySpaceAssignmentTest, ScopedAllocationWithDifferentOffset) {
8196
8199
}
8197
8200
8198
8201
TEST_F(MemorySpaceAssignmentTest,
8199
- ReduceReservedScopedVmemUpdatesPeakMemoryUsage ) {
8202
+ ReduceReservedScopedAllocationUpdatesPeakMemoryUsage ) {
8200
8203
// This test is designed to test that the peak_memory_usage_ is updated
8201
8204
// correctly after scoped memory allocation is updated. The test HLO module
8202
8205
// has two HLO values: a and b. The size of a is 64, and the size of b is 128.
0 commit comments