The list of changes include: . Added a new test manually generated from the nmt benchmark to stress test the optimization and to make sure the optimization can benefit a particular conditional in the benchmark. . In the Run() method of the pass, add a new variable "cleanup_changed", so that the variable "changed" only reflects whether conditional code motion has occured: the previous implementation mixed cleanup changes with the code motion changes, causing inconsistencies in various decision making in the Pass:Run, which uses the variable "changed" to keep track of whether code motion had occured. . In the Run method, fix bug in support for cloning conditional branches that are shared: the old implementation tried to re-analyze the optimization but with a boundary that has not be updated with the new cloned instructions. The fix adopts a different approach --- it passes in a CloneContext object to save the mappings between cloned instructions and then uses the info to update the analysis results directly, without re-analyzing anything. . In the Run() method of the pass, instead of breaking out after the first benefitial optimization if found, analyze all the move-out and move-in options, compare their benefits, and then selecting the mose beneficial one as the final decision. Add support for allowing multiple streams of boundaries to move into branches . Adjust the cost model (WorthHoisting and ReusesAcross methods) to allow more accurate calculation of profitability (for nmt). . Add support for moving instructions that have multiple users or operands, by keeping track of the number of times each instruction/boundary has been visited. Fixed problems in interfering code modifications that invalidate previously computed boundaries to be moved in later. Added a visited_count_ mapping to keep track of how many times each instruction has been visited due to different dependents, to allow previously moved instructions to be used in later transformations.Added a new test for safety constraint violations. PiperOrigin-RevId: 332551100 Change-Id: I61cff48b7007c626d140a2fbf21ce7ee9504cb06
125 lines
5.2 KiB
C++
125 lines
5.2 KiB
C++
/* Copyright 2017 The TensorFlow Authors. All Rights Reserved.
|
|
|
|
Licensed under the Apache License, Version 2.0 (the "License");
|
|
you may not use this file except in compliance with the License.
|
|
You may obtain a copy of the License at
|
|
|
|
http://www.apache.org/licenses/LICENSE-2.0
|
|
|
|
Unless required by applicable law or agreed to in writing, software
|
|
distributed under the License is distributed on an "AS IS" BASIS,
|
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
See the License for the specific language governing permissions and
|
|
limitations under the License.
|
|
==============================================================================*/
|
|
|
|
#ifndef TENSORFLOW_COMPILER_XLA_SERVICE_CONDITIONAL_CODE_MOTION_H_
|
|
#define TENSORFLOW_COMPILER_XLA_SERVICE_CONDITIONAL_CODE_MOTION_H_
|
|
|
|
#include "absl/strings/string_view.h"
|
|
#include "tensorflow/compiler/xla/service/hlo_module.h"
|
|
#include "tensorflow/compiler/xla/service/hlo_pass_interface.h"
|
|
#include "tensorflow/compiler/xla/statusor.h"
|
|
|
|
namespace xla {
|
|
|
|
namespace conditional_opt {
|
|
// At the conceptual level, a boundary can be thought of as representing a
|
|
// single virtual operation, except this virtual operation is conditionally
|
|
// instantiated into different concrete operations at each conditional branch.
|
|
// So a boundary is mapped to a single concrete operation if it is outside of
|
|
// conditional branches, and is mapped to a list of instructions if inside the
|
|
// branches. This data structure therefore allows a common data structure
|
|
// representation of the instructions to be moved, whether they are inside or
|
|
// outside of the branches. Subsequently, it allows a common implementation
|
|
// basis to be used for both moving instructions out of and for moving them
|
|
// inside branches.
|
|
class Boundary {
|
|
public:
|
|
enum class Position { kInsideBranch, kOutsideBranch, kUndefined };
|
|
Boundary() : position_(Position::kUndefined) {}
|
|
explicit Boundary(Position p) : position_(p) {}
|
|
std::vector<HloInstruction*>& mutable_operands() { return operands_; }
|
|
const std::vector<HloInstruction*>& operands() const { return operands_; }
|
|
bool IsInsideBranch() const { return position_ == Position::kInsideBranch; }
|
|
bool IsOutsideBranch() const { return position_ == Position::kOutsideBranch; }
|
|
Position GetPosition() const { return position_; }
|
|
bool IsEmpty() const { return operands_.empty(); }
|
|
std::string ToString() const {
|
|
std::string res;
|
|
for (HloInstruction* op : operands_) {
|
|
res += op->ToString() + ";";
|
|
}
|
|
return res;
|
|
}
|
|
bool operator==(const Boundary& that) {
|
|
return ContainersEqual(operands_, that.operands_);
|
|
}
|
|
|
|
private:
|
|
// Boundary instructions in the conditional branches, one from each branch
|
|
// of the conditional; or a single operand from outside the conditional.
|
|
std::vector<HloInstruction*> operands_;
|
|
Position position_;
|
|
};
|
|
|
|
// HLO pass that moves identical ops in/out of conditional.
|
|
// - The definition of identical are the shape of the operands are identical
|
|
// and their properties are identical.
|
|
// - Only the identical ops that won't share operands with other ops will
|
|
// be moved out of conditional.
|
|
class ConditionalCodeMotion : public HloModulePass {
|
|
public:
|
|
// If is_layout_sensitive is true, then the hoist process preserves layout
|
|
// during identical comparison. Otherwise, layout is ignored.
|
|
explicit ConditionalCodeMotion(bool is_layout_sensitive,
|
|
bool pursue_full_conditional_code_motion)
|
|
: is_layout_sensitive_(is_layout_sensitive),
|
|
pursue_full_conditional_code_motion_(
|
|
pursue_full_conditional_code_motion) {}
|
|
absl::string_view name() const override { return "conditional-code-motion"; }
|
|
StatusOr<bool> Run(HloModule* module) override;
|
|
|
|
// Optimization decision for each boundary of the conditional instruction.
|
|
class Decision {
|
|
public:
|
|
enum class Direction : uint8 {
|
|
kMoveOutOfBranch,
|
|
kMoveIntoBranch,
|
|
kNoChange
|
|
};
|
|
|
|
public:
|
|
Decision(Direction direction, int benefit)
|
|
: direction_(direction), benefit_(benefit) {}
|
|
Direction GetDirection() const { return direction_; }
|
|
int GetBenefit() const { return benefit_; }
|
|
|
|
private:
|
|
Direction direction_;
|
|
int benefit_;
|
|
};
|
|
// If the optimization decision is NO_CHANGE, new_boundary is set to nullptr;
|
|
// otherwise, it is set to the new boundary after proposed optimization.
|
|
virtual Decision ConsiderCodeMotion(
|
|
HloInstruction* conditional, const Boundary& cur_boundary,
|
|
std::vector<Boundary>& to_move, std::vector<Boundary>& new_boundaries,
|
|
absl::flat_hash_map<HloInstruction*, int>& visited_count);
|
|
|
|
private:
|
|
const bool is_layout_sensitive_;
|
|
const bool pursue_full_conditional_code_motion_;
|
|
|
|
StatusOr<bool> MoveInstructionOut(HloInstruction* conditional,
|
|
std::vector<Boundary>& to_move_out,
|
|
std::vector<Boundary>& new_boundaries);
|
|
StatusOr<bool> MoveInstructionIn(HloInstruction* conditional,
|
|
std::vector<Boundary>& to_move_in,
|
|
std::vector<Boundary>& new_boundaries);
|
|
};
|
|
} // namespace conditional_opt
|
|
|
|
} // namespace xla
|
|
|
|
#endif // TENSORFLOW_COMPILER_XLA_SERVICE_CONDITIONAL_CODE_MOTION_H_
|