1 /* Copyright 2021 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 // This file implements logic for lowering LHLO GPU dialect to TFRT CUDA
17 // dialect.
18
19 #include "tensorflow/compiler/mlir/tfrt/transforms/lhlo_gpu_to_tfrt_gpu/gpu_passes.h"
20
21 #include <memory>
22 #include <utility>
23
24 #include "mlir-hlo/Dialect/mhlo/IR/lhlo_ops.h"
25 #include "mlir/Dialect/GPU/GPUDialect.h"
26 #include "mlir/Dialect/MemRef/IR/MemRef.h"
27 #include "mlir/Dialect/StandardOps/IR/Ops.h"
28 #include "mlir/IR/BuiltinOps.h"
29 #include "mlir/IR/MLIRContext.h"
30 #include "mlir/IR/PatternMatch.h"
31 #include "mlir/IR/Types.h"
32 #include "mlir/Pass/Pass.h"
33 #include "mlir/Rewrite/PatternApplicator.h"
34 #include "mlir/Transforms/DialectConversion.h"
35 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
36 #include "llvm/ADT/ArrayRef.h"
37 #include "tensorflow/compiler/mlir/tfrt/transforms/lhlo_gpu_to_tfrt_gpu/PassDetail.h"
38 #include "tensorflow/compiler/mlir/tfrt/transforms/lhlo_gpu_to_tfrt_gpu/ccl_pattern.h"
39 #include "tensorflow/compiler/mlir/tfrt/transforms/lhlo_gpu_to_tfrt_gpu/gemm_pattern.h"
40 #include "tensorflow/compiler/xla/service/gpu/xlir_ops.h"
41 #include "tfrt/gpu/kernels/gpu_ops.h" // from @tf_runtime
42 #include "tfrt/gpu/pass/pass.h" // from @tf_runtime
43 #include "tfrt/basic_kernels/opdefs/basic_kernels.h" // from @tf_runtime
44
45 namespace tensorflow {
46 namespace {
47
48 struct LmhloGpuAsyncConversionPass
49 : public LmhloGpuAsyncConversionPassBase<LmhloGpuAsyncConversionPass> {
50 private:
runOnFunctiontensorflow::__anona52e31080111::LmhloGpuAsyncConversionPass51 void runOnFunction() override {
52 auto* context = &getContext();
53
54 TypeConverter converter;
55 converter.addConversion([](Type type) { return type; });
56 auto buffer_type = tfrt::gpu::BufferType::get(context);
57 converter.addConversion([&](BaseMemRefType) { return buffer_type; });
58
59 ConversionTarget target(*context);
60 target.addIllegalDialect<lmhlo_gpu::LmhloGpuDialect>();
61 target.addLegalDialect<tfrt::compiler::TFRTDialect, tfrt::gpu::GpuDialect,
62 xla::gpu::XlirDialect>();
63 target.addDynamicallyLegalOp<FuncOp>([&](FuncOp op) {
64 return converter.isSignatureLegal(op.getType()) &&
65 converter.isLegal(&op.body());
66 });
67 target.addDynamicallyLegalOp<tfrt::gpu::conversion::AsyncExecuteOp>(
68 [&](tfrt::gpu::conversion::AsyncExecuteOp op) {
69 return converter.isLegal(&op.body());
70 });
71
72 RewritePatternSet patterns(context);
73 populateCclConversionPattern(patterns);
74 populateGemmConversionPattern(patterns);
75 populateFuncOpTypeConversionPattern(patterns, converter);
76
77 ConversionTarget wrap_target(*context);
78 wrap_target.addLegalDialect<lmhlo_gpu::LmhloGpuDialect>();
79 wrap_target.addLegalOp<lmhlo::AllReduceOp>();
80 tfrt::gpu::populateGpuAsyncConversionPatterns(patterns, converter,
81 wrap_target);
82
83 if (failed(applyPartialConversion(getOperation(), target,
84 std::move(patterns))))
85 return signalPassFailure();
86 }
87 };
88
89 struct AsyncGpuTfrtConversionPass
90 : public AsyncGpuTfrtConversionPassBase<AsyncGpuTfrtConversionPass> {
91 private:
runOnFunctiontensorflow::__anona52e31080111::AsyncGpuTfrtConversionPass92 void runOnFunction() override {
93 auto* context = &getContext();
94
95 ConversionTarget target(*context);
96 target.addLegalDialect<tfrt::compiler::TFRTDialect, tfrt::gpu::GpuDialect,
97 xla::gpu::XlirDialect>();
98
99 RewritePatternSet patterns(context);
100 tfrt::gpu::populateTfrtConversionPatterns(patterns, target);
101
102 if (failed(applyPartialConversion(getOperation(), target,
103 std::move(patterns))))
104 return signalPassFailure();
105 }
106 };
107
108 } // namespace
109
createLmhloGpuAsyncConversionPass()110 std::unique_ptr<FunctionPass> createLmhloGpuAsyncConversionPass() {
111 return std::make_unique<LmhloGpuAsyncConversionPass>();
112 }
113
createAsyncGpuTfrtConversionPass()114 std::unique_ptr<FunctionPass> createAsyncGpuTfrtConversionPass() {
115 return std::make_unique<AsyncGpuTfrtConversionPass>();
116 }
117
118 } // namespace tensorflow
119