From 610139d2d9ce6746b3c617fb3e2f7886272d26ff Mon Sep 17 00:00:00 2001 From: Alex Zinenko Date: Thu, 16 Jun 2022 12:53:49 +0200 Subject: [PATCH] [mlir] replace 'emit_c_wrappers' func->llvm conversion option with a pass The 'emit_c_wrappers' option in the FuncToLLVM conversion requests C interface wrappers to be emitted for every builtin function in the module. While this has been useful to bootstrap the interface, it is problematic in the longer term as it may unintentionally affect the functions that should retain their existing interface, e.g., libm functions obtained by lowering math operations (see D126964 for an example). Since D77314, we have a finer-grain control over interface generation via an attribute that avoids the problem entirely. Remove the 'emit_c_wrappers' option. Introduce the '-llvm-request-c-wrappers' pass that can be run in any pipeline that needs blanket emission of functions to annotate all builtin functions with the attribute before performing the usual lowering that accounts for the attribute. Reviewed By: chelini Differential Revision: https://reviews.llvm.org/D127952 --- .../Conversion/LLVMCommon/LoweringOptions.h | 1 - mlir/include/mlir/Conversion/Passes.td | 3 -- .../include/mlir/Dialect/LLVMIR/LLVMOpBase.td | 5 ++++ .../mlir/Dialect/LLVMIR/Transforms/Passes.h | 1 + .../mlir/Dialect/LLVMIR/Transforms/Passes.td | 17 +++++++++-- .../LLVMIR/Transforms/RequestCWrappers.h | 22 ++++++++++++++ mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp | 22 ++++---------- .../GPUToNVVM/LowerGpuOpsToNVVMOps.cpp | 17 +++++++---- .../GPUToROCDL/LowerGpuOpsToROCDLOps.cpp | 7 ++++- .../LinalgToStandard/LinalgToStandard.cpp | 4 ++- .../ConvertLaunchFuncToLLVMCalls.cpp | 8 ++++- .../Dialect/LLVMIR/Transforms/CMakeLists.txt | 2 ++ .../Dialect/LLVMIR/Transforms/PassDetail.h | 3 ++ .../LLVMIR/Transforms/RequestCWrappers.cpp | 29 +++++++++++++++++++ .../Transforms/SparseTensorConversion.cpp | 3 +- .../FuncToLLVM/calling-convention.mlir | 2 +- .../emit-c-wrappers-for-external-callers.mlir | 2 +- ...mit-c-wrappers-for-external-functions.mlir | 2 +- .../GPUToNVVM/wmma-ops-to-nvvm.mlir | 6 ++-- .../Dialect/LLVMIR/request-c-wrappers.mlir | 9 ++++++ .../mlir-vulkan-runner/mlir-vulkan-runner.cpp | 3 +- .../llvm-project-overlay/mlir/BUILD.bazel | 3 ++ 22 files changed, 132 insertions(+), 39 deletions(-) create mode 100644 mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h create mode 100644 mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp rename mlir/test/Conversion/{StandardToLLVM => FuncToLLVM}/emit-c-wrappers-for-external-callers.mlir (98%) rename mlir/test/Conversion/{StandardToLLVM => FuncToLLVM}/emit-c-wrappers-for-external-functions.mlir (97%) create mode 100644 mlir/test/Dialect/LLVMIR/request-c-wrappers.mlir diff --git a/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h b/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h index 9214b333cfde09f..91f6f3d8addf5de 100644 --- a/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h +++ b/mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h @@ -33,7 +33,6 @@ class LowerToLLVMOptions { 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 index 6d9863e7234879b..ebe9dd7645f7b95 100644 --- a/mlir/include/mlir/Conversion/Passes.td +++ b/mlir/include/mlir/Conversion/Passes.td @@ -294,9 +294,6 @@ def ConvertFuncToLLVM : Pass<"convert-func-to-llvm", "ModuleOp"> { /*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 index f984493da7824f0..c747cb13f2c2bac 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td +++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td @@ -56,6 +56,11 @@ def LLVM_Dialect : Dialect { /// 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 index 39948557b55a6fd..b475488c7ee9799 100644 --- 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 index 060822603bc20f9..a22efd392d5e6a2 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td @@ -13,12 +13,25 @@ include "mlir/Pass/PassBase.td" 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 index 000000000000000..166ae6a35cd8e36 --- /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 index f162f779922f8a1..58b81c1f883863e 100644 --- a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp +++ b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp @@ -361,7 +361,6 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern { /// 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 @@ struct FuncOpConversion : public FuncOpConversionBase { 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 @@ namespace { 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 @@ struct ConvertFuncToLLVMPass 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 @@ mlir::createConvertFuncToLLVMPass(const LowerToLLVMOptions &options) { 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 index e6ae7c8126c0634..b7158f22e7e4d54 100644 --- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp +++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp @@ -166,18 +166,23 @@ struct LowerGpuOpsToNVVMOpsPass 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 index 6b48b4ee4230b0e..a149dda2be39b64 100644 --- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp +++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp @@ -63,11 +63,16 @@ struct LowerGpuOpsToROCDLOpsPass 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 index 0fa51c7938b47ec..b257d2a2a8642cc 100644 --- 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 @@ static FlatSymbolRefAttr getLibraryCallSymbolRef(Operation *op, // 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 index fa6930fe46c868a..663fa8958baba9b 100644 --- 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 @@ class LowerHostCodeToLLVM 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 index c68f14fa71dde7a..b1406be17bd5b50 100644 --- 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 index 8449f77d4457875..deb54ba082a5a43 100644 --- 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 index 000000000000000..d32f51053e8750b --- /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 index 0e7a96e03f08c0d..3ca7ff1c62cabb5 100644 --- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp @@ -65,7 +65,8 @@ static FlatSymbolRefAttr getFunc(Operation *op, StringRef name, 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 index 7e32ef5a45d2a10..8dc5dac7f2d5f52 100644 --- 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 similarity index 98% 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 index dd9fc392365620e..956924595dc313b 100644 --- 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 similarity index 97% 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 index 1e97d9548da44fa..1f6a0f5227a550e 100644 --- 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 index 1ca73dbd0f7c4a8..46be05f08e612fe 100644 --- 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 { 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 index 000000000000000..efc0a797220b7ba --- /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 index 2c672794e3d2887..7e61ea746b968cc 100644 --- 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 @@ static LogicalResult runMLIRPasses(ModuleOp module) { 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 index 5dc2b902704af49..89ca55ecacabdca 100644 --- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel +++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel @@ -3397,6 +3397,7 @@ cc_library( hdrs = glob(["include/mlir/Dialect/LLVMIR/Transforms/*.h"]), includes = ["include"], deps = [ + ":FuncDialect", ":IR", ":LLVMDialect", ":LLVMPassIncGen", @@ -6484,6 +6485,7 @@ cc_binary( ":GPUTransforms", ":LLVMCommonConversion", ":LLVMDialect", + ":LLVMIRTransforms", ":LLVMToLLVMIRTranslation", ":MemRefDialect", ":MemRefToLLVM", @@ -7273,6 +7275,7 @@ cc_library( ":ConversionPassIncGen", ":FuncDialect", ":IR", + ":LLVMDialect", ":LinalgDialect", ":LinalgTransforms", ":MemRefDialect",