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/buffer_liveness.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 #include "tensorflow/core/lib/gtl/flatmap.h" 25 26 namespace xla { 27 28 // Copy insertion is a legalization HLO pass which inserts copies (kCopy 29 // instructions) to eliminate several kinds of problems in the HLO module. 30 // 31 // (1) Entry parameter or a constant live out of the entry computation. Entry 32 // computation arguments and constants have different lifetimes than the 33 // computation result and cannot share the same allocation. Parameters and 34 // constants live out of non-entry computations do not need copies. 35 // 36 // (2) Different values which are simultaneously live and which must be held 37 // in the same buffer. This can occur in while bodies. Specifically, the 38 // while loop state (the arguments to the while instruction) is updated 39 // in-place and the update may clobber the value from the previous 40 // iteration before the previous value is dead. Computations called from 41 // kCall instructions do not need such copies because kCall has no update 42 // in-place semantics. 43 // 44 // (3) The buffer set of the root instruction of the entry computation must be 45 // unambiguous and distinct. That is, InstructionAliasSet::IsAmbiguous and 46 // InstructionAliasSet::IsDistinct return true. 47 class CopyInsertion : public HloPassInterface { 48 public: 49 tensorflow::StringPiece name() const override { return "copy-insertion"; } 50 51 // Run the pass on the given module. Returns whether the module was changed 52 // (copies were inserted). 53 StatusOr<bool> Run(HloModule* module) override; 54 55 // The CPU and GPU backend need additional copies added due to deficiencies in 56 // buffer assignment. Specifically, copies are needed for constants live-out 57 // of computations, and for values which are live-in and live-out of the same 58 // computation. These copies are needed because buffer-assignment uses a 59 // computation-scoped analyis (TuplePointsToAnalysis) and has limited 60 // visibility across computation boundaries. This method adds these necessary 61 // copies. Returns whether the module was modified. 62 // 63 // TODO(b/62548313): Remove this when buffer assignment is module-scoped. 64 static StatusOr<bool> AddCopiesForBufferAssignment(HloModule* module); 65 }; 66 67 } // namespace xla 68 69 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_COPY_INSERTION_H_ 70
