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_GPU_INSTRUCTION_FUSION_H_ 17 #define TENSORFLOW_COMPILER_XLA_SERVICE_GPU_INSTRUCTION_FUSION_H_ 18 19 #include "tensorflow/compiler/xla/service/hlo_instruction.h" 20 #include "tensorflow/compiler/xla/service/instruction_fusion.h" 21 22 namespace xla { 23 namespace gpu { 24 25 class GpuInstructionFusion : public InstructionFusion { 26 public: GpuInstructionFusion(bool may_duplicate)27 explicit GpuInstructionFusion(bool may_duplicate) 28 : InstructionFusion(GpuInstructionFusion::IsExpensive, may_duplicate) {} 29 30 // Maximum number of operands plus outputs allowed on a single fusion node. 31 // Exposed publicly mainly for tests. 32 static constexpr int64 kMaxOperandsAndOutputsPerFusion = 64; 33 34 // Determines whether the combination of `a` and `b` into a (possibly 35 // multi-output) fusion would be "too large" -- i.e., have more operands and 36 // outputs than is allowed. 37 // 38 // `ShouldFuse` and `ShouldFuseIntoMultiOutput` call this; it's public so that 39 // other fusion passes (e.g. GPU multi-output fusion) can also call this. 40 static bool FusionWouldBeTooLarge(const HloInstruction* a, 41 const HloInstruction* b); 42 43 static bool IsExpensive(const HloInstruction& instruction); 44 45 bool ShouldFuse(HloInstruction* consumer, int64 operand_index) override; 46 47 bool ShouldFuseIntoMultiOutput(HloInstruction* consumer, 48 int64 operand_index) override; 49 50 HloInstruction::FusionKind ChooseKind( 51 const HloInstruction* producer, const HloInstruction* consumer) override; 52 53 private: 54 // This method is called by ShouldFuse() to do all the computationally 55 // inexpensive checks whether we should fuse the operand into 'consumer'. 56 bool ShouldFuseInexpensiveChecks(HloInstruction* consumer, 57 int64 operand_index); 58 }; 59 60 } // namespace gpu 61 } // namespace xla 62 63 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_GPU_INSTRUCTION_FUSION_H_ 64