| 233 | } |
| 234 | |
| 235 | Status HloSchedule::Verify() const { |
| 236 | VLOG(2) << "VerifySchedule()"; |
| 237 | XLA_VLOG_LINES(2, ToString()); |
| 238 | |
| 239 | // Verify schedule contains exactly the same set of non-fusion computations as |
| 240 | // module currently does. |
| 241 | std::vector<HloComputation*> nonfusion_computations = |
| 242 | module_->MakeNonfusionComputations(); |
| 243 | TF_RET_CHECK(nonfusion_computations.size() == sequences_.size()) |
| 244 | << "Schedule has " << sequences_.size() << " sequences, but module has " |
| 245 | << nonfusion_computations.size() << " non-fusion computations"; |
| 246 | for (const HloComputation* computation : nonfusion_computations) { |
| 247 | TF_RET_CHECK(sequences_.contains(computation->unique_id())) |
| 248 | << "Computation " << computation->name() |
| 249 | << " missing from HLO schedule."; |
| 250 | } |
| 251 | |
| 252 | // For each computation verify the set of instructions is the same and that |
| 253 | // each dependency and control edge is honored. |
| 254 | for (const HloComputation* computation : nonfusion_computations) { |
| 255 | absl::flat_hash_map<const HloInstruction*, int> instruction_position; |
| 256 | int pos = 0; |
| 257 | for (const HloInstruction* instruction : |
| 258 | sequence(computation).instructions()) { |
| 259 | TF_RET_CHECK(instruction_position.insert({instruction, pos}).second) |
| 260 | << "Instruction " << instruction->name() |
| 261 | << " appears more than once in the schedule"; |
| 262 | pos++; |
| 263 | } |
| 264 | |
| 265 | TF_RET_CHECK(instruction_position.size() == |
| 266 | computation->instruction_count()) |
| 267 | << "Schedule for computation " << computation->name() << " has " |
| 268 | << instruction_position.size() << " instructions, expected " |
| 269 | << computation->instruction_count(); |
| 270 | for (const HloInstruction* instruction : computation->instructions()) { |
| 271 | TF_RET_CHECK(instruction_position.contains(instruction)) |
| 272 | << "Instruction " << instruction->name() << " is not in schedule"; |
| 273 | } |
| 274 | |
| 275 | for (const HloInstruction* instruction : computation->instructions()) { |
| 276 | for (const HloInstruction* operand : instruction->operands()) { |
| 277 | TF_RET_CHECK(instruction_position.at(operand) < |
| 278 | instruction_position.at(instruction)) |
| 279 | << "Instruction " << instruction->name() |
| 280 | << " is not scheduled after its operand " << operand->name(); |
| 281 | } |
| 282 | |
| 283 | for (const HloInstruction* pred : instruction->control_predecessors()) { |
| 284 | TF_RET_CHECK(instruction_position.at(pred) < |
| 285 | instruction_position.at(instruction)) |
| 286 | << "Instruction " << instruction->name() |
| 287 | << " is not scheduled after its control predecessor " |
| 288 | << pred->name(); |
| 289 | } |
| 290 | } |
| 291 | } |
| 292 | |