1 /* Copyright 2017 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 #ifndef TENSORFLOW_COMPILER_XLA_SERVICE_COPY_INSERTION_H_ 17 #define TENSORFLOW_COMPILER_XLA_SERVICE_COPY_INSERTION_H_ 18 19 #include "tensorflow/compiler/xla/service/hlo_alias_analysis.h" 20 #include "tensorflow/compiler/xla/service/hlo_computation.h" 21 #include "tensorflow/compiler/xla/service/hlo_instruction.h" 22 #include "tensorflow/compiler/xla/service/hlo_module.h" 23 #include "tensorflow/compiler/xla/service/hlo_pass_interface.h" 24 25 namespace xla { 26 27 // Copy insertion is a legalization HLO pass which inserts copies (kCopy 28 // instructions) to eliminate several kinds of problems in the HLO module. 29 // 30 // (1) Entry parameter or a constant live out of the entry computation. Entry 31 // computation arguments and constants have different lifetimes than the 32 // computation result and cannot share the same allocation. Parameters and 33 // constants live out of non-entry computations do not need copies. 34 // 35 // (2) Different values which are simultaneously live and which must be held 36 // in the same buffer. This can occur in while bodies. Specifically, the 37 // while loop state (the arguments to the while instruction) is updated 38 // in-place and the update may clobber the value from the previous 39 // iteration before the previous value is dead. Computations called from 40 // kCall instructions do not need such copies because kCall has no update 41 // in-place semantics. 42 // 43 // (3) The buffer set of the root instruction of the entry computation must be 44 // unambiguous and distinct. That is, InstructionAliasSet::IsAmbiguous and 45 // InstructionAliasSet::IsDistinct return true. 46 class CopyInsertion : public HloModulePass { 47 public: name()48 absl::string_view name() const override { return "copy-insertion"; } 49 static constexpr int64_t kUseRegionAnalysisLimit = 0; 50 51 // backend specific function that decides whether an instruction 52 // can share buffer with its operand. 53 // 54 // TODO(b/80315712): Find a better way to tell whether a fusion can share 55 // buffer. 56 explicit CopyInsertion( 57 const HloDataflowAnalysis::CanShareBuffer& can_share_buffer = nullptr, 58 int64_t use_region_based_live_range_analysis = kUseRegionAnalysisLimit) can_share_buffer_(can_share_buffer)59 : can_share_buffer_(can_share_buffer), 60 use_region_based_live_range_analysis_( 61 use_region_based_live_range_analysis) {} 62 63 // Run the pass on the given module. Returns whether the module was changed 64 // (copies were inserted). 65 using HloPassInterface::Run; 66 StatusOr<bool> Run( 67 HloModule* module, 68 const absl::flat_hash_set<absl::string_view>& execution_threads) override; 69 70 // Try to remove as many copies from the module as possible without 71 // introducing live range interference. Only copy instructions that are 72 // eligible for copy elision are considered for removal. 73 // If check_live_range_ordering is true, check that live ranges are ordered 74 // in all the existing aliased buffers. 75 Status RemoveUnnecessaryCopies( 76 HloOrdering* ordering, HloModule* module, 77 bool check_live_range_ordering = false, 78 const absl::flat_hash_set<absl::string_view>& execution_threads = {}); 79 80 // Add copies to address special constraints on the roots of computations not 81 // related to live range interference: 82 // 83 // (1) Entry computation root must be unambiguous and distinct. 84 // 85 // (2) Any computation called by a kCall instruction must have an 86 // unambiguous root. 87 // 88 // (3) Constants and parameters cannot be live out of the entry computation 89 // 90 Status AddSpecialCaseCopies( 91 HloModule* module, 92 const absl::flat_hash_set<absl::string_view>& execution_threads = {}); 93 94 protected: 95 // Override which requires the caller to pass in a call graph. 96 virtual Status AddSpecialCaseCopies( 97 const CallGraph& call_graph, 98 const absl::flat_hash_set<absl::string_view>& execution_threads, 99 HloModule* module); 100 101 // Add copies for conditional instructions. 102 virtual Status AddCopiesForConditional(const HloAliasAnalysis& alias_analysis, 103 HloInstruction* conditional); 104 105 // Backend specific function that decides whether an instruction can share 106 // buffer with its operand. 107 HloDataflowAnalysis::CanShareBuffer can_share_buffer_; 108 109 private: 110 Status AddCopiesToResolveInterference( 111 HloModule* module, 112 const absl::flat_hash_set<absl::string_view>& execution_threads); 113 int64_t use_region_based_live_range_analysis_; 114 }; 115 116 } // namespace xla 117 118 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_COPY_INSERTION_H_ 119