| 448 | public ::testing::WithParamInterface<bool> {}; |
| 449 | |
| 450 | TEST_P(IndirectUseTest, IndirectUseNotRematerialized) { |
| 451 | // Test that an rematerializable instruction is not rematerialized if it has |
| 452 | // an indirect use. Test is parameterized on whether the value has an indirect |
| 453 | // use, and the instruction should be rematerialized iff the value has no |
| 454 | // indirect use. Module: |
| 455 | // |
| 456 | // Entry computation: |
| 457 | // F32[] %param = {...} |
| 458 | // F32[1024] %bcast = broadcast(%param) |
| 459 | // F32[1024] %add_1 = add(%bcast, bcast) |
| 460 | // F32[1024] %call = call(Subcomputation, {%add_1}) |
| 461 | // F32[1024] %add_2 = add(%bcast, call) |
| 462 | // {F32[1024], F32[1024]} %tuple = tuple(%bcast, %add_2) |
| 463 | // F32[1024] %gte = GetTupleElement(%tuple, 0) |
| 464 | // F32[1024] %negate = negate(%gte) |
| 465 | // |
| 466 | // Subcomputation: |
| 467 | // F32[1024] %param = {...} |
| 468 | // F32[2048] %concat = concat({%param, %param}) |
| 469 | // F32[1024] %slice = slice(%concat) |
| 470 | // |
| 471 | // The value %bcast is live across the call and rematerialization of %bcast |
| 472 | // across that point would reduce peak memory use by 4KB. However, %bcast is |
| 473 | // used indirectly in the %negate so rematerialization should not happen. |
| 474 | // |
| 475 | // This test is parameterized on whether the broadcast has an indirect use or |
| 476 | // not. The indirect use is controlled by the index of the GetTupleElement |
| 477 | // instruction. If the element is 0, then the %negate operand aliases %bcast |
| 478 | // (ie %bcast is used indirectly by %negate), otherwise the %negate operand |
| 479 | // aliases %add_2. |
| 480 | const bool indirectly_used = GetParam(); |
| 481 | auto module = CreateNewVerifiedModule(); |
| 482 | |
| 483 | HloComputation* subcomputation = nullptr; |
| 484 | { |
| 485 | auto builder = HloComputation::Builder(TestName() + ".subcomputation"); |
| 486 | auto param = builder.AddInstruction( |
| 487 | HloInstruction::CreateParameter(0, vec1024_shape_, "param")); |
| 488 | auto concat = builder.AddInstruction(HloInstruction::CreateConcatenate( |
| 489 | ShapeUtil::MakeShape(xla::F32, {2048}), {param, param}, |
| 490 | /*dimension=*/0)); |
| 491 | builder.AddInstruction(HloInstruction::CreateSlice( |
| 492 | vec1024_shape_, concat, /*start_indices=*/{0}, |
| 493 | /*limit_indices=*/{1024}, /*strides=*/{1})); |
| 494 | subcomputation = module->AddEmbeddedComputation(builder.Build()); |
| 495 | } |
| 496 | |
| 497 | auto builder = HloComputation::Builder(TestName()); |
| 498 | auto param = builder.AddInstruction( |
| 499 | HloInstruction::CreateParameter(0, scalar_shape_, "param")); |
| 500 | auto bcast = builder.AddInstruction( |
| 501 | HloInstruction::CreateBroadcast(vec1024_shape_, param, {})); |
| 502 | auto add_1 = builder.AddInstruction(HloInstruction::CreateBinary( |
| 503 | vec1024_shape_, HloOpcode::kAdd, bcast, bcast)); |
| 504 | auto call_1 = builder.AddInstruction( |
| 505 | HloInstruction::CreateCall(vec1024_shape_, {add_1}, subcomputation)); |
| 506 | auto add_2 = builder.AddInstruction(HloInstruction::CreateBinary( |
| 507 | vec1024_shape_, HloOpcode::kAdd, bcast, call_1)); |
nothing calls this directly
no test coverage detected