1 /* Copyright 2020 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_MLIR_TOOLS_KERNEL_GEN_TRANSFORMS_PASSES_H_ 17 #define TENSORFLOW_COMPILER_MLIR_TOOLS_KERNEL_GEN_TRANSFORMS_PASSES_H_ 18 19 #include <memory> 20 21 #include "mlir/Dialect/GPU/GPUDialect.h" // from @llvm-project 22 #include "mlir/Dialect/LLVMIR/LLVMDialect.h" // from @llvm-project 23 #include "mlir/IR/BuiltinOps.h" // from @llvm-project 24 #include "mlir/Pass/Pass.h" // from @llvm-project 25 26 namespace mlir { 27 namespace kernel_gen { 28 namespace tf_framework { 29 30 // Pass to replace some of the Standard ops with TF Framework ops. 31 // * adds tf_framework::OpKernelContextType argument to the function 32 // * std.alloc becomes tf_framework.alloc_raw 33 // * std.dealloc becomes tf_framework.dealloc_raw 34 // * std.assert becomes tf_framework.assert 35 std::unique_ptr<OperationPass<ModuleOp>> CreateEmbedTFFrameworkPass(); 36 37 // Pass to convert tf_framework.assert operations to calls to 38 // tf_framework.report_error and create the required control flow to abort the 39 // function on failed execution. 40 std::unique_ptr<OperationPass<ModuleOp>> CreateRewriteTFFrameworkAssert(); 41 42 } // namespace tf_framework 43 44 namespace transforms { 45 46 // Pass to find and annotate candidates for buffer reuse. 47 std::unique_ptr<FunctionPass> CreateBufferReusePass(); 48 49 // Pass to rewrite all TF operations to JIT invocations through the TF 50 // framework. 51 std::unique_ptr<FunctionPass> CreateTFToJITInvocationPass( 52 llvm::ArrayRef<std::string> architectures = {}, 53 llvm::ArrayRef<int64_t> tile_sizes = {}, 54 llvm::ArrayRef<int64_t> unroll_factors = {}, int64_t max_supported_rank = 5, 55 bool enable_ftz = false, bool cpu_codegen = false); 56 57 // Pass for applying LLVM legalization patterns. 58 std::unique_ptr<OperationPass<ModuleOp>> CreateTFKernelToLLVMPass( 59 mlir::StringRef blob_annotation = {}); 60 61 // Pass to tranform shape computations in shape dialect to standard and scf 62 // using memref descriptors. 63 std::unique_ptr<OperationPass<ModuleOp>> CreateShapeToDescriptorsPass(); 64 65 // Pass to tranform compute computations (hlo and linalg) on values to their 66 // corresponding counterparts on buffers. Also bufferizes function signatures. 67 std::unique_ptr<OperationPass<ModuleOp>> CreateComputeOpAndFuncBufferizePass(); 68 69 // Pass to tranform computations on values to their corresponding parts on 70 // buffers. 71 std::unique_ptr<OperationPass<ModuleOp>> CreateFinalBufferizePass(); 72 73 // Pass to replace unsigned types with signless integers. 74 std::unique_ptr<OperationPass<ModuleOp>> CreateConvertToSignlessPass(); 75 76 // Pass to convert scf::ParallelOp to scf::ForOp. 77 std::unique_ptr<FunctionPass> CreateParallelLoopsToSequential(); 78 79 // Pass to annotate GPU Module with its PTX. 80 std::unique_ptr<OperationPass<gpu::GPUModuleOp>> CreateGpuKernelToBlobPass( 81 mlir::StringRef blob_annotation = {}, 82 ArrayRef<std::string> architectures = {}, bool print_ptx = false, 83 bool print_llvmir = false, bool enable_ftz = false); 84 85 // Pass to propagate tensorflow runtime ABI knowledge across kernel boundaries. 86 std::unique_ptr<FunctionPass> CreatePropagateTfAbiKnowledgeToKernels(); 87 88 // Pass to propagate shape equalities across kernel boundaries. 89 std::unique_ptr<FunctionPass> CreatePropagateShapeKnowledgeToKernels(); 90 91 // Pass to print content of memrefs. 92 std::unique_ptr<FunctionPass> CreateEmbedMemRefPrintsPass(); 93 94 /// Greedily maps loops to GPU hardware dimensions. 95 std::unique_ptr<mlir::FunctionPass> CreateMapParallelLoopsPass(); 96 97 /// We need to direct fusion to the inner loops. This cannot be done with 98 /// a passmanager alone ATM, as nested pass managers require operations to 99 /// be closed from above. 100 std::unique_ptr<mlir::FunctionPass> CreateFuseInnerParallelLoopsPass(); 101 102 /// Pass that transforms gpu modules in standard dialect to NNVM. 103 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>> 104 CreateGpuKernelToNvvmPass(); 105 106 /// Pass that transforms gpu modules in standard dialect to ROCDL. 107 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>> 108 CreateGpuKernelToRocdlPass(); 109 110 // Pass to lower index cast on tensors to tensor dialect. 111 std::unique_ptr<FunctionPass> CreateLowerIndexCastPass(); 112 113 // Pass to simplify shape ops. 114 std::unique_ptr<FunctionPass> CreateShapeSimplification(); 115 116 // Pass to create vectorized code for CPU. 117 std::unique_ptr<FunctionPass> CreateVectorizationPass(); 118 119 // Pass to remove unneeded code generated in VectorizationPass. 120 std::unique_ptr<FunctionPass> CreateVectorizationCleanupPass(); 121 122 // Pass to remove copies which are consumed by a GenericOp. 123 std::unique_ptr<FunctionPass> CreateCopyCleanupPass(); 124 125 } // namespace transforms 126 127 #define GEN_PASS_REGISTRATION 128 #include "tensorflow/compiler/mlir/tools/kernel_gen/transforms/kernel_gen_passes.h.inc" 129 130 } // namespace kernel_gen 131 } // namespace mlir 132 133 #endif // TENSORFLOW_COMPILER_MLIR_TOOLS_KERNEL_GEN_TRANSFORMS_PASSES_H_ 134