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_verifier.h"
17
18 #include <set>
19
20 #include "tensorflow/compiler/xla/service/hlo_computation.h"
21 #include "tensorflow/compiler/xla/service/hlo_domain_map.h"
22 #include "tensorflow/compiler/xla/service/hlo_graph_dumper.h"
23 #include "tensorflow/compiler/xla/service/hlo_instruction.h"
24 #include "tensorflow/compiler/xla/service/hlo_opcode.h"
25 #include "tensorflow/compiler/xla/types.h"
26
27 namespace xla {
28
29 class HloDomainVerifier::RunContext {
30 public:
RunContext(HloModule * module,HloDomainVerifier * verifier)31 RunContext(HloModule* module, HloDomainVerifier* verifier)
32 : module_(module), verifier_(verifier) {}
33
34 Status Run(const absl::flat_hash_set<absl::string_view>& execution_threads);
35
36 private:
37 // If the verifier caller passed an empty vector for kinds, we collect all the
38 // available domain types.
39 Status PopulateDomainKinds(
40 const absl::flat_hash_set<absl::string_view>& execution_threads);
41
42 HloModule* module_;
43 HloDomainVerifier* verifier_;
44 };
45
PopulateDomainKinds(const absl::flat_hash_set<absl::string_view> & execution_threads)46 Status HloDomainVerifier::RunContext::PopulateDomainKinds(
47 const absl::flat_hash_set<absl::string_view>& execution_threads) {
48 if (verifier_->kinds_.empty()) {
49 // The caller specified no domain kinds, collect all the ones available.
50 std::set<std::string> kinds;
51 for (HloComputation* computation :
52 module_->computations(execution_threads)) {
53 for (HloInstruction* instruction : computation->instructions()) {
54 if (instruction->opcode() == HloOpcode::kDomain) {
55 TF_RET_CHECK(instruction->user_side_metadata().Kind() ==
56 instruction->operand_side_metadata().Kind())
57 << instruction->ToString();
58 kinds.insert(std::string(instruction->user_side_metadata().Kind()));
59 }
60 }
61 }
62 verifier_->kinds_.insert(verifier_->kinds_.end(), kinds.begin(),
63 kinds.end());
64 }
65 return OkStatus();
66 }
67
Run(const absl::flat_hash_set<absl::string_view> & execution_threads)68 Status HloDomainVerifier::RunContext::Run(
69 const absl::flat_hash_set<absl::string_view>& execution_threads) {
70 VLOG(4) << "Running HLO Domain Verifier";
71 TF_RETURN_IF_ERROR(PopulateDomainKinds(execution_threads));
72 for (HloComputation* computation : module_->computations(execution_threads)) {
73 for (auto& kind : verifier_->kinds_) {
74 // First create the domain instruction sets. A domain instruction set is
75 // the set of instructions whose edges never cross a kDomain instruction.
76 TF_ASSIGN_OR_RETURN(std::unique_ptr<HloDomainMap> domain_map,
77 HloDomainMap::Create(computation, kind));
78 // Verify every domain populated within the map.
79 for (auto& domain : domain_map->GetDomains()) {
80 TF_RETURN_IF_ERROR(VerifyDomain(*domain).status());
81 }
82 }
83 }
84 return OkStatus();
85 }
86
Run(HloModule * module,const absl::flat_hash_set<absl::string_view> & execution_threads)87 StatusOr<bool> HloDomainVerifier::Run(
88 HloModule* module,
89 const absl::flat_hash_set<absl::string_view>& execution_threads) {
90 RunContext run_context(module, this);
91 TF_RETURN_IF_ERROR(run_context.Run(execution_threads));
92 return false;
93 }
94
VerifyDomain(const DomainMetadata::Domain & domain)95 StatusOr<const DomainMetadata*> HloDomainVerifier::VerifyDomain(
96 const DomainMetadata::Domain& domain) {
97 const DomainMetadata* ref_metadata = nullptr;
98 VLOG(4) << "Reach set:";
99 for (HloInstruction* instruction : domain.instructions) {
100 VLOG(4) << " " << instruction->name();
101 }
102 VLOG(4) << " Domains:";
103 for (HloInstruction* instruction : domain.enter_domains) {
104 const DomainMetadata& meta = instruction->user_side_metadata();
105 VLOG(4) << " User side: " << instruction->name();
106 VLOG(4) << " " << meta.ToString();
107 if (ref_metadata == nullptr) {
108 ref_metadata = &meta;
109 } else {
110 TF_RET_CHECK(meta.Matches(*ref_metadata))
111 << "Metadata mismatch at instruction " << instruction->name() << " : "
112 << meta.ToString() << " vs " << ref_metadata->ToString();
113 }
114 }
115 for (HloInstruction* instruction : domain.exit_domains) {
116 const DomainMetadata& meta = instruction->operand_side_metadata();
117 VLOG(4) << " Operand side: " << instruction->name();
118 VLOG(4) << " " << meta.ToString();
119 if (ref_metadata == nullptr) {
120 ref_metadata = &meta;
121 } else {
122 TF_RET_CHECK(meta.Matches(*ref_metadata))
123 << "Metadata mismatch at instruction " << instruction->name() << " : "
124 << meta.ToString() << " vs " << ref_metadata->ToString();
125 }
126 }
127 return ref_metadata;
128 }
129
130 } // namespace xla
131