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 #ifndef TENSORFLOW_COMPILER_XLA_SERVICE_WHILE_LOOP_INVARIANT_CODE_MOTION_H_ 17 #define TENSORFLOW_COMPILER_XLA_SERVICE_WHILE_LOOP_INVARIANT_CODE_MOTION_H_ 18 19 #include "tensorflow/compiler/xla/service/compile_time_cap.h" 20 #include "tensorflow/compiler/xla/service/hlo_module.h" 21 #include "tensorflow/compiler/xla/service/hlo_pass_interface.h" 22 #include "tensorflow/compiler/xla/statusor.h" 23 24 namespace xla { 25 26 // HLO pass that rewrites while loops to hoist loop invariant instructions in 27 // the while body into the computation that contains the while instruction. 28 29 class WhileLoopInvariantCodeMotion : public HloModulePass { 30 public: 31 using ShapeSizeFunction = std::function<int64_t(const Shape&)>; 32 // If `hoist_constants` is true then constants are always hoisted out of while 33 // loop bodies. Otherwise they are only hoisted out if they enable other 34 // non-trivial computations to be hoisted out. 35 // 36 // Setting `hoist_constants` to false can be help if LICM is run in the mid 37 // level HLO pipeline because hoisting constants out of while loop bodies can 38 // break optimizations like constant folding. 39 // 40 // Setting `hoist_other` and `hoist_reshapes` to false can be used to hoist 41 // only constants. If provided, `hoist_size_inflation_ratio` will forbid 42 // hoisting instructions where the ratio of the size of the output(s) to the 43 // input(s) is larger than hoist_size_inflation_ratio. This is useful on 44 // platforms on which it's important to prevent blow-ups in memory size. 45 // 46 // If `hoist_reshapes` is true, then reshapes are allowed to be hoisted out of 47 // while loop body by themselves. Otherwise, they are only hoisted out if they 48 // enable other non-trivial computations to be hoisted out. 49 // 50 // Setting `hoist_reshapes` to false can be useful when LICM is run in the 51 // mid level HLO pipeline because the reshapes will often get fused with 52 // consumer instructions, and won't cost anything if not hoisted. However, 53 // any stand alone reshapes after fusion will benefit from hoisting. 54 explicit WhileLoopInvariantCodeMotion( 55 bool hoist_constants = false, bool hoist_reshapes = false, 56 bool hoist_other = true, 57 std::optional<float> hoist_size_inflation_ratio = std::nullopt, 58 ShapeSizeFunction shape_size_function = ShapeUtil::ByteSizeOfElements) hoist_constants_(hoist_constants)59 : hoist_constants_(hoist_constants), 60 hoist_reshapes_(hoist_reshapes), 61 hoist_other_(hoist_other), 62 hoist_size_inflation_ratio_(hoist_size_inflation_ratio), 63 shape_size_function_(shape_size_function) {} 64 ~WhileLoopInvariantCodeMotion() override = default; 65 name()66 absl::string_view name() const override { 67 return "while-loop-invariant-code-motion"; 68 } 69 using HloPassInterface::Run; 70 StatusOr<bool> Run( 71 HloModule* module, 72 const absl::flat_hash_set<absl::string_view>& execution_threads) override; 73 74 private: 75 bool NotWorthHoistingIndividually(const HloInstruction& instruction); 76 StatusOr<bool> TryHoistingInvariantInstructionsFromWhileBody( 77 HloInstruction* while_instr, BoundNonLinearCompilerAnalysis* allowance); 78 79 bool hoist_constants_; 80 bool hoist_reshapes_; 81 bool hoist_other_; 82 std::optional<float> hoist_size_inflation_ratio_; 83 ShapeSizeFunction shape_size_function_; 84 }; 85 } // namespace xla 86 87 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_WHILE_LOOP_INVARIANT_CODE_MOTION_H_ 88