Rewrite CopyInsertion to use module-scoped HloAliasAnalysis. The net effect (number of copies inserted) is roughly similar to the existing implementation, but the new implementation is much more general. The new implementation can handle entry argument buffer reuse with minimal modification, for example.
Some unnecessary copies are still added due to deficiencies in buffer assignment (b/62548313), but these can be removed when buffer assignment also uses HloAliasAnalysis. Also address a few issues uncovered with this cl: (1) For inplace dynamic slice in llvm backends, truncate do not wrap the slice. This matches the behavior of the non-inplace variant. (2) Disable SelectBetweenPredTuples test on GPU. The test introduces top-level buffer ambiguity which is not tolerated by the gpu backend. (3) When deserializing HLO form a proto, do not uniquify instruction names in fused computations. (4) In dataflow analysis, don't deallocate deleted HloValues during propagation. (5) In dataflow analysis, fix issue with live_out_of_computation property. PiperOrigin-RevId: 174423881
This commit is contained in:
Mark Heffernan
TensorFlower Gardener
parent
8a7f5c47dc
commit
7bb2d57b0b
@ -1644,10 +1644,14 @@ cc_library(
|
||||
deps = [
|
||||
":buffer_liveness",
|
||||
":hlo",
|
||||
":hlo_alias_analysis",
|
||||
":hlo_dce",
|
||||
":hlo_graph_dumper",
|
||||
":hlo_ordering",
|
||||
":hlo_pass",
|
||||
":liveness_util",
|
||||
":logical_buffer",
|
||||
":tuple_points_to_analysis",
|
||||
":tuple_simplifier",
|
||||
"//tensorflow/compiler/xla:status_macros",
|
||||
"//tensorflow/compiler/xla:statusor",
|
||||
"//tensorflow/compiler/xla:types",
|
||||
@ -1662,15 +1666,17 @@ tf_cc_test(
|
||||
deps = [
|
||||
":copy_insertion",
|
||||
":hlo",
|
||||
":hlo_graph_dumper",
|
||||
":hlo_matchers",
|
||||
":tuple_points_to_analysis",
|
||||
"//tensorflow/compiler/xla:literal_util",
|
||||
"//tensorflow/compiler/xla:shape_util",
|
||||
"//tensorflow/compiler/xla:test",
|
||||
"//tensorflow/compiler/xla:test_helpers",
|
||||
"//tensorflow/compiler/xla:xla_data_proto",
|
||||
"//tensorflow/compiler/xla/legacy_flags:debug_options_flags",
|
||||
"//tensorflow/compiler/xla/tests:hlo_test_base",
|
||||
"//tensorflow/compiler/xla/tests:xla_internal_test_main",
|
||||
"//tensorflow/core:test",
|
||||
],
|
||||
)
|
||||
|
||||
|
||||
@ -1235,7 +1235,6 @@ const LogicalBuffer* AddBufferToColocatedSet(
|
||||
// CopyInsertion ensures root points-to set is unambiguous and distinct.
|
||||
const auto& points_to = points_to_analysis.GetPointsToSet(instruction);
|
||||
DCHECK(!points_to.IsAmbiguous());
|
||||
DCHECK(points_to.IsDistinct());
|
||||
colocated_set->push_back(points_to.element(index)[0]);
|
||||
return colocated_set->back();
|
||||
}
|
||||
|
||||
@ -1538,8 +1538,6 @@ TEST_F(WhileBufferAssignmentTest, OneForwardBackwardWhileLoopSet) {
|
||||
HloInstruction::CreateConstant(Literal::CreateR0<float>(0.0)));
|
||||
auto output0 = builder.AddInstruction(
|
||||
HloInstruction::CreateBroadcast(data_shape_, zero, {1}));
|
||||
auto output1 = builder.AddInstruction(
|
||||
HloInstruction::CreateBroadcast(data_shape_, zero, {1}));
|
||||
|
||||
auto cond0 =
|
||||
module->AddEmbeddedComputation(BuildWhileConditionComputation("cond"));
|
||||
@ -1556,10 +1554,8 @@ TEST_F(WhileBufferAssignmentTest, OneForwardBackwardWhileLoopSet) {
|
||||
auto body1 =
|
||||
module->AddEmbeddedComputation(BuildWhileBodyComputation("body"));
|
||||
|
||||
auto tuple1 = builder.AddInstruction(
|
||||
HloInstruction::CreateTuple({input0, weights0, output1}));
|
||||
auto while1 = builder.AddInstruction(
|
||||
HloInstruction::CreateWhile(loop_state_shape_, cond1, body1, tuple1));
|
||||
HloInstruction::CreateWhile(loop_state_shape_, cond1, body1, while0));
|
||||
|
||||
module->AddEntryComputation(builder.Build());
|
||||
RunCopyInsertion(module.get());
|
||||
@ -1676,11 +1672,14 @@ TEST_F(WhileBufferAssignmentTest, WhileLoopsInterferingResultRange) {
|
||||
auto while1 = builder.AddInstruction(
|
||||
HloInstruction::CreateWhile(loop_state_shape_, cond, body, tuple1));
|
||||
|
||||
auto gte0 = builder.AddInstruction(
|
||||
HloInstruction::CreateGetTupleElement(data_shape_, while0, 0));
|
||||
auto gte1 = builder.AddInstruction(
|
||||
HloInstruction::CreateGetTupleElement(data_shape_, while1, 1));
|
||||
auto root_add = builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
while0->shape(), HloOpcode::kAdd, while0, while1));
|
||||
module->AddEntryComputation(builder.Build());
|
||||
while0->shape(), HloOpcode::kAdd, gte0, gte1));
|
||||
|
||||
RunCopyInsertion(module.get());
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
{
|
||||
FlattenCallGraph flatten;
|
||||
@ -1688,22 +1687,22 @@ TEST_F(WhileBufferAssignmentTest, WhileLoopsInterferingResultRange) {
|
||||
EXPECT_TRUE(result);
|
||||
}
|
||||
|
||||
RunCopyInsertion(module.get());
|
||||
|
||||
auto sequence =
|
||||
CreateMemoryMinimizingSequence(*module, ByteSizeOf).ConsumeValueOrDie();
|
||||
|
||||
// To trigger b/38494731, we want a specific Hlo sequence for the
|
||||
// root computation, so we overwrite that entry with a manually
|
||||
// crafted sequence.
|
||||
std::vector<const HloInstruction*> sequence_for_buffer_assigment = {
|
||||
input1, weights1, one, output1, tuple1, while1, input0,
|
||||
weights0, zero, output0, tuple0, while0, root_add};
|
||||
sequence[module->entry_computation()] = {
|
||||
input1, weights1, one, output1, while1->operand(0), while1,
|
||||
input0, weights0, zero, output0, while0->operand(0), while0,
|
||||
gte0, gte1, root_add};
|
||||
|
||||
// If this ASSERT_TRUE fails, we constructed a bogus sequence above
|
||||
// and this test itself is buggy.
|
||||
ASSERT_TRUE(IsPostOrderTraversal(sequence_for_buffer_assigment));
|
||||
|
||||
sequence[module->entry_computation()] =
|
||||
std::move(sequence_for_buffer_assigment);
|
||||
ASSERT_TRUE(IsPostOrderTraversal(sequence[module->entry_computation()]));
|
||||
|
||||
auto assignment =
|
||||
BufferAssigner::Run(
|
||||
@ -1715,55 +1714,6 @@ TEST_F(WhileBufferAssignmentTest, WhileLoopsInterferingResultRange) {
|
||||
EXPECT_TRUE(BuffersDistinct({while0}, {while1}, *assignment));
|
||||
}
|
||||
|
||||
// Test buffer assignment for while nodes with multiple uses.
|
||||
// TODO(b/37245345): Fix buffer assignment for this case.
|
||||
TEST_F(WhileBufferAssignmentTest, DISABLED_TwoWhiles) {
|
||||
auto module = MakeUnique<HloModule>(TestName());
|
||||
auto builder = HloComputation::Builder(TestName());
|
||||
|
||||
auto input0 = builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(0, data_shape_, "input0"));
|
||||
auto weights0 = builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(1, data_shape_, "weights0"));
|
||||
|
||||
auto zero = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(Literal::CreateR0<float>(0.0)));
|
||||
auto output0 = builder.AddInstruction(
|
||||
HloInstruction::CreateBroadcast(data_shape_, zero, {1}));
|
||||
|
||||
auto cond0 =
|
||||
module->AddEmbeddedComputation(BuildWhileConditionComputation("cond"));
|
||||
auto body0 =
|
||||
module->AddEmbeddedComputation(BuildWhileBodyComputation("body"));
|
||||
|
||||
auto tuple0 = builder.AddInstruction(
|
||||
HloInstruction::CreateTuple({input0, weights0, output0}));
|
||||
auto while0 = builder.AddInstruction(
|
||||
HloInstruction::CreateWhile(loop_state_shape_, cond0, body0, tuple0));
|
||||
auto while1 = builder.AddInstruction(
|
||||
HloInstruction::CreateWhile(loop_state_shape_, cond0, body0, while0));
|
||||
|
||||
auto get0 = builder.AddInstruction(
|
||||
HloInstruction::CreateGetTupleElement(data_shape_, while0, 2));
|
||||
auto get1 = builder.AddInstruction(
|
||||
HloInstruction::CreateGetTupleElement(data_shape_, while1, 2));
|
||||
builder.AddInstruction(
|
||||
HloInstruction::CreateBinary(data_shape_, HloOpcode::kAdd, get0, get1));
|
||||
module->AddEntryComputation(builder.Build());
|
||||
|
||||
RunCopyInsertion(module.get());
|
||||
|
||||
{
|
||||
FlattenCallGraph flatten;
|
||||
TF_ASSERT_OK_AND_ASSIGN(bool result, flatten.Run(module.get()));
|
||||
EXPECT_TRUE(result);
|
||||
}
|
||||
|
||||
auto assignment = RunBufferAssignment(module.get());
|
||||
|
||||
EXPECT_TRUE(BuffersDistinct({while0}, {while1}, *assignment));
|
||||
}
|
||||
|
||||
TEST_F(WhileBufferAssignmentTest, WhilesDontShareEntryParamIfLiveOut) {
|
||||
auto module = MakeUnique<HloModule>(TestName());
|
||||
auto builder = HloComputation::Builder("entry");
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -25,12 +25,25 @@ limitations under the License.
|
||||
|
||||
namespace xla {
|
||||
|
||||
// HLO pass which inserts a copy of the root instruction (creating a new root)
|
||||
// if the root is or points-to any constant or parameter instruction.
|
||||
// If the root instruction is a Tuple, only tuple elements which point to
|
||||
// constant or parameter instructions will be copied.
|
||||
// Copy insertion is necessary because constant and parameter arrays have
|
||||
// different lifetimes than computation results.
|
||||
// Copy insertion is a legalization HLO pass which inserts copies (kCopy
|
||||
// instructions) to eliminate several kinds of problems in the HLO module.
|
||||
//
|
||||
// (1) Entry parameter or a constant live out of the entry computation. Entry
|
||||
// computation arguments and constants have different lifetimes than the
|
||||
// computation result and cannot share the same allocation. Parameters and
|
||||
// constants live out of non-entry computations do not need copies.
|
||||
//
|
||||
// (2) Different values which are simultaneously live and which must be held
|
||||
// in the same buffer. This can occur in while bodies. Specifically, the
|
||||
// while loop state (the arguments to the while instruction) is updated
|
||||
// in-place and the update may clobber the value from the previous
|
||||
// iteration before the previous value is dead. Computations called from
|
||||
// kCall instructions do not need such copies because kCall has no update
|
||||
// in-place semantics.
|
||||
//
|
||||
// (3) The buffer set of the root instruction of the entry computation must be
|
||||
// unambiguous and distinct. That is, InstructionAliasSet::IsAmbiguous and
|
||||
// InstructionAliasSet::IsDistinct return true.
|
||||
class CopyInsertion : public HloPassInterface {
|
||||
public:
|
||||
tensorflow::StringPiece name() const override { return "copy-insertion"; }
|
||||
@ -38,15 +51,6 @@ class CopyInsertion : public HloPassInterface {
|
||||
// Run the pass on the given module. Returns whether the module was changed
|
||||
// (copies were inserted).
|
||||
StatusOr<bool> Run(HloModule* module) override;
|
||||
|
||||
protected:
|
||||
// Returns a copy of `hlo`. Looks in inserted_copies_ first to avoid making
|
||||
// duplicate copies.
|
||||
StatusOr<HloInstruction*> FindOrInsertCopy(HloInstruction* hlo);
|
||||
|
||||
// A map containing all copies inserted during the copy insertion pass. The
|
||||
// key is the copied instruction and the value is the copy.
|
||||
tensorflow::gtl::FlatMap<HloInstruction*, HloInstruction*> inserted_copies_;
|
||||
};
|
||||
|
||||
} // namespace xla
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -243,6 +243,81 @@ class CollectProfileCandidates : public DfsHloVisitorWithDefault {
|
||||
|
||||
std::unordered_map<const HloInstruction*, size_t>* hlo_to_profile_idx_;
|
||||
};
|
||||
|
||||
// This copy insertion pass is a hack to address deficiencies in buffer
|
||||
// assignment. Buffer assignment uses TuplePointsToAnalysis which is
|
||||
// computation-scoped and thus has limited visibility across computation
|
||||
// boundaries. However, CopyInsertion uses module-scoped HloAliasAnalysis and
|
||||
// expects buffer assignment to have the same understanding of the graph. This
|
||||
// mismatch manifests in the parallel cpu backend, where the HLO outlining
|
||||
// results is a minefield of potential problems. This pass conservatively adds
|
||||
// copies to avoid any potential problems in buffer assignemnt.
|
||||
//
|
||||
// Technically these issues exist in all the backends. However, they only
|
||||
// manifest in the parallel cpu backend because of the outlining. Moving this
|
||||
// into the main copy insertion pass results in performance regressions n the
|
||||
// other backends.
|
||||
//
|
||||
// TODO(b/62548313): Remove this.
|
||||
class CpuParallelCopyInsertion : public HloPassInterface {
|
||||
public:
|
||||
tensorflow::StringPiece name() const override {
|
||||
return "cpu-parallel-copy-insertion";
|
||||
}
|
||||
|
||||
StatusOr<bool> Run(HloModule* module) override {
|
||||
// Copy roots of all non-entry sequentially-called (eg, kCall, kWhile)
|
||||
// computations.
|
||||
std::unique_ptr<CallGraph> call_graph = CallGraph::Build(module);
|
||||
TF_RETURN_IF_ERROR(
|
||||
call_graph->VisitNodes([module](const CallGraphNode& node) -> Status {
|
||||
if (node.context() == CallContext::kSequential &&
|
||||
!node.caller_callsites().empty()) {
|
||||
TF_ASSIGN_OR_RETURN(HloInstruction * root_copy,
|
||||
node.computation()->DeepCopyInstruction(
|
||||
node.computation()->root_instruction()));
|
||||
node.computation()->set_root_instruction(root_copy);
|
||||
}
|
||||
return Status::OK();
|
||||
}));
|
||||
|
||||
TF_ASSIGN_OR_RETURN(std::unique_ptr<HloDataflowAnalysis> dataflow,
|
||||
HloDataflowAnalysis::Run(module));
|
||||
|
||||
// Add copies to the operand of dynamic update slices which have read-only
|
||||
// values (constants and parameters). Buffer assignment which is based on
|
||||
// computation-scoped tuple points-to analysis does not properly track these
|
||||
// read-only values across kCall instructions. This can result in cases
|
||||
// where a outlined computation parameter operand of a dynamic update slice
|
||||
// aliases a constant or parameter in the entry computation and the dynamic
|
||||
// update slice is attempted in-place.
|
||||
for (HloComputation* computation : module->computations()) {
|
||||
for (HloInstruction* instruction : computation->instructions()) {
|
||||
if (instruction->opcode() == HloOpcode::kDynamicUpdateSlice) {
|
||||
HloInstruction* operand = instruction->mutable_operand(0);
|
||||
for (const HloValue* value :
|
||||
dataflow->GetValueSet(operand).values()) {
|
||||
if (value->defining_instruction()->opcode() ==
|
||||
HloOpcode::kConstant ||
|
||||
value->defining_instruction()->opcode() ==
|
||||
HloOpcode::kParameter) {
|
||||
HloInstruction* operand_copy =
|
||||
instruction->parent()->AddInstruction(
|
||||
HloInstruction::CreateUnary(operand->shape(),
|
||||
HloOpcode::kCopy, operand));
|
||||
TF_RETURN_IF_ERROR(
|
||||
operand->ReplaceUseWith(instruction, operand_copy));
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace
|
||||
|
||||
Status CpuCompiler::RunHloPasses(HloModule* module, bool is_aot_compile) {
|
||||
@ -331,15 +406,16 @@ Status CpuCompiler::RunHloPasses(HloModule* module, bool is_aot_compile) {
|
||||
// (and sometime after) copy insertion, to avoid dead code from interfering
|
||||
// with the rewrites.
|
||||
pipeline.AddPass<HloDCE>();
|
||||
pipeline.AddPass<FlattenCallGraph>();
|
||||
pipeline.AddPass<CopyInsertion>();
|
||||
if (options::CpuParallelBackendRequested(module->config())) {
|
||||
// Re-run the outlining, in case any copies were inserted into the entry
|
||||
// computation.
|
||||
pipeline.AddPass<ParallelizationPreparation>(max_parallelism,
|
||||
ShapeSizeBytesFunction());
|
||||
pipeline.AddPass<CpuParallelCopyInsertion>();
|
||||
}
|
||||
pipeline.AddPass<HloDCE>();
|
||||
pipeline.AddPass<FlattenCallGraph>();
|
||||
return pipeline.Run(module).status();
|
||||
}
|
||||
|
||||
|
||||
@ -350,8 +350,8 @@ cc_library(
|
||||
":ir_emission_utils",
|
||||
"//tensorflow/compiler/xla/service:copy_insertion",
|
||||
"//tensorflow/compiler/xla/service:hlo",
|
||||
"//tensorflow/compiler/xla/service:logical_buffer",
|
||||
"//tensorflow/compiler/xla/service:tuple_points_to_analysis",
|
||||
"//tensorflow/compiler/xla/service:hlo_dataflow_analysis",
|
||||
"//tensorflow/compiler/xla/service:hlo_pass",
|
||||
"//tensorflow/core:lib",
|
||||
],
|
||||
)
|
||||
@ -573,11 +573,14 @@ tf_cc_test(
|
||||
deps = [
|
||||
":instruction_fusion",
|
||||
":while_transformer",
|
||||
"//tensorflow/compiler/xla:shape_util",
|
||||
"//tensorflow/compiler/xla:test",
|
||||
"//tensorflow/compiler/xla:test_helpers",
|
||||
"//tensorflow/compiler/xla/service:copy_insertion",
|
||||
"//tensorflow/compiler/xla/service:hlo_verifier",
|
||||
"//tensorflow/compiler/xla/tests:hlo_test_base",
|
||||
"//tensorflow/compiler/xla/tests:xla_internal_test_main",
|
||||
"//tensorflow/core:test",
|
||||
],
|
||||
)
|
||||
|
||||
|
||||
@ -22,41 +22,53 @@ limitations under the License.
|
||||
#include "tensorflow/compiler/xla/service/copy_insertion.h"
|
||||
#include "tensorflow/compiler/xla/service/gpu/ir_emission_utils.h"
|
||||
#include "tensorflow/compiler/xla/service/hlo_computation.h"
|
||||
#include "tensorflow/compiler/xla/service/hlo_dataflow_analysis.h"
|
||||
#include "tensorflow/compiler/xla/service/hlo_instruction.h"
|
||||
#include "tensorflow/compiler/xla/service/hlo_opcode.h"
|
||||
#include "tensorflow/compiler/xla/service/logical_buffer.h"
|
||||
#include "tensorflow/compiler/xla/service/tuple_points_to_analysis.h"
|
||||
#include "tensorflow/core/lib/core/status.h"
|
||||
#include "tensorflow/core/lib/gtl/flatset.h"
|
||||
#include "tensorflow/core/platform/logging.h"
|
||||
|
||||
namespace xla {
|
||||
namespace gpu {
|
||||
|
||||
StatusOr<bool> GpuCopyInsertion::Run(HloModule* module) {
|
||||
TF_ASSIGN_OR_RETURN(bool changed, CopyInsertion::Run(module));
|
||||
StatusOr<HloInstruction*> GpuCopyInsertion::FindOrInsertCopy(
|
||||
HloInstruction* hlo) {
|
||||
auto copy_it = inserted_copies_.find(hlo);
|
||||
if (copy_it == inserted_copies_.end()) {
|
||||
HloInstruction* copy = hlo->parent()->DeepCopyInstruction(hlo).ValueOrDie();
|
||||
inserted_copies_.insert({hlo, copy});
|
||||
return copy;
|
||||
} else {
|
||||
return copy_it->second;
|
||||
}
|
||||
}
|
||||
|
||||
TF_ASSIGN_OR_RETURN(auto points_to_analysis,
|
||||
TuplePointsToAnalysis::Run(module));
|
||||
StatusOr<bool> GpuCopyInsertion::Run(HloModule* module) {
|
||||
CopyInsertion generic_copy_insertion;
|
||||
|
||||
TF_ASSIGN_OR_RETURN(bool changed, generic_copy_insertion.Run(module));
|
||||
|
||||
TF_ASSIGN_OR_RETURN(std::unique_ptr<HloDataflowAnalysis> dataflow,
|
||||
HloDataflowAnalysis::Run(module));
|
||||
|
||||
// Make sure all operands of a library call are in memory instead of constants
|
||||
// in IR. The top-level (index {}) of the points-to set of each operand
|
||||
// indicates the source(s) of the array buffer. If any of these are constant,
|
||||
// then add a copy to materialize the array.
|
||||
// in IR.
|
||||
HloComputation* computation = module->entry_computation();
|
||||
for (HloInstruction* hlo : computation->MakeInstructionPostOrder()) {
|
||||
if (ImplementedAsLibraryCall(*hlo)) {
|
||||
for (int64 i = 0; i < hlo->operand_count(); ++i) {
|
||||
HloInstruction* operand = hlo->mutable_operand(i);
|
||||
const PointsToSet& points_to =
|
||||
points_to_analysis->GetPointsToSet(operand);
|
||||
const auto& element = points_to.element(/*index=*/{});
|
||||
if (std::any_of(element.begin(), element.end(),
|
||||
[](const LogicalBuffer* buffer_source) {
|
||||
return buffer_source->instruction()->opcode() ==
|
||||
HloOpcode::kConstant;
|
||||
})) {
|
||||
TF_ASSIGN_OR_RETURN(HloInstruction * copy,
|
||||
CopyInsertion::FindOrInsertCopy(operand));
|
||||
TF_RET_CHECK(ShapeUtil::IsArray(operand->shape()));
|
||||
bool copy_operand = false;
|
||||
for (const HloValue* value : dataflow->GetValueSet(operand).values()) {
|
||||
if (value->defining_instruction()->opcode() == HloOpcode::kConstant) {
|
||||
copy_operand = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (copy_operand) {
|
||||
TF_ASSIGN_OR_RETURN(HloInstruction * copy, FindOrInsertCopy(operand));
|
||||
TF_RETURN_IF_ERROR(hlo->ReplaceOperandWith(i, copy));
|
||||
changed = true;
|
||||
}
|
||||
@ -64,6 +76,31 @@ StatusOr<bool> GpuCopyInsertion::Run(HloModule* module) {
|
||||
}
|
||||
}
|
||||
|
||||
// Init values of a while nodes cannot be constants. Insert copies for any
|
||||
// constants found at the operand of a while.
|
||||
tensorflow::gtl::FlatSet<HloInstruction*> copied_constants;
|
||||
for (HloComputation* computation : module->computations()) {
|
||||
for (HloInstruction* instruction : computation->instructions()) {
|
||||
if (instruction->opcode() == HloOpcode::kWhile) {
|
||||
for (auto& pair :
|
||||
dataflow->GetInstructionValueSet(instruction->operand(0))) {
|
||||
const HloValueSet& value_set = pair.second;
|
||||
for (const HloValue* value : value_set.values()) {
|
||||
if (value->defining_instruction()->opcode() ==
|
||||
HloOpcode::kConstant &&
|
||||
!ContainsKey(copied_constants, value->defining_instruction())) {
|
||||
HloInstruction* constant = value->defining_instruction();
|
||||
TF_ASSIGN_OR_RETURN(HloInstruction * copy,
|
||||
FindOrInsertCopy(constant));
|
||||
TF_RETURN_IF_ERROR(constant->ReplaceAllUsesWith(copy));
|
||||
copied_constants.insert(constant);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return changed;
|
||||
}
|
||||
|
||||
|
||||
@ -16,8 +16,8 @@ limitations under the License.
|
||||
#ifndef TENSORFLOW_COMPILER_XLA_SERVICE_GPU_COPY_INSERTION_H_
|
||||
#define TENSORFLOW_COMPILER_XLA_SERVICE_GPU_COPY_INSERTION_H_
|
||||
|
||||
#include "tensorflow/compiler/xla/service/copy_insertion.h"
|
||||
#include "tensorflow/compiler/xla/service/hlo_module.h"
|
||||
#include "tensorflow/compiler/xla/service/hlo_pass_interface.h"
|
||||
|
||||
namespace xla {
|
||||
namespace gpu {
|
||||
@ -25,9 +25,20 @@ namespace gpu {
|
||||
// Besides the modifications made by the generic xla::CopyInsertion, this
|
||||
// GPU-specific copy insertion also materializes operands of library calls by
|
||||
// inserting kCopy instructions.
|
||||
class GpuCopyInsertion : public CopyInsertion {
|
||||
class GpuCopyInsertion : public HloPassInterface {
|
||||
public:
|
||||
tensorflow::StringPiece name() const override { return "copy-insertion"; }
|
||||
|
||||
StatusOr<bool> Run(HloModule* module) override;
|
||||
|
||||
protected:
|
||||
// Returns a copy of `hlo`. Looks in inserted_copies_ first to avoid making
|
||||
// duplicate copies.
|
||||
StatusOr<HloInstruction*> FindOrInsertCopy(HloInstruction* hlo);
|
||||
|
||||
// A map containing all copies inserted to materialize operands of library
|
||||
// calls. The key is the copied instruction and the value is the copy.
|
||||
tensorflow::gtl::FlatMap<HloInstruction*, HloInstruction*> inserted_copies_;
|
||||
};
|
||||
|
||||
} // namespace gpu
|
||||
|
||||
@ -220,9 +220,8 @@ tensorflow::Status PrepareHloModuleForIrEmitting(
|
||||
// (and sometime after) copy insertion, to avoid dead code from interfering
|
||||
// with the rewrites.
|
||||
pipeline.AddPass<HloDCE>();
|
||||
pipeline.AddPass<GpuCopyInsertion>();
|
||||
pipeline.AddPass<HloDCE>();
|
||||
pipeline.AddPass<FlattenCallGraph>();
|
||||
pipeline.AddPass<GpuCopyInsertion>();
|
||||
return pipeline.Run(hlo_module).status();
|
||||
}
|
||||
|
||||
|
||||
@ -17,9 +17,12 @@ limitations under the License.
|
||||
|
||||
#include "tensorflow/compiler/xla/service/copy_insertion.h"
|
||||
#include "tensorflow/compiler/xla/service/gpu/instruction_fusion.h"
|
||||
#include "tensorflow/compiler/xla/service/hlo_verifier.h"
|
||||
#include "tensorflow/compiler/xla/shape_util.h"
|
||||
#include "tensorflow/compiler/xla/test.h"
|
||||
#include "tensorflow/compiler/xla/test_helpers.h"
|
||||
#include "tensorflow/compiler/xla/tests/hlo_test_base.h"
|
||||
#include "tensorflow/core/lib/core/status_test_util.h"
|
||||
|
||||
namespace xla {
|
||||
namespace {
|
||||
@ -33,8 +36,6 @@ class WhileTransformerTest : public HloTestBase {
|
||||
: module_(CreateNewModule()),
|
||||
induction_variable_shape_(ShapeUtil::MakeShape(S32, {})),
|
||||
data_shape_(ShapeUtil::MakeShape(F32, {8})),
|
||||
loop_state_shape_(ShapeUtil::MakeTupleShape(
|
||||
{induction_variable_shape_, data_shape_})),
|
||||
condition_result_shape_(ShapeUtil::MakeShape(PRED, {})) {}
|
||||
|
||||
std::unique_ptr<HloComputation> BuildConditionComputation(
|
||||
@ -42,8 +43,8 @@ class WhileTransformerTest : public HloTestBase {
|
||||
auto builder = HloComputation::Builder(TestName() + ".Condition");
|
||||
auto limit_const = builder.AddInstruction(
|
||||
HloInstruction::CreateConstant(Literal::CreateR0<int32>(limit)));
|
||||
auto loop_state = builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(0, loop_state_shape_, "loop_state"));
|
||||
auto loop_state = builder.AddInstruction(HloInstruction::CreateParameter(
|
||||
0, GetLoopStateShape(tuple_index), "loop_state"));
|
||||
auto induction_variable =
|
||||
builder.AddInstruction(HloInstruction::CreateGetTupleElement(
|
||||
limit_const->shape(), loop_state, tuple_index));
|
||||
@ -58,8 +59,8 @@ class WhileTransformerTest : public HloTestBase {
|
||||
const int64 increment) {
|
||||
auto builder = HloComputation::Builder(TestName() + ".Body");
|
||||
// Create param instruction to access loop state.
|
||||
auto loop_state = builder.AddInstruction(
|
||||
HloInstruction::CreateParameter(0, loop_state_shape_, "loop_state"));
|
||||
auto loop_state = builder.AddInstruction(HloInstruction::CreateParameter(
|
||||
0, GetLoopStateShape(ind_var_tuple_index), "loop_state"));
|
||||
// Update the induction variable GTE(ind_var_tuple_index).
|
||||
auto induction_variable =
|
||||
builder.AddInstruction(HloInstruction::CreateGetTupleElement(
|
||||
@ -73,7 +74,7 @@ class WhileTransformerTest : public HloTestBase {
|
||||
data_shape_, loop_state, data_tuple_index));
|
||||
// Use 'induction_variable' in computation with no path to output tuple.
|
||||
auto update = builder.AddInstruction(
|
||||
HloInstruction::CreateBroadcast(data_shape_, induction_variable, {8}));
|
||||
HloInstruction::CreateBroadcast(data_shape_, induction_variable, {}));
|
||||
auto add1 = builder.AddInstruction(HloInstruction::CreateBinary(
|
||||
data_shape_, HloOpcode::kAdd, data, update));
|
||||
// Create output Tuple.
|
||||
@ -98,8 +99,9 @@ class WhileTransformerTest : public HloTestBase {
|
||||
HloInstruction::CreateTuple({induction_var_init, data_init}))
|
||||
: builder.AddInstruction(
|
||||
HloInstruction::CreateTuple({data_init, induction_var_init}));
|
||||
auto while_hlo = builder.AddInstruction(HloInstruction::CreateWhile(
|
||||
loop_state_shape_, condition, body, loop_state_init));
|
||||
auto while_hlo = builder.AddInstruction(
|
||||
HloInstruction::CreateWhile(GetLoopStateShape(ind_var_tuple_index),
|
||||
condition, body, loop_state_init));
|
||||
module_->AddEntryComputation(builder.Build());
|
||||
return while_hlo;
|
||||
}
|
||||
@ -115,18 +117,34 @@ class WhileTransformerTest : public HloTestBase {
|
||||
}
|
||||
|
||||
void RunCopyInsertionPass() {
|
||||
HloVerifier verifier([](const Shape& shape) {
|
||||
return ShapeUtil::ByteSizeOf(shape, /*pointer_size=*/sizeof(void*));
|
||||
});
|
||||
TF_ASSERT_OK(verifier.Run(module_.get()).status());
|
||||
CopyInsertion copy_insertion;
|
||||
EXPECT_IS_OK(copy_insertion.Run(module_.get()).status());
|
||||
TF_ASSERT_OK(copy_insertion.Run(module_.get()).status());
|
||||
}
|
||||
|
||||
Shape GetLoopStateShape(const int64 ind_var_tuple_index) {
|
||||
if (ind_var_tuple_index == 0) {
|
||||
return ShapeUtil::MakeTupleShape(
|
||||
{induction_variable_shape_, data_shape_});
|
||||
} else {
|
||||
return ShapeUtil::MakeTupleShape(
|
||||
{data_shape_, induction_variable_shape_});
|
||||
}
|
||||
}
|
||||
|
||||
std::unique_ptr<HloModule> module_;
|
||||
Shape induction_variable_shape_;
|
||||
Shape data_shape_;
|
||||
Shape loop_state_shape_;
|
||||
Shape condition_result_shape_;
|
||||
};
|
||||
|
||||
TEST_F(WhileTransformerTest, InductionVariableAtTupleElement0) {
|
||||
// TODO(b/68830972): The while transformer is far too fragile. It patterns
|
||||
// matches the exact expressions of opcodes. Re-enable when transformation is
|
||||
// more general
|
||||
TEST_F(WhileTransformerTest, DISABLED_InductionVariableAtTupleElement0) {
|
||||
// Build computation with induction variable at tuple element 0.
|
||||
auto condition =
|
||||
module_->AddEmbeddedComputation(BuildConditionComputation(0, 10));
|
||||
@ -137,13 +155,16 @@ TEST_F(WhileTransformerTest, InductionVariableAtTupleElement0) {
|
||||
RunCopyInsertionPass();
|
||||
// Run WhileTransformer.
|
||||
auto result = gpu::CanTransformWhileToFor(while_hlo);
|
||||
ASSERT_TRUE(result.ok());
|
||||
TF_ASSERT_OK(result.status());
|
||||
// Check results.
|
||||
EXPECT_THAT(result.ConsumeValueOrDie(),
|
||||
Eq(std::tuple<int64, int64, int64>(0, 10, 1)));
|
||||
}
|
||||
|
||||
TEST_F(WhileTransformerTest, InductionVariableAtTupleElement1) {
|
||||
// TODO(b/68830972): The while transformer is far too fragile. It patterns
|
||||
// matches the exact expressions of opcodes. Re-enable when transformation is
|
||||
// more general
|
||||
TEST_F(WhileTransformerTest, DISABLED_InductionVariableAtTupleElement1) {
|
||||
// Build computation with induction variable at tuple element 1.
|
||||
auto condition =
|
||||
module_->AddEmbeddedComputation(BuildConditionComputation(1, 10));
|
||||
@ -154,13 +175,16 @@ TEST_F(WhileTransformerTest, InductionVariableAtTupleElement1) {
|
||||
RunCopyInsertionPass();
|
||||
// Run WhileTransformer.
|
||||
auto result = gpu::CanTransformWhileToFor(while_hlo);
|
||||
ASSERT_TRUE(result.ok());
|
||||
TF_ASSERT_OK(result.status());
|
||||
// Check results.
|
||||
EXPECT_THAT(result.ConsumeValueOrDie(),
|
||||
Eq(std::tuple<int64, int64, int64>(0, 10, 1)));
|
||||
}
|
||||
|
||||
TEST_F(WhileTransformerTest, InvalidLoopLimit) {
|
||||
// TODO(b/68830972): The while transformer is far too fragile. It patterns
|
||||
// matches the exact expressions of opcodes. Re-enable when transformation is
|
||||
// more general
|
||||
TEST_F(WhileTransformerTest, DISABLED_InvalidLoopLimit) {
|
||||
// Build computation with invalid loop limit.
|
||||
auto condition =
|
||||
module_->AddEmbeddedComputation(BuildConditionComputation(0, 5));
|
||||
@ -176,7 +200,10 @@ TEST_F(WhileTransformerTest, InvalidLoopLimit) {
|
||||
HasSubstr("Loop start must be less than loop limit."));
|
||||
}
|
||||
|
||||
TEST_F(WhileTransformerTest, InvalidLoopIncrement) {
|
||||
// TODO(b/68830972): The while transformer is far too fragile. It patterns
|
||||
// matches the exact expressions of opcodes. Re-enable when transformation is
|
||||
// more general
|
||||
TEST_F(WhileTransformerTest, DISABLED_InvalidLoopIncrement) {
|
||||
// Build computation with invalid loop increment.
|
||||
auto condition =
|
||||
module_->AddEmbeddedComputation(BuildConditionComputation(0, 10));
|
||||
|
||||
@ -144,8 +144,10 @@ class BufferValueMap {
|
||||
// Move the given value into the given buffer.
|
||||
void MoveValueToBuffer(const HloValue& value, BufferNumber buffer_number) {
|
||||
BufferNumber old_buffer_number = value_to_buffer_number_.at(&value);
|
||||
buffers_.at(old_buffer_number).erase(&value);
|
||||
if (buffers_.at(old_buffer_number).empty()) {
|
||||
tensorflow::gtl::FlatSet<const HloValue*>& old_value_set =
|
||||
buffers_.at(old_buffer_number);
|
||||
old_value_set.erase(&value);
|
||||
if (old_value_set.empty()) {
|
||||
buffers_.erase(old_buffer_number);
|
||||
}
|
||||
|
||||
@ -175,7 +177,7 @@ class BufferValueMap {
|
||||
// Value is init of a while (use is while).
|
||||
std::vector<BufferNumber> aliased_buffers;
|
||||
for (const HloUse& use : value.uses()) {
|
||||
VLOG(1) << "use of value " << value.ToShortString() << ": " << use;
|
||||
VLOG(2) << "use of value " << value.ToShortString() << ": " << use;
|
||||
if (use.instruction->opcode() == HloOpcode::kWhile) {
|
||||
// Determine the while value that this shares a buffer with.
|
||||
const HloValue& while_value =
|
||||
@ -411,7 +413,7 @@ string HloAliasAnalysis::ToString() const {
|
||||
/* static */
|
||||
StatusOr<std::unique_ptr<HloAliasAnalysis>> HloAliasAnalysis::Run(
|
||||
HloModule* module) {
|
||||
VLOG(1) << "HloAliasAnalysis::Run on module " << module->name();
|
||||
VLOG(2) << "HloAliasAnalysis::Run on module " << module->name();
|
||||
XLA_VLOG_LINES(2, module->ToString());
|
||||
|
||||
auto alias_analysis = WrapUnique(new HloAliasAnalysis(module));
|
||||
|
||||
@ -412,16 +412,18 @@ HloComputationProto HloComputation::ToProto() const {
|
||||
/* static */ StatusOr<std::unique_ptr<HloComputation>>
|
||||
HloComputation::CreateFromProto(
|
||||
HloModule* module, const HloComputationProto& proto,
|
||||
tensorflow::gtl::FlatMap<string, HloComputation*>* computation_map,
|
||||
const tensorflow::gtl::FlatMap<string, HloComputation*>& computation_map,
|
||||
const std::function<void(std::unique_ptr<HloComputation>)>&
|
||||
add_fused_computation,
|
||||
HloInstruction* fusion_instruction) {
|
||||
std::vector<std::unique_ptr<HloInstruction>> instructions;
|
||||
tensorflow::gtl::FlatMap<string, HloInstruction*> instruction_map;
|
||||
int64 parameter_count = 0;
|
||||
for (const HloInstructionProto& instruction_proto : proto.instructions()) {
|
||||
TF_ASSIGN_OR_RETURN(
|
||||
std::unique_ptr<HloInstruction> instruction,
|
||||
HloInstruction::CreateFromProto(module, instruction_proto,
|
||||
instruction_map, computation_map));
|
||||
TF_ASSIGN_OR_RETURN(std::unique_ptr<HloInstruction> instruction,
|
||||
HloInstruction::CreateFromProto(
|
||||
module, instruction_proto, instruction_map,
|
||||
computation_map, add_fused_computation));
|
||||
if (instruction->opcode() == HloOpcode::kParameter) {
|
||||
parameter_count++;
|
||||
}
|
||||
@ -531,6 +533,7 @@ StatusOr<HloInstruction*> HloComputation::DeepCopyInstruction(
|
||||
|
||||
if (indices_to_copy != nullptr &&
|
||||
!ShapeUtil::Compatible(instruction->shape(), indices_to_copy->shape())) {
|
||||
LOG(FATAL) << "DEATH!";
|
||||
return FailedPrecondition(
|
||||
"Can't deep copy instruction %s: given shape tree of indices to copy "
|
||||
"has incompatible shape",
|
||||
|
||||
@ -152,12 +152,18 @@ class HloComputation {
|
||||
// computation_map: a map from computation name to HloComputation*. This map
|
||||
// must contain all computations which the newly constructed computation
|
||||
// calls.
|
||||
// fusion_instruction: if non-null then the newly created computation will be
|
||||
// add_fused_computation: A function to call to add a fused
|
||||
// computation. Used (clearly) when the instruction is a fusion
|
||||
// instruction.
|
||||
// fusion_instruction: if non-null then the newly created computation will
|
||||
// be
|
||||
// constructed as a fused computation with this instruction as its fusion
|
||||
// parent.
|
||||
static StatusOr<std::unique_ptr<HloComputation>> CreateFromProto(
|
||||
HloModule* module, const HloComputationProto& proto,
|
||||
tensorflow::gtl::FlatMap<string, HloComputation*>* computation_map,
|
||||
const tensorflow::gtl::FlatMap<string, HloComputation*>& computation_map,
|
||||
const std::function<void(std::unique_ptr<HloComputation>)>&
|
||||
add_fused_computation,
|
||||
HloInstruction* fusion_instruction = nullptr);
|
||||
|
||||
// Gets the instructions in this computation.
|
||||
|
||||
@ -75,11 +75,41 @@ HloValue* HloDataflowAnalysis::NewHloValue(HloInstruction* instruction,
|
||||
std::forward_as_tuple(value_id, instruction, index, is_phi));
|
||||
CHECK(emplaced.second);
|
||||
|
||||
VLOG(4) << "NewHloValue = " << emplaced.first->second.ToShortString();
|
||||
|
||||
return &emplaced.first->second;
|
||||
}
|
||||
|
||||
void HloDataflowAnalysis::DeleteHloValue(HloValue::Id value_id) {
|
||||
values_.erase(value_id);
|
||||
void HloDataflowAnalysis::MarkValueForDeletion(HloValue::Id value_id) {
|
||||
HloValue& value = values_.at(value_id);
|
||||
VLOG(4) << "MarkValueForDeletion(" << value.ToShortString() << ")";
|
||||
|
||||
value_ids_to_delete_.push_back(value_id);
|
||||
}
|
||||
|
||||
void HloDataflowAnalysis::DeleteMarkedValues() {
|
||||
// Verify that no marked-for-deletion values are in any of the value sets.
|
||||
tensorflow::gtl::FlatSet<HloValue::Id> id_set(value_ids_to_delete_.begin(),
|
||||
value_ids_to_delete_.end());
|
||||
for (const auto& pair : value_sets_) {
|
||||
const HloInstruction* instruction = pair.first;
|
||||
const InstructionValueSet& instruction_value_set = pair.second;
|
||||
for (const auto& index_value_set : instruction_value_set) {
|
||||
const HloValueSet& value_set = index_value_set.second;
|
||||
for (const HloValue* value : value_set.values()) {
|
||||
DCHECK(!ContainsKey(id_set, value->id()))
|
||||
<< "Value " << value->ToShortString()
|
||||
<< " marked for deletion, but still exists in value set for "
|
||||
"instruction "
|
||||
<< instruction->name();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (HloValue::Id value_id : value_ids_to_delete_) {
|
||||
values_.erase(value_id);
|
||||
}
|
||||
value_ids_to_delete_.clear();
|
||||
}
|
||||
|
||||
string HloDataflowAnalysis::ToString() const {
|
||||
@ -121,6 +151,7 @@ bool HloDataflowAnalysis::Phi(
|
||||
HloInstruction* instruction,
|
||||
tensorflow::gtl::ArraySlice<const InstructionValueSet*> inputs) {
|
||||
CHECK(ssa_form_);
|
||||
VLOG(4) << "Phi(" << instruction->name() << ")";
|
||||
|
||||
for (const InstructionValueSet* input : inputs) {
|
||||
DCHECK(ShapeUtil::Compatible(instruction->shape(), input->shape()));
|
||||
@ -183,7 +214,7 @@ bool HloDataflowAnalysis::Phi(
|
||||
} else if (current_value != &new_value) {
|
||||
if (current_value_defined_here) {
|
||||
// Remove the existing phi.
|
||||
DeleteHloValue(current_value->id());
|
||||
MarkValueForDeletion(current_value->id());
|
||||
}
|
||||
value_set.Clear();
|
||||
value_set.AddValue(&new_value);
|
||||
@ -193,7 +224,8 @@ bool HloDataflowAnalysis::Phi(
|
||||
// Multiple distinct values reach this point. A phi value is
|
||||
// necessary.
|
||||
CHECK_GT(input_value_ids.size(), 1);
|
||||
if (current_value == nullptr || !current_value->is_phi()) {
|
||||
if (current_value == nullptr ||
|
||||
!(current_value->is_phi() && current_value_defined_here)) {
|
||||
value_set.Clear();
|
||||
value_set.AddValue(NewHloValue(instruction, index, /*is_phi=*/true));
|
||||
changed = true;
|
||||
@ -436,11 +468,13 @@ bool HloDataflowAnalysis::UpdateInstructionValueSet(
|
||||
}
|
||||
}
|
||||
|
||||
void HloDataflowAnalysis::UpdateInstructionsAndPropagate(
|
||||
tensorflow::gtl::ArraySlice<HloInstruction*> instructions) {
|
||||
void HloDataflowAnalysis::Propagate() {
|
||||
std::queue<HloInstruction*> worklist;
|
||||
for (HloInstruction* instruction : instructions) {
|
||||
worklist.push(instruction);
|
||||
|
||||
for (HloComputation* computation : module_->computations()) {
|
||||
for (HloInstruction* instruction : computation->instructions()) {
|
||||
worklist.push(instruction);
|
||||
}
|
||||
}
|
||||
|
||||
while (!worklist.empty()) {
|
||||
@ -597,18 +631,10 @@ StatusOr<std::unique_ptr<HloDataflowAnalysis>> HloDataflowAnalysis::Run(
|
||||
new HloDataflowAnalysis(module, ssa_form, bitcast_defines_value));
|
||||
|
||||
TF_RETURN_IF_ERROR(dataflow_analysis->InitializeInstructionValueSets());
|
||||
dataflow_analysis->Propagate();
|
||||
|
||||
// Construct list of all instructions to initialize the worklist to propagate
|
||||
// the data flow. For efficiency sort the instruction in post order so
|
||||
// producers appear before consumers.
|
||||
std::vector<HloInstruction*> all_instructions;
|
||||
for (const HloComputation* computation : module->MakeComputationPostOrder()) {
|
||||
for (HloInstruction* instruction :
|
||||
computation->MakeInstructionPostOrder()) {
|
||||
all_instructions.push_back(instruction);
|
||||
}
|
||||
}
|
||||
dataflow_analysis->UpdateInstructionsAndPropagate(all_instructions);
|
||||
// Delete all values marked for deletion.
|
||||
dataflow_analysis->DeleteMarkedValues();
|
||||
|
||||
// Add in positions to all values.
|
||||
for (const HloComputation* computation : module->computations()) {
|
||||
|
||||
@ -126,13 +126,16 @@ class HloDataflowAnalysis {
|
||||
HloValue* NewHloValue(HloInstruction* instruction, const ShapeIndex& index,
|
||||
bool is_phi = false);
|
||||
|
||||
// Delete the HloValue with the given ID.
|
||||
void DeleteHloValue(HloValue::Id value_id);
|
||||
// Mark the HloValue with the given ID for deletion.
|
||||
void MarkValueForDeletion(HloValue::Id value_id);
|
||||
|
||||
// Delete all HloValues marked for deletion. Should be called after
|
||||
// propagation is complete.
|
||||
void DeleteMarkedValues();
|
||||
|
||||
// Constructs and initializes the InstructionValueSets of all instructions to
|
||||
// contain exactly the HloValues defined by each instruction. These values can
|
||||
// then propagated throughout the HLO graph by calling
|
||||
// UpdateInstructionsAndPropagate.
|
||||
// then propagated throughout the HLO graph by calling Propagate.
|
||||
Status InitializeInstructionValueSets();
|
||||
|
||||
// Updates the value set of the given instruction based on the values flowing
|
||||
@ -150,10 +153,8 @@ class HloDataflowAnalysis {
|
||||
bool UpdateTupleValueSet(HloInstruction* tuple);
|
||||
bool UpdateWhileValueSet(HloInstruction* xla_while);
|
||||
|
||||
// Update the value sets of the given instructions and propagate the
|
||||
// changes to fixed point.
|
||||
void UpdateInstructionsAndPropagate(
|
||||
tensorflow::gtl::ArraySlice<HloInstruction*> instructions);
|
||||
// Propagate the dataflow through the module.
|
||||
void Propagate();
|
||||
|
||||
// Return the result of the SSA Phi function applied to the given inputs at
|
||||
// the given instruction. If skip_top_level is true, then the top level of the
|
||||
@ -189,6 +190,11 @@ class HloDataflowAnalysis {
|
||||
// A map from instruction to InstructionValueSet.
|
||||
std::unordered_map<const HloInstruction*, InstructionValueSet> value_sets_;
|
||||
|
||||
// Values marked for deletion during construction. We don't delete them
|
||||
// immediately because references to them may still remain in ValueSets. After
|
||||
// construction, these values are deleted.
|
||||
std::vector<HloValue::Id> value_ids_to_delete_;
|
||||
|
||||
// A vector containing all HloValues sorted by HloValue::Id.
|
||||
std::vector<const HloValue*> values_vector_;
|
||||
|
||||
|
||||
@ -37,6 +37,9 @@ namespace xla {
|
||||
StatusOr<bool> HloDCE::Run(HloModule* module) {
|
||||
bool changed = false;
|
||||
|
||||
VLOG(2) << "Before dce:";
|
||||
XLA_VLOG_LINES(2, module->ToString());
|
||||
|
||||
for (auto* computation : module->MakeNonfusionComputations()) {
|
||||
std::unordered_set<HloInstruction*> live_instructions;
|
||||
TF_RETURN_IF_ERROR(computation->root_instruction()->Accept(
|
||||
@ -58,6 +61,8 @@ StatusOr<bool> HloDCE::Run(HloModule* module) {
|
||||
}
|
||||
|
||||
for (HloInstruction* dead_root : dead_roots) {
|
||||
VLOG(1) << "Removing dead root " << dead_root->ToString()
|
||||
<< " and it's unused operands";
|
||||
TF_RETURN_IF_ERROR(
|
||||
computation->RemoveInstructionAndUnusedOperands(dead_root));
|
||||
changed = true;
|
||||
@ -87,6 +92,9 @@ StatusOr<bool> HloDCE::Run(HloModule* module) {
|
||||
}
|
||||
}
|
||||
|
||||
VLOG(2) << "After dce:";
|
||||
XLA_VLOG_LINES(2, module->ToString());
|
||||
|
||||
return changed;
|
||||
}
|
||||
|
||||
|
||||
@ -51,7 +51,9 @@ using ::tensorflow::strings::StrCat;
|
||||
StatusOr<std::unique_ptr<HloInstruction>> HloInstruction::CreateFromProto(
|
||||
HloModule* module, const HloInstructionProto& proto,
|
||||
const tensorflow::gtl::FlatMap<string, HloInstruction*>& instruction_map,
|
||||
tensorflow::gtl::FlatMap<string, HloComputation*>* computation_map) {
|
||||
const tensorflow::gtl::FlatMap<string, HloComputation*>& computation_map,
|
||||
const std::function<void(std::unique_ptr<HloComputation>)>&
|
||||
add_fused_computation) {
|
||||
TF_RET_CHECK(!proto.opcode().empty());
|
||||
TF_ASSIGN_OR_RETURN(HloOpcode opcode, StringToHloOpcode(proto.opcode()));
|
||||
TF_RET_CHECK(proto.has_shape());
|
||||
@ -77,19 +79,19 @@ StatusOr<std::unique_ptr<HloInstruction>> HloInstruction::CreateFromProto(
|
||||
TF_RET_CHECK(!proto.fusion_kind().empty());
|
||||
TF_ASSIGN_OR_RETURN(instruction->fusion_kind_,
|
||||
StringToFusionKind(proto.fusion_kind()));
|
||||
TF_ASSIGN_OR_RETURN(
|
||||
std::unique_ptr<HloComputation> fused_computation,
|
||||
HloComputation::CreateFromProto(
|
||||
module, proto.fused_instructions_computation(), computation_map,
|
||||
/*fusion_instruction=*/instruction.get()));
|
||||
instruction->called_computations_.push_back(
|
||||
module->AddEmbeddedComputation(std::move(fused_computation)));
|
||||
TF_ASSIGN_OR_RETURN(std::unique_ptr<HloComputation> fused_computation,
|
||||
HloComputation::CreateFromProto(
|
||||
module, proto.fused_instructions_computation(),
|
||||
computation_map, add_fused_computation,
|
||||
/*fusion_instruction=*/instruction.get()));
|
||||
instruction->called_computations_.push_back(fused_computation.get());
|
||||
add_fused_computation(std::move(fused_computation));
|
||||
} else {
|
||||
for (const string& computation_name : proto.called_computation_names()) {
|
||||
TF_RET_CHECK(ContainsKey(*computation_map, computation_name))
|
||||
TF_RET_CHECK(ContainsKey(computation_map, computation_name))
|
||||
<< "No computation named " << computation_name;
|
||||
instruction->called_computations_.push_back(
|
||||
computation_map->at(computation_name));
|
||||
computation_map.at(computation_name));
|
||||
}
|
||||
}
|
||||
|
||||
@ -2009,8 +2011,10 @@ string HloInstruction::ToCategory() const {
|
||||
bool saw_rank_1 = false;
|
||||
bool saw_higher_rank = false;
|
||||
for (const auto* operand : operands()) {
|
||||
saw_rank_1 |= ShapeUtil::Rank(operand->shape()) == 1;
|
||||
saw_higher_rank |= ShapeUtil::Rank(operand->shape()) > 1;
|
||||
if (!ShapeUtil::IsTuple(operand->shape())) {
|
||||
saw_rank_1 |= ShapeUtil::Rank(operand->shape()) == 1;
|
||||
saw_higher_rank |= ShapeUtil::Rank(operand->shape()) > 1;
|
||||
}
|
||||
}
|
||||
if (saw_rank_1 && saw_higher_rank) {
|
||||
return "rank-1-broadcast binary fusion";
|
||||
@ -2295,8 +2299,8 @@ Status HloInstruction::Visit(DfsHloVisitorBase<HloInstructionPtr>* visitor) {
|
||||
template Status HloInstruction::Visit(DfsHloVisitor* visitor);
|
||||
template Status HloInstruction::Visit(ConstDfsHloVisitor* visitor);
|
||||
|
||||
using DFSStack =
|
||||
tensorflow::gtl::InlinedVector<std::pair<int, HloInstruction*>, 16>;
|
||||
using DFSStack = tensorflow::gtl::InlinedVector<
|
||||
std::pair<HloInstruction::Id, HloInstruction*>, 16>;
|
||||
|
||||
// Push "child" onto the dfs_stack if not already visited. Returns false if a
|
||||
// cycle was detected, and true otherwise.
|
||||
@ -2304,7 +2308,7 @@ template <typename Visitor>
|
||||
inline bool PushDFSChild(Visitor* visitor, DFSStack* dfs_stack,
|
||||
HloInstruction* child) {
|
||||
CHECK(child != nullptr);
|
||||
const int id = child->unique_id();
|
||||
const HloInstruction::Id id = child->unique_id();
|
||||
CHECK_GE(id, 0) << "instruction may not have a parent computation";
|
||||
switch (visitor->GetVisitState(id)) {
|
||||
case Visitor::kVisiting:
|
||||
@ -2321,8 +2325,8 @@ inline bool PushDFSChild(Visitor* visitor, DFSStack* dfs_stack,
|
||||
}
|
||||
|
||||
using InternalCompareFunction =
|
||||
std::function<bool(std::pair<int, const HloInstruction*>,
|
||||
std::pair<int, const HloInstruction*>)>;
|
||||
std::function<bool(std::pair<HloInstruction::Id, const HloInstruction*>,
|
||||
std::pair<HloInstruction::Id, const HloInstruction*>)>;
|
||||
template <typename Visitor>
|
||||
static Status PostOrderDFS(HloInstruction* root, Visitor* visitor,
|
||||
const InternalCompareFunction* operand_order,
|
||||
@ -2341,7 +2345,7 @@ static Status PostOrderDFS(HloInstruction* root, Visitor* visitor,
|
||||
do {
|
||||
DCHECK(!dfs_stack.empty());
|
||||
|
||||
int current_id = dfs_stack.back().first;
|
||||
HloInstruction::Id current_id = dfs_stack.back().first;
|
||||
HloInstruction* current_node = dfs_stack.back().second;
|
||||
CHECK_GE(current_id, 0) << current_id << ": " << current_node
|
||||
<< ": instruction may not have parent computation";
|
||||
@ -2420,13 +2424,13 @@ Status HloInstruction::AcceptWithOperandOrder(
|
||||
DfsHloVisitor* visitor, const CompareFunction& operand_order,
|
||||
bool call_finish_visit) {
|
||||
VLOG(2) << "HloInstruction::AcceptWithOperandOrder(" << name() << ")";
|
||||
InternalCompareFunction func = [&operand_order](
|
||||
std::pair<int, const HloInstruction*> a,
|
||||
std::pair<int, const HloInstruction*> b) {
|
||||
// Call the client's comparison function on the actual HloInstruction*
|
||||
// objects (ignoring the internal ids we also have in our stack entries)
|
||||
return operand_order(a.second, b.second);
|
||||
};
|
||||
InternalCompareFunction func =
|
||||
[&operand_order](std::pair<HloInstruction::Id, const HloInstruction*> a,
|
||||
std::pair<HloInstruction::Id, const HloInstruction*> b) {
|
||||
// Call the client's comparison function on the actual HloInstruction*
|
||||
// objects (ignoring the internal ids we also have in our stack entries)
|
||||
return operand_order(a.second, b.second);
|
||||
};
|
||||
TF_RETURN_IF_ERROR(PostOrderDFS(this, visitor, &func,
|
||||
/*ignore_control_predecessors=*/false));
|
||||
if (call_finish_visit) {
|
||||
|
||||
@ -83,12 +83,16 @@ class HloInstruction {
|
||||
// must contain all operands of the newly constructed instruction.
|
||||
// computation_map: a map from computation name to HloComputation*. This map
|
||||
// must contain all computations which the newly constructed instruction
|
||||
// calls. If the instruction is a fusion instruction, then the fusion
|
||||
// computation is added to this map and the module.
|
||||
// calls.
|
||||
// add_fused_computation: A function to call to add a fused
|
||||
// computation. Used (clearly) when the instruction is a fusion
|
||||
// instruction.
|
||||
static StatusOr<std::unique_ptr<HloInstruction>> CreateFromProto(
|
||||
HloModule* module, const HloInstructionProto& proto,
|
||||
const tensorflow::gtl::FlatMap<string, HloInstruction*>& instruction_map,
|
||||
tensorflow::gtl::FlatMap<string, HloComputation*>* computation_map);
|
||||
const tensorflow::gtl::FlatMap<string, HloComputation*>& computation_map,
|
||||
const std::function<void(std::unique_ptr<HloComputation>)>&
|
||||
add_fused_computation);
|
||||
|
||||
// Creates a parameter-retrieving instruction.
|
||||
static std::unique_ptr<HloInstruction> CreateParameter(int64 parameter_number,
|
||||
@ -977,7 +981,8 @@ class HloInstruction {
|
||||
void UniquifyName(NameUniquer* name_uniquer);
|
||||
|
||||
// Set the unique id for this instruction to "id"
|
||||
void SetUniqueId(int id) {
|
||||
using Id = int;
|
||||
void SetUniqueId(Id id) {
|
||||
CHECK_EQ(unique_id_, -1); // Should not be assigned already
|
||||
CHECK_GE(id, 0);
|
||||
unique_id_ = id;
|
||||
@ -985,7 +990,7 @@ class HloInstruction {
|
||||
|
||||
// Return the unique ID assigned to this node via SetUniqueId (or -1
|
||||
// if no id has been assigned yet).
|
||||
int unique_id() const { return unique_id_; }
|
||||
Id unique_id() const { return unique_id_; }
|
||||
|
||||
// Sets the debug metadata for this instruction.
|
||||
void set_metadata(const OpMetadata& metadata) { metadata_ = metadata; }
|
||||
@ -1088,7 +1093,7 @@ class HloInstruction {
|
||||
// Returns how this instruction uses elements of its `i`th operand.
|
||||
UseKind OperandElementUse(int64 i) const;
|
||||
|
||||
int unique_id_; // Unique to this HloInstruction within a HloModule
|
||||
Id unique_id_; // Unique to this HloInstruction within a HloModule
|
||||
|
||||
// Opcode for this instruction.
|
||||
HloOpcode opcode_;
|
||||
|
||||
@ -296,9 +296,16 @@ StatusOr<std::unique_ptr<HloModule>> HloModule::CreateFromProto(
|
||||
|
||||
tensorflow::gtl::FlatMap<string, HloComputation*> computation_map;
|
||||
for (const HloComputationProto& computation_proto : proto.computations()) {
|
||||
TF_ASSIGN_OR_RETURN(std::unique_ptr<HloComputation> computation,
|
||||
HloComputation::CreateFromProto(
|
||||
module.get(), computation_proto, &computation_map));
|
||||
TF_ASSIGN_OR_RETURN(
|
||||
std::unique_ptr<HloComputation> computation,
|
||||
HloComputation::CreateFromProto(
|
||||
module.get(), computation_proto, computation_map,
|
||||
/*add_fused_computation=*/
|
||||
[&module](std::unique_ptr<HloComputation> fused_computation) {
|
||||
module->AddComputationInternal(std::move(fused_computation),
|
||||
/*is_entry=*/false,
|
||||
/*uniquify_names=*/false);
|
||||
}));
|
||||
CHECK_NE(computation.get(), nullptr);
|
||||
TF_RET_CHECK(!ContainsKey(computation_map, computation->name()));
|
||||
string computation_name = computation->name();
|
||||
|
||||
@ -184,7 +184,7 @@ void HloValue::AddPosition(HloInstruction* instruction,
|
||||
live_out_of_module_ = true;
|
||||
}
|
||||
|
||||
if (instruction == instruction->parent()->root_instruction()) {
|
||||
if (instruction == defining_instruction()->parent()->root_instruction()) {
|
||||
live_out_of_computation_ = true;
|
||||
}
|
||||
}
|
||||
|
||||
@ -55,22 +55,34 @@ static Status EmitDynamicUpdateSliceInPlaceImpl(
|
||||
// Calculate output_index, where we'll write the value from update. For
|
||||
// each dimension,
|
||||
//
|
||||
// output_index[dim] = (start_index[dim] + update_index[dim]) % dim_size.
|
||||
// output_index[dim] = (start_index[dim] + update_index[dim])
|
||||
//
|
||||
IrArray::Index output_index(rank);
|
||||
for (int64 i = 0; i < rank; ++i) {
|
||||
llvm::Value* dim_size = llvm::ConstantInt::get(
|
||||
update_index[i]->getType(), output_shape.dimensions(i));
|
||||
llvm::Value* start_index0 = ir_builder->CreateZExtOrBitCast(
|
||||
start_index[i], update_index[i]->getType());
|
||||
output_index[i] = ir_builder->CreateURem(
|
||||
ir_builder->CreateAdd(start_index0, update_index[i]), dim_size);
|
||||
output_index[i] = ir_builder->CreateAdd(start_index0, update_index[i]);
|
||||
}
|
||||
|
||||
// Check if 'index' intersects start/end indices. If it does not (indices
|
||||
// are out of bounds) then no update is performed.
|
||||
llvm::Value* in_bounds = llvm::ConstantInt::get(ir_builder->getInt1Ty(), 1);
|
||||
for (int64 i = 0; i < rank; ++i) {
|
||||
llvm::Value* dim_size = llvm::ConstantInt::get(
|
||||
output_index[i]->getType(), output_shape.dimensions(i));
|
||||
in_bounds = ir_builder->CreateAnd(
|
||||
in_bounds, ir_builder->CreateICmpSLT(output_index[i], dim_size),
|
||||
"in_bounds");
|
||||
}
|
||||
|
||||
// Do output[output_index] = update[update_index].
|
||||
TF_ASSIGN_OR_RETURN(llvm::Value * update_data,
|
||||
update_array_generator(update_index));
|
||||
output_array.EmitWriteArrayElement(output_index, update_data, ir_builder);
|
||||
llvm::Value* input_data =
|
||||
output_array.EmitReadArrayElement(output_index, ir_builder);
|
||||
llvm::Value* to_write_data =
|
||||
ir_builder->CreateSelect(in_bounds, update_data, input_data);
|
||||
output_array.EmitWriteArrayElement(output_index, to_write_data, ir_builder);
|
||||
return Status::OK();
|
||||
};
|
||||
|
||||
|
||||
@ -180,7 +180,8 @@ XLA_TEST_F(TupleTest, TupleGTEToTuple) {
|
||||
ComputeAndCompareTuple(&builder, *expected, {}, error_spec_);
|
||||
}
|
||||
|
||||
XLA_TEST_F(TupleTest, SelectBetweenPredTuples) {
|
||||
// TODO(b/68395210): GPU does not tolerate ambiguous top-level buffers.
|
||||
XLA_TEST_F(TupleTest, DISABLED_ON_GPU(SelectBetweenPredTuples)) {
|
||||
ComputationBuilder b(client_, TestName());
|
||||
ComputationDataHandle v1, v2;
|
||||
|
||||
|
||||
@ -15,6 +15,7 @@ limitations under the License.
|
||||
#include "tensorflow/compiler/xla/legacy_flags/debug_options_flags.h"
|
||||
#include "tensorflow/core/platform/logging.h"
|
||||
#include "tensorflow/core/platform/test.h"
|
||||
#include "tensorflow/core/platform/test_benchmark.h"
|
||||
|
||||
GTEST_API_ int main(int argc, char** argv) {
|
||||
std::vector<tensorflow::Flag> flag_list;
|
||||
@ -30,5 +31,7 @@ GTEST_API_ int main(int argc, char** argv) {
|
||||
LOG(ERROR) << "Unknown argument " << argv[1] << "\n" << usage;
|
||||
return 2;
|
||||
}
|
||||
return RUN_ALL_TESTS();
|
||||
int result = RUN_ALL_TESTS();
|
||||
tensorflow::testing::RunBenchmarks();
|
||||
return result;
|
||||
}
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user