diff --git a/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h b/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h --- a/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h +++ b/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h @@ -33,7 +33,6 @@ LowerToLLVMOptions(MLIRContext *ctx, const DataLayout &dl); bool useBarePtrCallConv = false; - bool emitCWrappers = false; enum class AllocLowering { /// Use malloc for for heap allocations. diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td --- a/mlir/include/mlir/Conversion/Passes.td +++ b/mlir/include/mlir/Conversion/Passes.td @@ -294,9 +294,6 @@ /*default=*/"false", "Replace FuncOp's MemRef arguments with bare pointers to the MemRef " "element types">, - Option<"emitCWrappers", "emit-c-wrappers", "bool", /*default=*/"false", - "Emit wrappers for C-compatible pointer-to-struct memref " - "descriptors">, Option<"indexBitwidth", "index-bitwidth", "unsigned", /*default=kDeriveIndexBitwidthFromDataLayout*/"0", "Bitwidth of the index type, 0 to use size of machine word">, diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td --- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td +++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td @@ -56,6 +56,11 @@ /// Name of the target triple attribute. static StringRef getTargetTripleAttrName() { return "llvm.target_triple"; } + + /// Name of the C wrapper emission attribute. + static StringRef getEmitCWrapperAttrName() { + return "llvm.emit_c_interface"; + } }]; let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed; diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h @@ -11,6 +11,7 @@ #include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h" #include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h" +#include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h" #include "mlir/Pass/Pass.h" namespace mlir { diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td @@ -13,12 +13,25 @@ def LLVMLegalizeForExport : Pass<"llvm-legalize-for-export"> { let summary = "Legalize LLVM dialect to be convertible to LLVM IR"; - let constructor = "mlir::LLVM::createLegalizeForExportPass()"; + let constructor = "::mlir::LLVM::createLegalizeForExportPass()"; +} + +def LLVMRequestCWrappers + : Pass<"llvm-request-c-wrappers", "::mlir::func::FuncOp"> { + let summary = "Request C wrapper emission for all functions"; + let description = [{ + Annotate every builtin function in the module with the LLVM dialect + attribute that instructs the conversion to LLVM to emit the C wrapper for + the function. This pass is expected to be applied immediately before the + conversion of builtin functions to LLVM to avoid the attribute being + dropped by other passes. + }]; + let constructor = "::mlir::LLVM::createRequestCWrappersPass()"; } def NVVMOptimizeForTarget : Pass<"llvm-optimize-for-nvvm-target"> { let summary = "Optimize NVVM IR"; - let constructor = "mlir::NVVM::createOptimizeForTargetPass()"; + let constructor = "::mlir::NVVM::createOptimizeForTargetPass()"; } #endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_PASSES diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h new file mode 100644 --- /dev/null +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h @@ -0,0 +1,22 @@ +//===- RequestCWrappers.h - Annotate funcs with wrap attributes -*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_LLVMIR_TRANSFORMS_REQUESTCWRAPPERS_H +#define MLIR_DIALECT_LLVMIR_TRANSFORMS_REQUESTCWRAPPERS_H + +#include + +namespace mlir { +class Pass; + +namespace LLVM { +std::unique_ptr createRequestCWrappersPass(); +} // namespace LLVM +} // namespace mlir + +#endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_REQUESTCWRAPPERS_H diff --git a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp --- a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp +++ b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp @@ -361,7 +361,6 @@ /// FuncOp legalization pattern that converts MemRef arguments to pointers to /// MemRef descriptors (LLVM struct data types) containing all the MemRef type /// information. -static constexpr StringRef kEmitIfaceAttrName = "llvm.emit_c_interface"; struct FuncOpConversion : public FuncOpConversionBase { FuncOpConversion(LLVMTypeConverter &converter) : FuncOpConversionBase(converter) {} @@ -373,8 +372,8 @@ if (!newFuncOp) return failure(); - if (getTypeConverter()->getOptions().emitCWrappers || - funcOp->getAttrOfType(kEmitIfaceAttrName)) { + if (funcOp->getAttrOfType( + LLVM::LLVMDialect::getEmitCWrapperAttrName())) { if (newFuncOp.isExternal()) wrapExternalFunction(rewriter, funcOp.getLoc(), *getTypeConverter(), funcOp, newFuncOp); @@ -676,24 +675,16 @@ struct ConvertFuncToLLVMPass : public ConvertFuncToLLVMBase { ConvertFuncToLLVMPass() = default; - ConvertFuncToLLVMPass(bool useBarePtrCallConv, bool emitCWrappers, - unsigned indexBitwidth, bool useAlignedAlloc, + ConvertFuncToLLVMPass(bool useBarePtrCallConv, unsigned indexBitwidth, + bool useAlignedAlloc, const llvm::DataLayout &dataLayout) { this->useBarePtrCallConv = useBarePtrCallConv; - this->emitCWrappers = emitCWrappers; this->indexBitwidth = indexBitwidth; this->dataLayout = dataLayout.getStringRepresentation(); } /// Run the dialect converter on the module. void runOnOperation() override { - if (useBarePtrCallConv && emitCWrappers) { - getOperation().emitError() - << "incompatible conversion options: bare-pointer calling convention " - "and C wrapper emission"; - signalPassFailure(); - return; - } if (failed(LLVM::LLVMDialect::verifyDataLayoutString( this->dataLayout, [this](const Twine &message) { getOperation().emitError() << message.str(); @@ -708,7 +699,6 @@ LowerToLLVMOptions options(&getContext(), dataLayoutAnalysis.getAtOrAbove(m)); options.useBarePtrCallConv = useBarePtrCallConv; - options.emitCWrappers = emitCWrappers; if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout) options.overrideIndexBitwidth(indexBitwidth); options.dataLayout = llvm::DataLayout(this->dataLayout); @@ -747,6 +737,6 @@ bool useAlignedAlloc = (allocLowering == LowerToLLVMOptions::AllocLowering::AlignedAlloc); return std::make_unique( - options.useBarePtrCallConv, options.emitCWrappers, - options.getIndexBitwidth(), useAlignedAlloc, options.dataLayout); + options.useBarePtrCallConv, options.getIndexBitwidth(), useAlignedAlloc, + options.dataLayout); } diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp --- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp +++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp @@ -166,18 +166,23 @@ void runOnOperation() override { gpu::GPUModuleOp m = getOperation(); - /// Customize the bitwidth used for the device side index computations. + // Request C wrapper emission. + for (auto func : m.getOps()) { + func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), + UnitAttr::get(&getContext())); + } + + // Customize the bitwidth used for the device side index computations. LowerToLLVMOptions options( m.getContext(), DataLayout(cast(m.getOperation()))); - options.emitCWrappers = true; if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout) options.overrideIndexBitwidth(indexBitwidth); - /// MemRef conversion for GPU to NVVM lowering. The GPU dialect uses memory - /// space 5 for private memory attributions, but NVVM represents private - /// memory allocations as local `alloca`s in the default address space. This - /// converter drops the private memory space to support the use case above. + // MemRef conversion for GPU to NVVM lowering. The GPU dialect uses memory + // space 5 for private memory attributions, but NVVM represents private + // memory allocations as local `alloca`s in the default address space. This + // converter drops the private memory space to support the use case above. LLVMTypeConverter converter(m.getContext(), options); converter.addConversion([&](MemRefType type) -> Optional { if (type.getMemorySpaceAsInt() != diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp --- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp +++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp @@ -63,11 +63,16 @@ void runOnOperation() override { gpu::GPUModuleOp m = getOperation(); + // Request C wrapper emission. + for (auto func : m.getOps()) { + func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), + UnitAttr::get(&getContext())); + } + /// Customize the bitwidth used for the device side index computations. LowerToLLVMOptions options( m.getContext(), DataLayout(cast(m.getOperation()))); - options.emitCWrappers = true; if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout) options.overrideIndexBitwidth(indexBitwidth); LLVMTypeConverter converter(m.getContext(), options); diff --git a/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp b/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp --- a/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp +++ b/mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp @@ -11,6 +11,7 @@ #include "../PassDetail.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/Transforms/Transforms.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -71,7 +72,8 @@ // Insert a function attribute that will trigger the emission of the // corresponding `_mlir_ciface_xxx` interface so that external libraries see // a normalized ABI. This interface is added during std to llvm conversion. - funcOp->setAttr("llvm.emit_c_interface", UnitAttr::get(op->getContext())); + funcOp->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), + UnitAttr::get(op->getContext())); funcOp.setPrivate(); return fnNameAttr; } diff --git a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp --- a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp +++ b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp @@ -20,6 +20,7 @@ #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" #include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h" #include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" @@ -280,9 +281,14 @@ llvm::make_early_inc_range(module.getOps())) gpuModule.erase(); + // Request C wrapper emission. + for (auto func : module.getOps()) { + func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), + UnitAttr::get(&getContext())); + } + // Specify options to lower to LLVM and pull in the conversion patterns. LowerToLLVMOptions options(module.getContext()); - options.emitCWrappers = true; auto *context = module.getContext(); RewritePatternSet patterns(context); LLVMTypeConverter typeConverter(context, options); diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt --- a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt +++ b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt @@ -1,12 +1,14 @@ add_mlir_dialect_library(MLIRLLVMIRTransforms LegalizeForExport.cpp OptimizeForNVVM.cpp + RequestCWrappers.cpp DEPENDS MLIRLLVMPassIncGen LINK_LIBS PUBLIC MLIRIR + MLIRFuncDialect MLIRLLVMDialect MLIRPass MLIRTransforms diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h b/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h --- a/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h +++ b/mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h @@ -12,6 +12,9 @@ #include "mlir/Pass/Pass.h" namespace mlir { +namespace func { +class FuncOp; +} // namespace func #define GEN_PASS_CLASSES #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp new file mode 100644 --- /dev/null +++ b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp @@ -0,0 +1,29 @@ +//===- RequestCWrappers.cpp - Annotate funcs with wrap attributes ---------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h" +#include "PassDetail.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" + +using namespace mlir; + +namespace { +class RequestCWrappersPass + : public LLVMRequestCWrappersBase { +public: + void runOnOperation() override { + getOperation()->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), + UnitAttr::get(&getContext())); + } +}; +} // namespace + +std::unique_ptr mlir::LLVM::createRequestCWrappersPass() { + return std::make_unique(); +} diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp --- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp @@ -65,7 +65,8 @@ FunctionType::get(context, operands.getTypes(), resultType)); func.setPrivate(); if (static_cast(emitCInterface)) - func->setAttr("llvm.emit_c_interface", UnitAttr::get(context)); + func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), + UnitAttr::get(context)); } return result; } diff --git a/mlir/test/Conversion/FuncToLLVM/calling-convention.mlir b/mlir/test/Conversion/FuncToLLVM/calling-convention.mlir --- a/mlir/test/Conversion/FuncToLLVM/calling-convention.mlir +++ b/mlir/test/Conversion/FuncToLLVM/calling-convention.mlir @@ -1,4 +1,4 @@ -// RUN: mlir-opt -convert-memref-to-llvm -convert-func-to-llvm='emit-c-wrappers=1' -reconcile-unrealized-casts %s | FileCheck %s +// RUN: mlir-opt -convert-memref-to-llvm -llvm-request-c-wrappers -convert-func-to-llvm -reconcile-unrealized-casts %s | FileCheck %s // RUN: mlir-opt -convert-memref-to-llvm -convert-func-to-llvm -reconcile-unrealized-casts %s | FileCheck %s --check-prefix=EMIT_C_ATTRIBUTE // This tests the default memref calling convention and the emission of C diff --git a/mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-callers.mlir b/mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-callers.mlir rename from mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-callers.mlir rename to mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-callers.mlir --- a/mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-callers.mlir +++ b/mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-callers.mlir @@ -1,4 +1,4 @@ -// RUN: mlir-opt -convert-func-to-llvm='emit-c-wrappers=1' %s | FileCheck %s +// RUN: mlir-opt -llvm-request-c-wrappers -convert-func-to-llvm %s | FileCheck %s // CHECK: llvm.func @res_attrs_with_memref_return() -> (!llvm.struct{{.*}} {test.returnOne}) // CHECK-LABEL: llvm.func @_mlir_ciface_res_attrs_with_memref_return diff --git a/mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-functions.mlir b/mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-functions.mlir rename from mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-functions.mlir rename to mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-functions.mlir --- a/mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-functions.mlir +++ b/mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-functions.mlir @@ -1,4 +1,4 @@ -// RUN: mlir-opt -convert-func-to-llvm='emit-c-wrappers=1' %s | FileCheck %s +// RUN: mlir-opt -llvm-request-c-wrappers -convert-func-to-llvm %s | FileCheck %s // CHECK: llvm.func @res_attrs_with_memref_return() -> (!llvm.struct{{.*}} {test.returnOne}) // CHECK-LABEL: llvm.func @_mlir_ciface_res_attrs_with_memref_return diff --git a/mlir/test/Conversion/GPUToNVVM/wmma-ops-to-nvvm.mlir b/mlir/test/Conversion/GPUToNVVM/wmma-ops-to-nvvm.mlir --- a/mlir/test/Conversion/GPUToNVVM/wmma-ops-to-nvvm.mlir +++ b/mlir/test/Conversion/GPUToNVVM/wmma-ops-to-nvvm.mlir @@ -4,7 +4,7 @@ gpu.module @test_module { // CHECK-LABEL: func @gpu_wmma_load_op() -> - // CHECK-SAME: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> { + // CHECK-SAME: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> // CHECK32-LABEL: func @gpu_wmma_load_op() -> func.func @gpu_wmma_load_op() -> (!gpu.mma_matrix<16x16xf16, "AOp">) { %wg = memref.alloca() {alignment = 32} : memref<32x32xf16, 3> @@ -43,9 +43,9 @@ gpu.module @test_module { // CHECK-LABEL: func @gpu_wmma_store_op - // CHECK-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>) { + // CHECK-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>) // CHECK32-LABEL: func @gpu_wmma_store_op - // CHECK32-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>) { + // CHECK32-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>) func.func @gpu_wmma_store_op(%arg0 : !gpu.mma_matrix<16x16xf16, "COp">) -> () { %sg = memref.alloca(){alignment = 32} : memref<32x32xf16, 3> %i = arith.constant 16 : index diff --git a/mlir/test/Dialect/LLVMIR/request-c-wrappers.mlir b/mlir/test/Dialect/LLVMIR/request-c-wrappers.mlir new file mode 100644 --- /dev/null +++ b/mlir/test/Dialect/LLVMIR/request-c-wrappers.mlir @@ -0,0 +1,9 @@ +// RUN: mlir-opt %s -llvm-request-c-wrappers | FileCheck %s + +// CHECK: func.func private @foo() attributes {llvm.emit_c_interface} +func.func private @foo() + +// CHECK: func.func @bar() attributes {llvm.emit_c_interface} +func.func @bar() { + return +} diff --git a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp --- a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp +++ b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp @@ -24,6 +24,7 @@ #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/GPU/Transforms/Passes.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" @@ -52,8 +53,8 @@ modulePM.addPass(spirv::createUpdateVersionCapabilityExtensionPass()); passManager.addPass(createConvertGpuLaunchFuncToVulkanLaunchFuncPass()); LowerToLLVMOptions llvmOptions(module.getContext(), DataLayout(module)); - llvmOptions.emitCWrappers = true; passManager.addPass(createMemRefToLLVMPass()); + passManager.nest().addPass(LLVM::createRequestCWrappersPass()); passManager.addPass(createConvertFuncToLLVMPass(llvmOptions)); passManager.addPass(createReconcileUnrealizedCastsPass()); passManager.addPass(createConvertVulkanLaunchFuncToVulkanCallsPass()); diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel --- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel +++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel @@ -3397,6 +3397,7 @@ hdrs = glob(["include/mlir/Dialect/LLVMIR/Transforms/*.h"]), includes = ["include"], deps = [ + ":FuncDialect", ":IR", ":LLVMDialect", ":LLVMPassIncGen", @@ -6484,6 +6485,7 @@ ":GPUTransforms", ":LLVMCommonConversion", ":LLVMDialect", + ":LLVMIRTransforms", ":LLVMToLLVMIRTranslation", ":MemRefDialect", ":MemRefToLLVM", @@ -7273,6 +7275,7 @@ ":ConversionPassIncGen", ":FuncDialect", ":IR", + ":LLVMDialect", ":LinalgDialect", ":LinalgTransforms", ":MemRefDialect",