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_TRANSPOSE_FOLDING_H_ 17 #define TENSORFLOW_COMPILER_XLA_SERVICE_TRANSPOSE_FOLDING_H_ 18 19 #include "tensorflow/compiler/xla/service/hlo_module.h" 20 #include "tensorflow/compiler/xla/service/hlo_pass_interface.h" 21 22 namespace xla { 23 24 // HLO pass that folds transpose operators into Dot operators, where the Dot 25 // operator is implemented by a GEMM kernel that can transpose its inputs. 26 class TransposeFolding : public HloModulePass { 27 public: 28 using OperandIndices = std::vector<int64>; 29 30 // Returns the set of foldable operands for a given HLO and some candidate 31 // operands. 32 using FoldableOperands = std::function<OperandIndices(const HloInstruction&, 33 const OperandIndices&)>; 34 using TransposableGemmOperandsFn = FoldableOperands; 35 using TransposableConvOperandsFn = FoldableOperands; 36 37 // Helper function to explicitly not fold transposes. NeverFoldTranspose(const HloInstruction &,const OperandIndices &)38 static OperandIndices NeverFoldTranspose(const HloInstruction&, 39 const OperandIndices&) { 40 return {}; 41 } 42 43 // Helper function to always fold transposes. AlwaysFoldTranspose(const HloInstruction &,const OperandIndices & ids)44 static OperandIndices AlwaysFoldTranspose(const HloInstruction&, 45 const OperandIndices& ids) { 46 return ids; 47 } 48 49 // transposable_gemm_operands returns the set of operands it wants to fold if 50 // the instruction argument is implemented as a GEMM kernel that supports 51 // transposing its arguments. 52 // 53 // transposable_conv_operands returns the set of operands it wants to fold if 54 // the instruction argument is implemented as a convolution that supports 55 // transposing its arguments. 56 explicit TransposeFolding( 57 TransposableGemmOperandsFn transposable_gemm_operands = 58 AlwaysFoldTranspose, 59 TransposableConvOperandsFn transposable_conv_operands = 60 AlwaysFoldTranspose); name()61 absl::string_view name() const override { return "transpose-folding"; } 62 63 StatusOr<bool> Run(HloModule* module) override; 64 65 private: 66 TransposableGemmOperandsFn transposable_gemm_operands_; 67 TransposableConvOperandsFn transposable_conv_operands_; 68 }; 69 70 } // namespace xla 71 72 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_TRANSPOSE_FOLDING_H_ 73