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::unique_ptr<OperationPass<ModuleOp>> 35 CreateEmbedTFFrameworkFunctionAndAllocPass(); 36 37 // Pass to convert std.assert operations to calls to tf_framework.report_error 38 // and create the required control flow to abort the function on failed 39 // execution. 40 std::unique_ptr<OperationPass<ModuleOp>> CreateEmbedTFFrameworkAssertPass(); 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 for applying LLVM legalization patterns. 50 std::unique_ptr<OperationPass<ModuleOp>> CreateTFKernelToLLVMPass( 51 mlir::StringRef blob_annotation = {}); 52 53 // Pass to tranform shape computations in shape dialect to standard and scf 54 // using memref descriptors. 55 std::unique_ptr<OperationPass<ModuleOp>> CreateShapeToDescriptorsPass(); 56 57 // Pass to tranform hlo-level computations on values to their corresponding 58 // parts on buffers. 59 std::unique_ptr<OperationPass<ModuleOp>> CreateHloBufferizePass(); 60 61 // Pass to tranform computations on values to their corresponding parts on 62 // buffers. 63 std::unique_ptr<OperationPass<ModuleOp>> CreateFinalBufferizePass(); 64 65 // Pass to convert scf::ParallelOp to scf::ForOp. 66 std::unique_ptr<FunctionPass> CreateParallelLoopsToSequential(); 67 68 // Pass to annotate GPU Module with its PTX. 69 std::unique_ptr<OperationPass<gpu::GPUModuleOp>> CreateGpuKernelToBlobPass( 70 mlir::StringRef blob_annotation = {}, 71 ArrayRef<std::string> architectures = {}, bool generate_fatbin = true, 72 bool print_ptx = false, bool enable_ftz = false); 73 74 // Pass to propagate tensorflow runtime ABI knowledge across kernel boundaries. 75 std::unique_ptr<FunctionPass> CreatePropagateTfAbiKnowledgeToKernels(); 76 77 // Pass to propagate shape equalities across kernel boundaries. 78 std::unique_ptr<FunctionPass> CreatePropagateShapeKnowledgeToKernels(); 79 80 // Pass to print content of memrefs. 81 std::unique_ptr<FunctionPass> CreateEmbedMemRefPrintsPass(); 82 83 /// Greedily maps loops to GPU hardware dimensions. 84 std::unique_ptr<mlir::FunctionPass> CreateMapParallelLoopsPass(); 85 86 /// We need to direct fusion to the inner loops. This cannot be done with 87 /// a passmanager alone ATM, as nested pass managers require operations to 88 /// be closed from above. 89 std::unique_ptr<mlir::FunctionPass> CreateFuseInnerParallelLoopsPass(); 90 91 /// Pass that transforms gpu modules in standard dialect to NNVM. 92 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>> 93 CreateGpuKernelToNvvmPass(); 94 95 /// Pass that transforms gpu modules in standard dialect to ROCDL. 96 std::unique_ptr<OperationPass<mlir::gpu::GPUModuleOp>> 97 CreateGpuKernelToRocdlPass(); 98 99 } // namespace transforms 100 101 #define GEN_PASS_REGISTRATION 102 #include "tensorflow/compiler/mlir/tools/kernel_gen/transforms/kernel_gen_passes.h.inc" 103 104 } // namespace kernel_gen 105 } // namespace mlir 106 107 #endif // TENSORFLOW_COMPILER_MLIR_TOOLS_KERNEL_GEN_TRANSFORMS_PASSES_H_ 108