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/MemRefToLLVM/MemRefToLLVM.h" 22 #include "mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h" 23 #include "mlir/Conversion/NVVMToLLVM/NVVMToLLVM.h" 24 #include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h" 25 #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" 26 #include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.h" 27 #include "mlir/Conversion/VectorToSCF/VectorToSCF.h" 28 #include "mlir/Dialect/Func/IR/FuncOps.h" 29 #include "mlir/Dialect/GPU/IR/GPUDialect.h" 30 #include "mlir/Dialect/GPU/Pipelines/Passes.h" 31 #include "mlir/Dialect/GPU/Transforms/Passes.h" 32 #include "mlir/Dialect/LLVMIR/LLVMDialect.h" 33 #include "mlir/Dialect/Linalg/Passes.h" 34 #include "mlir/Dialect/MemRef/Transforms/Passes.h" 35 #include "mlir/Pass/PassManager.h" 36 #include "mlir/Pass/PassOptions.h" 37 #include "mlir/Transforms/Passes.h" 38 39 using namespace mlir; 40 41 namespace { 42 43 //===----------------------------------------------------------------------===// 44 // Common pipeline 45 //===----------------------------------------------------------------------===// 46 void buildCommonPassPipeline( 47 OpPassManager &pm, const mlir::gpu::GPUToNVVMPipelineOptions &options) { 48 pm.addPass(createConvertNVGPUToNVVMPass()); 49 pm.addPass(createGpuKernelOutliningPass()); 50 pm.addPass(createConvertVectorToSCFPass()); 51 pm.addPass(createConvertSCFToCFPass()); 52 pm.addPass(createConvertNVVMToLLVMPass()); 53 pm.addPass(createConvertFuncToLLVMPass()); 54 pm.addPass(memref::createExpandStridedMetadataPass()); 55 56 GpuNVVMAttachTargetOptions nvvmTargetOptions; 57 nvvmTargetOptions.triple = options.cubinTriple; 58 nvvmTargetOptions.chip = options.cubinChip; 59 nvvmTargetOptions.features = options.cubinFeatures; 60 nvvmTargetOptions.optLevel = options.optLevel; 61 pm.addPass(createGpuNVVMAttachTarget(nvvmTargetOptions)); 62 pm.addPass(createLowerAffinePass()); 63 pm.addPass(createArithToLLVMConversionPass()); 64 ConvertIndexToLLVMPassOptions convertIndexToLLVMPassOpt; 65 convertIndexToLLVMPassOpt.indexBitwidth = options.indexBitWidth; 66 pm.addPass(createConvertIndexToLLVMPass(convertIndexToLLVMPassOpt)); 67 pm.addPass(createCanonicalizerPass()); 68 pm.addPass(createCSEPass()); 69 } 70 71 //===----------------------------------------------------------------------===// 72 // GPUModule-specific stuff. 73 //===----------------------------------------------------------------------===// 74 void buildGpuPassPipeline(OpPassManager &pm, 75 const mlir::gpu::GPUToNVVMPipelineOptions &options) { 76 ConvertGpuOpsToNVVMOpsOptions opt; 77 opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv; 78 opt.indexBitwidth = options.indexBitWidth; 79 pm.addNestedPass<gpu::GPUModuleOp>(createConvertGpuOpsToNVVMOps(opt)); 80 pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass()); 81 pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass()); 82 pm.addNestedPass<gpu::GPUModuleOp>(createReconcileUnrealizedCastsPass()); 83 } 84 85 //===----------------------------------------------------------------------===// 86 // Host Post-GPU pipeline 87 //===----------------------------------------------------------------------===// 88 void buildHostPostPipeline(OpPassManager &pm, 89 const mlir::gpu::GPUToNVVMPipelineOptions &options) { 90 GpuToLLVMConversionPassOptions opt; 91 opt.hostBarePtrCallConv = options.hostUseBarePtrCallConv; 92 opt.kernelBarePtrCallConv = options.kernelUseBarePtrCallConv; 93 pm.addPass(createGpuToLLVMConversionPass(opt)); 94 95 GpuModuleToBinaryPassOptions gpuModuleToBinaryPassOptions; 96 gpuModuleToBinaryPassOptions.compilationTarget = options.cubinFormat; 97 pm.addPass(createGpuModuleToBinaryPass(gpuModuleToBinaryPassOptions)); 98 pm.addPass(createConvertMathToLLVMPass()); 99 pm.addPass(createCanonicalizerPass()); 100 pm.addPass(createCSEPass()); 101 pm.addPass(createReconcileUnrealizedCastsPass()); 102 } 103 104 } // namespace 105 106 void mlir::gpu::buildLowerToNVVMPassPipeline( 107 OpPassManager &pm, const GPUToNVVMPipelineOptions &options) { 108 // Common pipelines 109 buildCommonPassPipeline(pm, options); 110 111 // GPUModule-specific stuff 112 buildGpuPassPipeline(pm, options); 113 114 // Host post-GPUModule-specific stuff 115 buildHostPostPipeline(pm, options); 116 } 117 118 void mlir::gpu::registerGPUToNVVMPipeline() { 119 PassPipelineRegistration<GPUToNVVMPipelineOptions>( 120 "gpu-lower-to-nvvm-pipeline", 121 "The default pipeline lowers main dialects (arith, memref, scf, " 122 "vector, gpu, and nvgpu) to NVVM. It starts by lowering GPU code to the " 123 "specified compilation target (default is fatbin) then lowers the host " 124 "code.", 125 buildLowerToNVVMPassPipeline); 126 } 127