1 /* Copyright 2019 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_DYNAMIC_PADDER_H_ 17 #define TENSORFLOW_COMPILER_XLA_SERVICE_DYNAMIC_PADDER_H_ 18 19 #include <functional> 20 21 #include "tensorflow/compiler/xla/service/dynamic_dimension_inference.h" 22 #include "tensorflow/compiler/xla/service/hlo_pass_interface.h" 23 24 namespace xla { 25 26 // With bounded shapes, only part of the shape contains effective data and the 27 // rest contains padded data, whose value can be anything depending on the 28 // source of the data. When a bounded shape is directly consumed by an 29 // instruction that collapses dimensions (reduce for example), the padding data 30 // would affect result of the instruction. 31 // 32 // DynamicPadder uses DynamicDimensionInference to detect bounded shapes in a 33 // hlo module, it then inserts certain instructions to reset the padding into an 34 // identity value so that in doesn't affect the result of subsequent 35 // instruction. For example, it'd reset the padding to 0 before a bounded shape 36 // is consumed by a reduce-sum. 37 // 38 // Dynamic_padder removes dynamic shapes from the entry computation, and inserts 39 // custom calls (with dynamic shapes), which are lowered by specialized 40 // emitters: PadToStatic and SliceToDynamic. 41 42 // Each instruction can have one of the three modes in supporting dynamic 43 // lowering. 44 enum OpDynamismSupport { 45 // There is no support for dynamic lowering -- dynamic padder will make sure 46 // the input to that op has static bound by rewriting the op (e.g, extra space 47 // in reduce_sum will be padded with 0). 48 kNoSupport = 0, 49 // The op can take either dynamic input or static input. 50 kOptional, 51 // The op only has a dynamic lowering, dynamic padder will make sure the input 52 // to this op is in dynamic form. 53 kRequired, 54 }; 55 56 struct DynamicPadderOptions { 57 // Returns true if given instruction supports native dynamic lowering. If 58 // so, dynamic padder will not attempt to pad it. 59 using OpSupportsDynamismHandler = 60 std::function<OpDynamismSupport(HloInstruction*)>; 61 62 OpSupportsDynamismHandler op_supports_dynamism_handler = nullptr; 63 64 // Instruct how to inference output dynamic dimensions of custom calls. 65 DynamicDimensionInference::CustomCallInferenceHandler custom_call_handler = 66 nullptr; 67 68 // If `slice_dynamic_output` is true, insert 'slice_to_dynamic' ops to all 69 // outputs that are inferred to be dynamic. 70 bool slice_dynamic_output = true; 71 72 // Assertion generator for shape checks, only used if shape check mode is 73 // "runtime". 74 DynamicDimensionInference::AssertionGenerator assertion_generator; 75 76 // If set to true, pessimisticly assumes runtime shape checks may fail and 77 // returns a compile-time error. 78 DynamicDimensionInference::ShapeCheckMode shape_check_mode = 79 DynamicDimensionInference::ShapeCheckMode::kIgnore; 80 }; 81 82 class DynamicPadder : public HloModulePass { 83 public: 84 explicit DynamicPadder(DynamicPadderOptions options = DynamicPadderOptions()) options_(options)85 : options_(options) {} 86 name()87 absl::string_view name() const override { return "dynamic_padder"; } 88 89 using HloPassInterface::Run; 90 StatusOr<bool> Run( 91 HloModule* module, 92 const absl::flat_hash_set<absl::string_view>& execution_threads) override; 93 94 private: 95 DynamicPadderOptions options_; 96 }; 97 98 } // namespace xla 99 100 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_DYNAMIC_PADDER_H_ 101