1//===- GPUToNVVMPipeline.cpp - Test lowering to NVVM as a sink pass -------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This file implements a pass for testing the lowering to NVVM as a generally
10// usable sink pass.
11//
12//===----------------------------------------------------------------------===//
13
14#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
15#include "mlir/Conversion/ArithToLLVM/ArithToLLVM.h"
16#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h"
17#include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
18#include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
19#include "mlir/Conversion/IndexToLLVM/IndexToLLVM.h"
20#include "mlir/Conversion/MathToLLVM/MathToLLVM.h"
21#include "mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h"
22#include "mlir/Conversion/NVVMToLLVM/NVVMToLLVM.h"
23#include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h"
24#include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h"
25#include "mlir/Conversion/VectorToSCF/VectorToSCF.h"
26#include "mlir/Dialect/GPU/IR/GPUDialect.h"
27#include "mlir/Dialect/GPU/Pipelines/Passes.h"
28#include "mlir/Dialect/GPU/Transforms/Passes.h"
29#include "mlir/Dialect/MemRef/Transforms/Passes.h"
30#include "mlir/Pass/PassManager.h"
31#include "mlir/Pass/PassOptions.h"
32#include "mlir/Transforms/Passes.h"
33
34using namespace mlir;
35
36namespace {
37
38//===----------------------------------------------------------------------===//
39// Common pipeline
40//===----------------------------------------------------------------------===//
41void buildCommonPassPipeline(
42 OpPassManager &pm, const mlir::gpu::GPUToNVVMPipelineOptions &options) {
43 pm.addPass(pass: createConvertNVGPUToNVVMPass());
44 pm.addPass(pass: createGpuKernelOutliningPass());
45 pm.addPass(pass: createConvertVectorToSCFPass());
46 pm.addPass(pass: createSCFToControlFlowPass());
47 pm.addPass(pass: createConvertNVVMToLLVMPass());
48 pm.addPass(pass: createConvertFuncToLLVMPass());
49 pm.addPass(pass: memref::createExpandStridedMetadataPass());
50
51 GpuNVVMAttachTargetOptions nvvmTargetOptions;
52 nvvmTargetOptions.triple = options.cubinTriple;
53 nvvmTargetOptions.chip = options.cubinChip;
54 nvvmTargetOptions.features = options.cubinFeatures;
55 nvvmTargetOptions.optLevel = options.optLevel;
56 nvvmTargetOptions.cmdOptions = options.cmdOptions;
57 pm.addPass(pass: createGpuNVVMAttachTarget(options: nvvmTargetOptions));
58 pm.addPass(pass: createLowerAffinePass());
59 pm.addPass(pass: createArithToLLVMConversionPass());
60 ConvertIndexToLLVMPassOptions convertIndexToLLVMPassOpt;
61 convertIndexToLLVMPassOpt.indexBitwidth = options.indexBitWidth;
62 pm.addPass(pass: createConvertIndexToLLVMPass(options: convertIndexToLLVMPassOpt));
63 pm.addPass(pass: createCanonicalizerPass());
64 pm.addPass(pass: createCSEPass());
65}
66
67//===----------------------------------------------------------------------===//
68// GPUModule-specific stuff.
69//===----------------------------------------------------------------------===//
70void buildGpuPassPipeline(OpPassManager &pm,
71 const mlir::gpu::GPUToNVVMPipelineOptions &options) {
72 ConvertGpuOpsToNVVMOpsOptions opt;
73 opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv;
74 opt.indexBitwidth = options.indexBitWidth;
75 pm.addNestedPass<gpu::GPUModuleOp>(pass: createConvertGpuOpsToNVVMOps(options: opt));
76 pm.addNestedPass<gpu::GPUModuleOp>(pass: createCanonicalizerPass());
77 pm.addNestedPass<gpu::GPUModuleOp>(pass: createCSEPass());
78 pm.addNestedPass<gpu::GPUModuleOp>(pass: createReconcileUnrealizedCastsPass());
79}
80
81//===----------------------------------------------------------------------===//
82// Host Post-GPU pipeline
83//===----------------------------------------------------------------------===//
84void buildHostPostPipeline(OpPassManager &pm,
85 const mlir::gpu::GPUToNVVMPipelineOptions &options) {
86 GpuToLLVMConversionPassOptions opt;
87 opt.hostBarePtrCallConv = options.hostUseBarePtrCallConv;
88 opt.kernelBarePtrCallConv = options.kernelUseBarePtrCallConv;
89 pm.addPass(pass: createGpuToLLVMConversionPass(options: opt));
90
91 GpuModuleToBinaryPassOptions gpuModuleToBinaryPassOptions;
92 gpuModuleToBinaryPassOptions.compilationTarget = options.cubinFormat;
93 pm.addPass(pass: createGpuModuleToBinaryPass(options: gpuModuleToBinaryPassOptions));
94 pm.addPass(pass: createConvertMathToLLVMPass());
95 pm.addPass(pass: createCanonicalizerPass());
96 pm.addPass(pass: createCSEPass());
97 pm.addPass(pass: createReconcileUnrealizedCastsPass());
98}
99
100} // namespace
101
102void mlir::gpu::buildLowerToNVVMPassPipeline(
103 OpPassManager &pm, const GPUToNVVMPipelineOptions &options) {
104 // Common pipelines
105 buildCommonPassPipeline(pm, options);
106
107 // GPUModule-specific stuff
108 buildGpuPassPipeline(pm, options);
109
110 // Host post-GPUModule-specific stuff
111 buildHostPostPipeline(pm, options);
112}
113
114void mlir::gpu::registerGPUToNVVMPipeline() {
115 PassPipelineRegistration<GPUToNVVMPipelineOptions>(
116 "gpu-lower-to-nvvm-pipeline",
117 "The default pipeline lowers main dialects (arith, memref, scf, "
118 "vector, gpu, and nvgpu) to NVVM. It starts by lowering GPU code to the "
119 "specified compilation target (default is fatbin) then lowers the host "
120 "code.",
121 buildLowerToNVVMPassPipeline);
122}
123

source code of mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp