[mlir] replace 'emit_c_wrappers' func->llvm conversion option with a pass

[ftynse]

Summary

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.

[zero9178]

Very big fan of this change. That option in particular is very specific to a single scenario (memrefs as arguments), and being able to move it out of a generic options class is a lot cleaner.

The layering here is a bit odd to me however. The way llvm.emit_c_interface is currently implemented, is only applicable to func.func to llvm.func conversions (which is fair), yet the convenience pass to add the llvm.emit_c_wrapper attribute to all functions lives in LLVM Dialect transforms and depends on func.func.
The way I see it either:

• llvm.emit_c_interface is meant as a general purpose indicator that any dialect may use during type lowering. In that case the pass should probably be a generic operation pass and not only apply to mlir::func::FuncOp. This would also drop the link dependency on MLIRFuncDialect from LLVMIRTransforms
• llvm.emit_c_interface shouldn't be part of the LLVM dialect but rather of the func dialect. There is currently no code that actually checks llvm.emit_c_interface except in FuncToLLVM.cpp

[chelini]

In D127952#3588590, @zero9178 wrote:

Very big fan of this change. That option in particular is very specific to a single scenario (memrefs as arguments), and being able to move it out of a generic options class is a lot cleaner.

The layering here is a bit odd to me however. The way llvm.emit_c_interface is currently implemented, is only applicable to func.func to llvm.func conversions (which is fair), yet the convenience pass to add the llvm.emit_c_wrapper attribute to all functions lives in LLVM Dialect transforms and depends on func.func.
The way I see it either:

• llvm.emit_c_interface is meant as a general purpose indicator that any dialect may use during type lowering. In that case the pass should probably be a generic operation pass and not only apply to mlir::func::FuncOp. This would also drop the link dependency on MLIRFuncDialect from LLVMIRTransforms
• llvm.emit_c_interface shouldn't be part of the LLVM dialect but rather of the func dialect. There is currently no code that actually checks llvm.emit_c_interface except in FuncToLLVM.cpp

My two cents on this: I'm not too fond of options one as I don't see many other use cases for the attribute. Option two may have the problem that any pass before conversion to LLVM can or cannot discard the attribute. I believe that's why we want to do that right before conversion to LLVM.

[zero9178]

In D127952#3588640, @chelini wrote:

In D127952#3588590, @zero9178 wrote:

Very big fan of this change. That option in particular is very specific to a single scenario (memrefs as arguments), and being able to move it out of a generic options class is a lot cleaner.

The layering here is a bit odd to me however. The way llvm.emit_c_interface is currently implemented, is only applicable to func.func to llvm.func conversions (which is fair), yet the convenience pass to add the llvm.emit_c_wrapper attribute to all functions lives in LLVM Dialect transforms and depends on func.func.
The way I see it either:

• llvm.emit_c_interface is meant as a general purpose indicator that any dialect may use during type lowering. In that case the pass should probably be a generic operation pass and not only apply to mlir::func::FuncOp. This would also drop the link dependency on MLIRFuncDialect from LLVMIRTransforms
• llvm.emit_c_interface shouldn't be part of the LLVM dialect but rather of the func dialect. There is currently no code that actually checks llvm.emit_c_interface except in FuncToLLVM.cpp

My two cents on this: I'm not too fond of options one as I don't see many other use cases for the attribute. Option two may have the problem that any pass before conversion to LLVM can or cannot discard the attribute. I believe that's why we want to do that right before conversion to LLVM.

Regarding option one: Other dialects with a custom function op might want to have the same behaviour when lowering to LLVM in the future. This option is also a tiny change from the current state of the patch.

Regarding option two: Passes can discard the attribute regardless of whether it is part of the LLVM dialect of the func dialect. The second option would not change anything in that regard except that we'd probably want the pass to live in the Func Dialect instead of the LLVM Dialect. It being run before the FuncToLLVM pass would remain a requirement.

[chelini]

In D127952#3588685, @zero9178 wrote:

In D127952#3588640, @chelini wrote:

In D127952#3588590, @zero9178 wrote:

Very big fan of this change. That option in particular is very specific to a single scenario (memrefs as arguments), and being able to move it out of a generic options class is a lot cleaner.

The layering here is a bit odd to me however. The way llvm.emit_c_interface is currently implemented, is only applicable to func.func to llvm.func conversions (which is fair), yet the convenience pass to add the llvm.emit_c_wrapper attribute to all functions lives in LLVM Dialect transforms and depends on func.func.
The way I see it either:

• llvm.emit_c_interface is meant as a general purpose indicator that any dialect may use during type lowering. In that case the pass should probably be a generic operation pass and not only apply to mlir::func::FuncOp. This would also drop the link dependency on MLIRFuncDialect from LLVMIRTransforms
• llvm.emit_c_interface shouldn't be part of the LLVM dialect but rather of the func dialect. There is currently no code that actually checks llvm.emit_c_interface except in FuncToLLVM.cpp

My two cents on this: I'm not too fond of options one as I don't see many other use cases for the attribute. Option two may have the problem that any pass before conversion to LLVM can or cannot discard the attribute. I believe that's why we want to do that right before conversion to LLVM.

Regarding option one: Other dialects with a custom function op might want to have the same behaviour when lowering to LLVM in the future. This option is also a tiny change from the current state of the patch.

Right my bad; I didn't understand from your previous comment how you want to scope the attribute if we move to a generic operation pass. This attribute is fairly specific and should apply only for function-like operations. So if this is what you have in mind, I agree. If we move to an operation pass, I suggest keeping the attribute in the LLVM dialect.

Regarding option two: Passes can discard the attribute regardless of whether it is part of the LLVM dialect of the func dialect. The second option would not change anything in that regard except that we'd probably want the pass to live in the Func Dialect instead of the LLVM Dialect. It being run before the FuncToLLVM pass would remain a requirement.

[ftynse]

I went for the least worst solution here. FuncTransforms library already depends on "higher level" dialects (SCF, MemRef and co.), adding the "lower level" dialect sounds undesirable. LLVMTransforms, on the other hand, can have a dependency on the "higher level" Func dialect and be consistent. (Note that there is no dependency between dialects themselves.)

The attribute is only relevant for the func->llvm conversion, so the independent place for it would be somewhere in lib/Conversions/FuncToLLVM/AdditionalPasses, but that would be a significant architectural change, even though with a small amount of code. Making it into an operation pass will allow the attribute to be attached to any operation, where it doesn't make sense and will be confusing. As for some hypothetical other dialect that may want this, I would rather not overgeneralize prematurely. When and if this dialect appears, we can reconsider. FWIW, the only reason the pass is introduced is some backwards compatibility for downstreams that may depend on this behavior to be configurable from the textual pass pipeline. Otherwise, the code to add the attribute to all functions takes three lines, it can be duplicated at will.

[chelini]

It makes sense to me.

[chelini]

mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td

24

requires -> instructs

[zero9178]

In D127952#3588973, @ftynse wrote:

I went for the least worst solution here. FuncTransforms library already depends on "higher level" dialects (SCF, MemRef and co.), adding the "lower level" dialect sounds undesirable. LLVMTransforms, on the other hand, can have a dependency on the "higher level" Func dialect and be consistent. (Note that there is no dependency between dialects themselves.)

The attribute is only relevant for the func->llvm conversion, so the independent place for it would be somewhere in lib/Conversions/FuncToLLVM/AdditionalPasses, but that would be a significant architectural change, even though with a small amount of code. Making it into an operation pass will allow the attribute to be attached to any operation, where it doesn't make sense and will be confusing. As for some hypothetical other dialect that may want this, I would rather not overgeneralize prematurely. When and if this dialect appears, we can reconsider. FWIW, the only reason the pass is introduced is some backwards compatibility for downstreams that may depend on this behavior to be configurable from the textual pass pipeline. Otherwise, the code to add the attribute to all functions takes three lines, it can be duplicated at will.

I see thank you

mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h

 /// Options to control the LLVM lowering. The struct is used to share lowering
 /// options between passes, patterns, and type converter.
 class LowerToLLVMOptions {
 public:
   explicit LowerToLLVMOptions(MLIRContext *ctx);
   LowerToLLVMOptions(MLIRContext *ctx, const DataLayout &dl);
 
   bool useBarePtrCallConv = false;
-  bool emitCWrappers = false;
 
   enum class AllocLowering {
     /// Use malloc for for heap allocations.
     Malloc,
 
     /// Use aligned_alloc for heap allocations.
     AlignedAlloc,
 

mlir/include/mlir/Conversion/Passes.td

def ConvertFuncToLLVM : Pass<"convert-func-to-llvm", "ModuleOp"> {
   }];
   let constructor = "mlir::createConvertFuncToLLVMPass()";
   let dependentDialects = ["LLVM::LLVMDialect"];
   let options = [
     Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool",
            /*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">,
     Option<"dataLayout", "data-layout", "std::string",
            /*default=*/"\"\"",
            "String description (LLVM format) of the data layout that is "
            "expected on the produced module">
   ];

mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td

let extraClassDeclaration = [{
 
     /// Verifies if the given string is a well-formed data layout descriptor.
     /// Uses `reportError` to report errors.
     static LogicalResult verifyDataLayoutString(
         StringRef descr, llvm::function_ref<void (const Twine &)> reportError);
 
     /// 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;
 }
 
 //===----------------------------------------------------------------------===//
 // LLVM dialect type constraints.
 //===----------------------------------------------------------------------===//

mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h

 //===- Passes.h - LLVM Pass Construction and Registration -----------------===//
 //
 // 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_PASSES_H
 #define MLIR_DIALECT_LLVMIR_TRANSFORMS_PASSES_H
 
 #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 {
 
 namespace LLVM {
 
 /// Generate the code for registering conversion passes.
 #define GEN_PASS_REGISTRATION
 #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
 
 } // namespace LLVM
 } // namespace mlir
 
 #endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_PASSES_H

mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td

 //===-- Passes.td - LLVM pass definition file --------------*- tablegen -*-===//
 //
 // 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_PASSES
 #define MLIR_DIALECT_LLVMIR_TRANSFORMS_PASSES
 
 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

mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h

+//===- 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 <memory>
+
+namespace mlir {
+class Pass;
+
+namespace LLVM {
+std::unique_ptr<Pass> createRequestCWrappersPass();
+} // namespace LLVM
+} // namespace mlir
+
+#endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_REQUESTCWRAPPERS_H

mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp

convertFuncOpToLLVMFuncOp(func::FuncOp funcOp,
 
     return newFuncOp;
   }
 };
 
 /// 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) {}
 
   LogicalResult
   matchAndRewrite(func::FuncOp funcOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto newFuncOp = convertFuncOpToLLVMFuncOp(funcOp, rewriter);
     if (!newFuncOp)
       return failure();
 
-    if (getTypeConverter()->getOptions().emitCWrappers ||
-        funcOp->getAttrOfType<UnitAttr>(kEmitIfaceAttrName)) {
+    if (funcOp->getAttrOfType<UnitAttr>(
+            LLVM::LLVMDialect::getEmitCWrapperAttrName())) {
       if (newFuncOp.isExternal())
         wrapExternalFunction(rewriter, funcOp.getLoc(), *getTypeConverter(),
                              funcOp, newFuncOp);
       else
         wrapForExternalCallers(rewriter, funcOp.getLoc(), *getTypeConverter(),
                                funcOp, newFuncOp);
     }
 
void mlir::populateFuncToLLVMConversionPatterns(LLVMTypeConverter &converter,
   // clang-format on
 }
 
 namespace {
 /// A pass converting Func operations into the LLVM IR dialect.
 struct ConvertFuncToLLVMPass
     : public ConvertFuncToLLVMBase<ConvertFuncToLLVMPass> {
   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();
             }))) {
       signalPassFailure();
       return;
     }
 
     ModuleOp m = getOperation();
     const auto &dataLayoutAnalysis = getAnalysis<DataLayoutAnalysis>();
 
     LowerToLLVMOptions options(&getContext(),
                                dataLayoutAnalysis.getAtOrAbove(m));
     options.useBarePtrCallConv = useBarePtrCallConv;
-    options.emitCWrappers = emitCWrappers;
     if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout)
       options.overrideIndexBitwidth(indexBitwidth);
     options.dataLayout = llvm::DataLayout(this->dataLayout);
 
     LLVMTypeConverter typeConverter(&getContext(), options,
                                     &dataLayoutAnalysis);
 
     RewritePatternSet patterns(&getContext());
mlir::createConvertFuncToLLVMPass(const LowerToLLVMOptions &options) {
   auto allocLowering = options.allocLowering;
   // There is no way to provide additional patterns for pass, so
   // AllocLowering::None will always fail.
   assert(allocLowering != LowerToLLVMOptions::AllocLowering::None &&
          "ConvertFuncToLLVMPass doesn't support AllocLowering::None");
   bool useAlignedAlloc =
       (allocLowering == LowerToLLVMOptions::AllocLowering::AlignedAlloc);
   return std::make_unique<ConvertFuncToLLVMPass>(
-      options.useBarePtrCallConv, options.emitCWrappers,
-      options.getIndexBitwidth(), useAlignedAlloc, options.dataLayout);
+      options.useBarePtrCallConv, options.getIndexBitwidth(), useAlignedAlloc,
+      options.dataLayout);
 }

mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

struct LowerGpuOpsToNVVMOpsPass
   LowerGpuOpsToNVVMOpsPass() = default;
   LowerGpuOpsToNVVMOpsPass(unsigned indexBitwidth) {
     this->indexBitwidth = indexBitwidth;
   }
 
   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::FuncOp>()) {
+      func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+                    UnitAttr::get(&getContext()));
+    }
+
+    // Customize the bitwidth used for the device side index computations.
     LowerToLLVMOptions options(
         m.getContext(),
         DataLayout(cast<DataLayoutOpInterface>(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<Type> {
       if (type.getMemorySpaceAsInt() !=
           gpu::GPUDialect::getPrivateAddressSpace())
         return llvm::None;
       return converter.convertType(MemRefType::Builder(type).setMemorySpace(0));
     });
     // Lowering for MMAMatrixType.

mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp

struct LowerGpuOpsToROCDLOpsPass
   LowerGpuOpsToROCDLOpsPass(unsigned indexBitwidth, gpu::amd::Runtime runtime) {
     this->indexBitwidth = indexBitwidth;
     this->runtime = runtime;
   }
 
   void runOnOperation() override {
     gpu::GPUModuleOp m = getOperation();
 
+    // Request C wrapper emission.
+    for (auto func : m.getOps<func::FuncOp>()) {
+      func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+                    UnitAttr::get(&getContext()));
+    }
+
     /// Customize the bitwidth used for the device side index computations.
     LowerToLLVMOptions options(
         m.getContext(),
         DataLayout(cast<DataLayoutOpInterface>(m.getOperation())));
-    options.emitCWrappers = true;
     if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout)
       options.overrideIndexBitwidth(indexBitwidth);
     LLVMTypeConverter converter(m.getContext(), options);
 
     RewritePatternSet patterns(m.getContext());
     RewritePatternSet llvmPatterns(m.getContext());
 
     populateGpuRewritePatterns(patterns);

mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp

 //===- LinalgToStandard.cpp - conversion from Linalg to Standard dialect --===//
 //
 // 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/Conversion/LinalgToStandard/LinalgToStandard.h"
 
 #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"
 #include "mlir/Dialect/SCF/SCF.h"
 
 using namespace mlir;
 using namespace mlir::linalg;
 
static FlatSymbolRefAttr getLibraryCallSymbolRef(Operation *op,
   // Insert before module terminator.
   rewriter.setInsertionPoint(module.getBody(),
                              std::prev(module.getBody()->end()));
   func::FuncOp funcOp = rewriter.create<func::FuncOp>(
       op->getLoc(), fnNameAttr.getValue(), libFnType);
   // 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;
 }
 
 static SmallVector<Value, 4>
 createTypeCanonicalizedMemRefOperands(OpBuilder &b, Location loc,
                                       ValueRange operands) {
   SmallVector<Value, 4> res;

mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp

 #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h"
 #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h"
 #include "mlir/Conversion/LLVMCommon/LoweringOptions.h"
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
 #include "mlir/Conversion/LLVMCommon/TypeConverter.h"
 #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"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Transforms/DialectConversion.h"
 
 #include "llvm/ADT/DenseMap.h"
public:
   void runOnOperation() override {
     ModuleOp module = getOperation();
 
     // Erase the GPU module.
     for (auto gpuModule :
          llvm::make_early_inc_range(module.getOps<gpu::GPUModuleOp>()))
       gpuModule.erase();
 
+    // Request C wrapper emission.
+    for (auto func : module.getOps<func::FuncOp>()) {
+      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);
     mlir::arith::populateArithmeticToLLVMConversionPatterns(typeConverter,
                                                             patterns);
     populateMemRefToLLVMConversionPatterns(typeConverter, patterns);
     populateFuncToLLVMConversionPatterns(typeConverter, patterns);
     patterns.add<GPULaunchLowering>(typeConverter);

mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt

 add_mlir_dialect_library(MLIRLLVMIRTransforms
   LegalizeForExport.cpp
   OptimizeForNVVM.cpp
+  RequestCWrappers.cpp
 
   DEPENDS
   MLIRLLVMPassIncGen
 
   LINK_LIBS PUBLIC
   MLIRIR
+  MLIRFuncDialect
   MLIRLLVMDialect
   MLIRPass
   MLIRTransforms
   MLIRNVVMDialect
   )

mlir/lib/Dialect/LLVMIR/Transforms/PassDetail.h

 //===- PassDetail.h - LLVM Pass class details -------------------*- 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 DIALECT_LLVMIR_TRANSFORMS_PASSDETAIL_H_
 #define DIALECT_LLVMIR_TRANSFORMS_PASSDETAIL_H_
 
 #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"
 
 } // namespace mlir
 
 #endif // DIALECT_LLVMIR_TRANSFORMS_PASSDETAIL_H_

mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp

+//===- 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<RequestCWrappersPass> {
+public:
+  void runOnOperation() override {
+    getOperation()->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+                            UnitAttr::get(&getContext()));
+  }
+};
+} // namespace
+
+std::unique_ptr<Pass> mlir::LLVM::createRequestCWrappersPass() {
+  return std::make_unique<RequestCWrappersPass>();
+}

mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp

static FlatSymbolRefAttr getFunc(Operation *op, StringRef name,
   auto func = module.lookupSymbol<func::FuncOp>(result.getAttr());
   if (!func) {
     OpBuilder moduleBuilder(module.getBodyRegion());
     func = moduleBuilder.create<func::FuncOp>(
         op->getLoc(), name,
         FunctionType::get(context, operands.getTypes(), resultType));
     func.setPrivate();
     if (static_cast<bool>(emitCInterface))
-      func->setAttr("llvm.emit_c_interface", UnitAttr::get(context));
+      func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+                    UnitAttr::get(context));
   }
   return result;
 }
 
 /// Creates a `CallOp` to the function reference returned by `getFunc()`.
 static func::CallOp createFuncCall(OpBuilder &builder, Operation *op,
                                    StringRef name, TypeRange resultType,
                                    ValueRange operands,

mlir/test/Conversion/FuncToLLVM/calling-convention.mlir

-// 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
 // wrappers. We don't need to separate runs because the wrapper-emission
 // version subsumes the calling convention and only adds new functions, that we
 // can also file-check in the same run.
 
 // An external function is transformed into the glue around calling an interface function.

mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-callers.mlir

-// 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
 // CHECK: %{{.*}}: !llvm.ptr{{.*}} {test.returnOne}
 func.func @res_attrs_with_memref_return() -> (memref<f32> {test.returnOne}) {
   %0 = memref.alloc() : memref<f32>
   return %0 : memref<f32>
 }

mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-functions.mlir

-// 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
 // CHECK: !llvm.ptr{{.*}} {test.returnOne}
 func.func private @res_attrs_with_memref_return() -> (memref<f32> {test.returnOne})
 
 // CHECK: llvm.func @res_attrs_with_value_return() -> (f32 {test.returnOne = 1 : i64})
 // CHECK-LABEL: llvm.func @_mlir_ciface_res_attrs_with_value_return

mlir/test/Conversion/GPUToNVVM/wmma-ops-to-nvvm.mlir

 // RUN: mlir-opt --convert-gpu-to-nvvm --split-input-file %s | FileCheck %s
 // RUN: mlir-opt --convert-gpu-to-nvvm="index-bitwidth=32" --split-input-file %s | FileCheck --check-prefix=CHECK32 %s
 
 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>
     %i = arith.constant 16 : index
     %j = arith.constant 16 : index
     %0 = gpu.subgroup_mma_load_matrix %wg[%i, %j] {leadDimension = 32 : index} : memref<32x32xf16, 3> -> !gpu.mma_matrix<16x16xf16, "AOp">
     // CHECK:  %[[INX:.*]] = llvm.mlir.constant(16 : index) : i64
     // CHECK: %{{.*}} = llvm.insertvalue %{{.*}}, %{{.*}}[{{.*}}, {{.*}}]
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
     %j = arith.constant 16 : index
     gpu.subgroup_mma_store_matrix %arg0, %sg[%i,%j] {leadDimension= 32 : index} : !gpu.mma_matrix<16x16xf16, "COp">, memref<32x32xf16, 3>
     // CHECK:  %[[INX:.*]] = llvm.mlir.constant(16 : index) : i64
     // CHECK: %{{.*}} = llvm.insertvalue %{{.*}}, %{{.*}}[{{.*}}, {{.*}}]
     // CHECK:  %[[EL1:.*]] = llvm.extractvalue %[[D]][0 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>

mlir/test/Dialect/LLVMIR/request-c-wrappers.mlir

+// 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
+}

mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp

 #include "mlir/Conversion/LLVMCommon/LoweringOptions.h"
 #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
 #include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h"
 #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #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"
 #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
 #include "mlir/Dialect/SPIRV/Transforms/Passes.h"
 #include "mlir/ExecutionEngine/JitRunner.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
 #include "mlir/Pass/Pass.h"
static LogicalResult runMLIRPasses(ModuleOp module) {
   passManager.addPass(createGpuKernelOutliningPass());
   passManager.addPass(memref::createFoldSubViewOpsPass());
   passManager.addPass(createConvertGPUToSPIRVPass());
   OpPassManager &modulePM = passManager.nest<spirv::ModuleOp>();
   modulePM.addPass(spirv::createLowerABIAttributesPass());
   modulePM.addPass(spirv::createUpdateVersionCapabilityExtensionPass());
   passManager.addPass(createConvertGpuLaunchFuncToVulkanLaunchFuncPass());
   LowerToLLVMOptions llvmOptions(module.getContext(), DataLayout(module));
-  llvmOptions.emitCWrappers = true;
   passManager.addPass(createMemRefToLLVMPass());
+  passManager.nest<func::FuncOp>().addPass(LLVM::createRequestCWrappersPass());
   passManager.addPass(createConvertFuncToLLVMPass(llvmOptions));
   passManager.addPass(createReconcileUnrealizedCastsPass());
   passManager.addPass(createConvertVulkanLaunchFuncToVulkanCallsPass());
   return passManager.run(module);
 }
 
 int main(int argc, char **argv) {
   llvm::llvm_shutdown_obj x;

utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

cc_library(
     name = "LLVMIRTransforms",
     srcs = glob([
         "lib/Dialect/LLVMIR/Transforms/*.cpp",
         "lib/Dialect/LLVMIR/Transforms/*.h",
     ]),
     hdrs = glob(["include/mlir/Dialect/LLVMIR/Transforms/*.h"]),
     includes = ["include"],
     deps = [
+        ":FuncDialect",
         ":IR",
         ":LLVMDialect",
         ":LLVMPassIncGen",
         ":NVVMDialect",
         ":Pass",
         ":Transforms",
     ],
 )
deps = [
         ":FuncToLLVM",
         ":FuncToSPIRV",
         ":GPUDialect",
         ":GPUToSPIRV",
         ":GPUToVulkanTransforms",
         ":GPUTransforms",
         ":LLVMCommonConversion",
         ":LLVMDialect",
+        ":LLVMIRTransforms",
         ":LLVMToLLVMIRTranslation",
         ":MemRefDialect",
         ":MemRefToLLVM",
         ":MemRefTransforms",
         ":MlirJitRunner",
         ":Pass",
         ":ReconcileUnrealizedCasts",
         ":SPIRVDialect",
hdrs = glob([
         "include/mlir/Conversion/LinalgToStandard/*.h",
     ]),
     includes = ["include"],
     deps = [
         ":AffineDialect",
         ":ConversionPassIncGen",
         ":FuncDialect",
         ":IR",
+        ":LLVMDialect",
         ":LinalgDialect",
         ":LinalgTransforms",
         ":MemRefDialect",
         ":Pass",
         ":SCFDialect",
         ":Support",
         ":Transforms",
         "//llvm:Core",