124 lines · cpp
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-exception6//7//===----------------------------------------------------------------------===//8//9// This file implements a pass for testing the lowering to NVVM as a generally10// 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 pipeline40//===----------------------------------------------------------------------===//41void buildCommonPassPipeline(42 OpPassManager &pm, const mlir::gpu::GPUToNVVMPipelineOptions &options) {43 pm.addPass(createConvertNVGPUToNVVMPass());44 pm.addPass(createGpuKernelOutliningPass());45 pm.addPass(createConvertVectorToSCFPass());46 pm.addPass(createSCFToControlFlowPass());47 pm.addPass(createConvertNVVMToLLVMPass());48 pm.addPass(createConvertFuncToLLVMPass());49 pm.addPass(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(createGpuNVVMAttachTarget(nvvmTargetOptions));58 pm.addPass(createLowerAffinePass());59 pm.addPass(createArithToLLVMConversionPass());60 ConvertIndexToLLVMPassOptions convertIndexToLLVMPassOpt;61 convertIndexToLLVMPassOpt.indexBitwidth = options.indexBitWidth;62 pm.addPass(createConvertIndexToLLVMPass(convertIndexToLLVMPassOpt));63 pm.addPass(createCanonicalizerPass());64 pm.addPass(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 opt.allowPatternRollback = options.allowPatternRollback;76 pm.addNestedPass<gpu::GPUModuleOp>(createConvertGpuOpsToNVVMOps(opt));77 pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass());78 pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());79 pm.addNestedPass<gpu::GPUModuleOp>(createReconcileUnrealizedCastsPass());80}81 82//===----------------------------------------------------------------------===//83// Host Post-GPU pipeline84//===----------------------------------------------------------------------===//85void buildHostPostPipeline(OpPassManager &pm,86 const mlir::gpu::GPUToNVVMPipelineOptions &options) {87 GpuToLLVMConversionPassOptions opt;88 opt.hostBarePtrCallConv = options.hostUseBarePtrCallConv;89 opt.kernelBarePtrCallConv = options.kernelUseBarePtrCallConv;90 pm.addPass(createGpuToLLVMConversionPass(opt));91 92 GpuModuleToBinaryPassOptions gpuModuleToBinaryPassOptions;93 gpuModuleToBinaryPassOptions.compilationTarget = options.cubinFormat;94 pm.addPass(createGpuModuleToBinaryPass(gpuModuleToBinaryPassOptions));95 pm.addPass(createConvertMathToLLVMPass());96 pm.addPass(createCanonicalizerPass());97 pm.addPass(createCSEPass());98 pm.addPass(createReconcileUnrealizedCastsPass());99}100 101} // namespace102 103void mlir::gpu::buildLowerToNVVMPassPipeline(104 OpPassManager &pm, const GPUToNVVMPipelineOptions &options) {105 // Common pipelines106 buildCommonPassPipeline(pm, options);107 108 // GPUModule-specific stuff109 buildGpuPassPipeline(pm, options);110 111 // Host post-GPUModule-specific stuff112 buildHostPostPipeline(pm, options);113}114 115void mlir::gpu::registerGPUToNVVMPipeline() {116 PassPipelineRegistration<GPUToNVVMPipelineOptions>(117 "gpu-lower-to-nvvm-pipeline",118 "The default pipeline lowers main dialects (arith, memref, scf, "119 "vector, gpu, and nvgpu) to NVVM. It starts by lowering GPU code to the "120 "specified compilation target (default is fatbin) then lowers the host "121 "code.",122 buildLowerToNVVMPassPipeline);123}124