1 /* Copyright 2018 The TensorFlow Authors. All Rights Reserved.
2
3 Licensed under the Apache License, Version 2.0 (the "License");
4 you may not use this file except in compliance with the License.
5 You may obtain a copy of the License at
6
7 http://www.apache.org/licenses/LICENSE-2.0
8
9 Unless required by applicable law or agreed to in writing, software
10 distributed under the License is distributed on an "AS IS" BASIS,
11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 See the License for the specific language governing permissions and
13 limitations under the License.
14 ==============================================================================*/
15
16 #include "tensorflow/compiler/xla/service/hlo_domain_remover.h"
17
18 #include "tensorflow/compiler/xla/service/hlo_computation.h"
19 #include "tensorflow/compiler/xla/service/hlo_domain_map.h"
20 #include "tensorflow/compiler/xla/service/hlo_domain_verifier.h"
21 #include "tensorflow/compiler/xla/service/hlo_graph_dumper.h"
22 #include "tensorflow/compiler/xla/service/hlo_instruction.h"
23 #include "tensorflow/compiler/xla/service/hlo_opcode.h"
24 #include "tensorflow/compiler/xla/types.h"
25
26 namespace xla {
27
28 class HloDomainRemover::RunContext {
29 public:
RunContext(HloModule * module,HloDomainRemover * remover)30 RunContext(HloModule* module, HloDomainRemover* remover)
31 : module_(module), remover_(remover) {}
32
33 StatusOr<bool> Run(
34 const absl::flat_hash_set<absl::string_view>& execution_threads);
35
36 private:
37 // Verifies the consistency of the domain, and normalizes the instructions
38 // within it.
39 Status VerifyAndNormalizeDomain(const DomainMetadata::Domain& domain);
40
41 HloModule* module_;
42 HloDomainRemover* remover_;
43 };
44
VerifyAndNormalizeDomain(const DomainMetadata::Domain & domain)45 Status HloDomainRemover::RunContext::VerifyAndNormalizeDomain(
46 const DomainMetadata::Domain& domain) {
47 TF_ASSIGN_OR_RETURN(const DomainMetadata* ref_metadata,
48 HloDomainVerifier::VerifyDomain(domain));
49 if (ref_metadata != nullptr) {
50 VLOG(4) << "Applying domain normalization: " << ref_metadata->ToString();
51 TF_RETURN_IF_ERROR(remover_->normalizer_(domain, ref_metadata));
52 } else {
53 // No kDomain instruction was present within this domain, so call the
54 // generic normalization functions and have them apply their heuristic.
55 VLOG(2) << "Applying domain-less normalization";
56 TF_RETURN_IF_ERROR(remover_->normalizer_(domain, nullptr));
57 }
58 return OkStatus();
59 }
60
Run(const absl::flat_hash_set<absl::string_view> & execution_threads)61 StatusOr<bool> HloDomainRemover::RunContext::Run(
62 const absl::flat_hash_set<absl::string_view>& execution_threads) {
63 VLOG(4) << "Processing metadata domain: '" << remover_->kind_ << "'";
64 int64_t removed_domains = 0;
65 for (HloComputation* computation : module_->computations(execution_threads)) {
66 // First create the domain instruction sets. A domain instruction set is
67 // the set of instructions whose edges never cross a kDomain instruction.
68 TF_ASSIGN_OR_RETURN(std::unique_ptr<HloDomainMap> domain_map,
69 HloDomainMap::Create(computation, remover_->kind_));
70 // Verify and normalize every domain populated within the map.
71 for (auto& domain : domain_map->GetDomains()) {
72 TF_RETURN_IF_ERROR(VerifyAndNormalizeDomain(*domain));
73 }
74
75 // Now remove all the kDomain instructions of the kind specified by the
76 // remover, that are within the currently processed computation from the
77 // graph.
78 for (HloInstruction* instruction :
79 computation->MakeInstructionPostOrder()) {
80 for (HloInstruction* operand : instruction->unique_operands()) {
81 if (domain_map->IsDomainInstruction(operand)) {
82 VLOG(5) << "Removing " << operand->name();
83 TF_RETURN_IF_ERROR(
84 operand->ReplaceAllUsesWith(operand->mutable_operand(0)));
85 TF_RETURN_IF_ERROR(computation->RemoveInstruction(operand));
86 ++removed_domains;
87 }
88 }
89 }
90 HloInstruction* root = computation->root_instruction();
91 if (root != nullptr && domain_map->IsDomainInstruction(root)) {
92 VLOG(5) << "Removing " << root->name();
93 computation->set_root_instruction(root->mutable_operand(0));
94 TF_RETURN_IF_ERROR(computation->RemoveInstruction(root));
95 ++removed_domains;
96 }
97 }
98 VLOG(3) << "Removed " << removed_domains << " kDomain instructions of '"
99 << remover_->kind_ << "' kind";
100 return removed_domains > 0;
101 }
102
RemoveExitDomains(HloInstruction * instruction,absl::string_view domain_kind)103 StatusOr<int64_t> HloDomainRemover::RemoveExitDomains(
104 HloInstruction* instruction, absl::string_view domain_kind) {
105 int64_t removed_domains = 0;
106 HloComputation* computation = instruction->parent();
107 // Make a const copy of instruction's users to loop through later, as the
108 // users vector could be changed during the loop(e.g. ReplaceAllUsesWith).
109 const std::vector<HloInstruction*> users(instruction->users());
110 for (HloInstruction* user : users) {
111 if (user->opcode() == HloOpcode::kDomain &&
112 user->user_side_metadata().Kind() == domain_kind &&
113 user->operand_side_metadata().Kind() == domain_kind) {
114 VLOG(5) << "Removing exit domain " << user->name();
115 TF_RETURN_IF_ERROR(user->ReplaceAllUsesWith(instruction));
116 TF_RETURN_IF_ERROR(computation->RemoveInstruction(user));
117 ++removed_domains;
118 }
119 }
120 return removed_domains;
121 }
122
Run(HloModule * module,const absl::flat_hash_set<absl::string_view> & execution_threads)123 StatusOr<bool> HloDomainRemover::Run(
124 HloModule* module,
125 const absl::flat_hash_set<absl::string_view>& execution_threads) {
126 RunContext run_context(module, this);
127 return run_context.Run(execution_threads);
128 }
129
130 } // namespace xla
131