Skip to content

Commit 610139d

Browse files
ftynseftynse
committed
[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
1 parent c263669 commit 610139d

File tree

22 files changed

+132
-39
lines changed

22 files changed

+132
-39
lines changed

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,6 @@ class LowerToLLVMOptions {
3333
LowerToLLVMOptions(MLIRContext *ctx, const DataLayout &dl);
3434

3535
bool useBarePtrCallConv = false;
36-
bool emitCWrappers = false;
3736

3837
enum class AllocLowering {
3938
/// Use malloc for for heap allocations.

‎mlir/include/mlir/Conversion/Passes.td

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -294,9 +294,6 @@ def ConvertFuncToLLVM : Pass<"convert-func-to-llvm", "ModuleOp"> {
294294
/*default=*/"false",
295295
"Replace FuncOp's MemRef arguments with bare pointers to the MemRef "
296296
"element types">,
297-
Option<"emitCWrappers", "emit-c-wrappers", "bool", /*default=*/"false",
298-
"Emit wrappers for C-compatible pointer-to-struct memref "
299-
"descriptors">,
300297
Option<"indexBitwidth", "index-bitwidth", "unsigned",
301298
/*default=kDeriveIndexBitwidthFromDataLayout*/"0",
302299
"Bitwidth of the index type, 0 to use size of machine word">,

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,11 @@ def LLVM_Dialect : Dialect {
5656

5757
/// Name of the target triple attribute.
5858
static StringRef getTargetTripleAttrName() { return "llvm.target_triple"; }
59+
60+
/// Name of the C wrapper emission attribute.
61+
static StringRef getEmitCWrapperAttrName() {
62+
return "llvm.emit_c_interface";
63+
}
5964
}];
6065

6166
let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed;

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@
1111

1212
#include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h"
1313
#include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h"
14+
#include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
1415
#include "mlir/Pass/Pass.h"
1516

1617
namespace mlir {

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

Lines changed: 15 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -13,12 +13,25 @@ include "mlir/Pass/PassBase.td"
1313

1414
def LLVMLegalizeForExport : Pass<"llvm-legalize-for-export"> {
1515
let summary = "Legalize LLVM dialect to be convertible to LLVM IR";
16-
let constructor = "mlir::LLVM::createLegalizeForExportPass()";
16+
let constructor = "::mlir::LLVM::createLegalizeForExportPass()";
17+
}
18+
19+
def LLVMRequestCWrappers
20+
: Pass<"llvm-request-c-wrappers", "::mlir::func::FuncOp"> {
21+
let summary = "Request C wrapper emission for all functions";
22+
let description = [{
23+
Annotate every builtin function in the module with the LLVM dialect
24+
attribute that instructs the conversion to LLVM to emit the C wrapper for
25+
the function. This pass is expected to be applied immediately before the
26+
conversion of builtin functions to LLVM to avoid the attribute being
27+
dropped by other passes.
28+
}];
29+
let constructor = "::mlir::LLVM::createRequestCWrappersPass()";
1730
}
1831

1932
def NVVMOptimizeForTarget : Pass<"llvm-optimize-for-nvvm-target"> {
2033
let summary = "Optimize NVVM IR";
21-
let constructor = "mlir::NVVM::createOptimizeForTargetPass()";
34+
let constructor = "::mlir::NVVM::createOptimizeForTargetPass()";
2235
}
2336

2437
#endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_PASSES

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===- RequestCWrappers.h - Annotate funcs with wrap attributes -*- C++ -*-===//
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+
#ifndef MLIR_DIALECT_LLVMIR_TRANSFORMS_REQUESTCWRAPPERS_H
10+
#define MLIR_DIALECT_LLVMIR_TRANSFORMS_REQUESTCWRAPPERS_H
11+
12+
#include <memory>
13+
14+
namespace mlir {
15+
class Pass;
16+
17+
namespace LLVM {
18+
std::unique_ptr<Pass> createRequestCWrappersPass();
19+
} // namespace LLVM
20+
} // namespace mlir
21+
22+
#endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_REQUESTCWRAPPERS_H

‎mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp

Lines changed: 6 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -361,7 +361,6 @@ struct FuncOpConversionBase : public ConvertOpToLLVMPattern<func::FuncOp> {
361361
/// FuncOp legalization pattern that converts MemRef arguments to pointers to
362362
/// MemRef descriptors (LLVM struct data types) containing all the MemRef type
363363
/// information.
364-
static constexpr StringRef kEmitIfaceAttrName = "llvm.emit_c_interface";
365364
struct FuncOpConversion : public FuncOpConversionBase {
366365
FuncOpConversion(LLVMTypeConverter &converter)
367366
: FuncOpConversionBase(converter) {}
@@ -373,8 +372,8 @@ struct FuncOpConversion : public FuncOpConversionBase {
373372
if (!newFuncOp)
374373
return failure();
375374

376-
if (getTypeConverter()->getOptions().emitCWrappers ||
377-
funcOp->getAttrOfType<UnitAttr>(kEmitIfaceAttrName)) {
375+
if (funcOp->getAttrOfType<UnitAttr>(
376+
LLVM::LLVMDialect::getEmitCWrapperAttrName())) {
378377
if (newFuncOp.isExternal())
379378
wrapExternalFunction(rewriter, funcOp.getLoc(), *getTypeConverter(),
380379
funcOp, newFuncOp);
@@ -676,24 +675,16 @@ namespace {
676675
struct ConvertFuncToLLVMPass
677676
: public ConvertFuncToLLVMBase<ConvertFuncToLLVMPass> {
678677
ConvertFuncToLLVMPass() = default;
679-
ConvertFuncToLLVMPass(bool useBarePtrCallConv, bool emitCWrappers,
680-
unsigned indexBitwidth, bool useAlignedAlloc,
678+
ConvertFuncToLLVMPass(bool useBarePtrCallConv, unsigned indexBitwidth,
679+
bool useAlignedAlloc,
681680
const llvm::DataLayout &dataLayout) {
682681
this->useBarePtrCallConv = useBarePtrCallConv;
683-
this->emitCWrappers = emitCWrappers;
684682
this->indexBitwidth = indexBitwidth;
685683
this->dataLayout = dataLayout.getStringRepresentation();
686684
}
687685

688686
/// Run the dialect converter on the module.
689687
void runOnOperation() override {
690-
if (useBarePtrCallConv && emitCWrappers) {
691-
getOperation().emitError()
692-
<< "incompatible conversion options: bare-pointer calling convention "
693-
"and C wrapper emission";
694-
signalPassFailure();
695-
return;
696-
}
697688
if (failed(LLVM::LLVMDialect::verifyDataLayoutString(
698689
this->dataLayout, [this](const Twine &message) {
699690
getOperation().emitError() << message.str();
@@ -708,7 +699,6 @@ struct ConvertFuncToLLVMPass
708699
LowerToLLVMOptions options(&getContext(),
709700
dataLayoutAnalysis.getAtOrAbove(m));
710701
options.useBarePtrCallConv = useBarePtrCallConv;
711-
options.emitCWrappers = emitCWrappers;
712702
if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout)
713703
options.overrideIndexBitwidth(indexBitwidth);
714704
options.dataLayout = llvm::DataLayout(this->dataLayout);
@@ -747,6 +737,6 @@ mlir::createConvertFuncToLLVMPass(const LowerToLLVMOptions &options) {
747737
bool useAlignedAlloc =
748738
(allocLowering == LowerToLLVMOptions::AllocLowering::AlignedAlloc);
749739
return std::make_unique<ConvertFuncToLLVMPass>(
750-
options.useBarePtrCallConv, options.emitCWrappers,
751-
options.getIndexBitwidth(), useAlignedAlloc, options.dataLayout);
740+
options.useBarePtrCallConv, options.getIndexBitwidth(), useAlignedAlloc,
741+
options.dataLayout);
752742
}

‎mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -166,18 +166,23 @@ struct LowerGpuOpsToNVVMOpsPass
166166
void runOnOperation() override {
167167
gpu::GPUModuleOp m = getOperation();
168168

169-
/// Customize the bitwidth used for the device side index computations.
169+
// Request C wrapper emission.
170+
for (auto func : m.getOps<func::FuncOp>()) {
171+
func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
172+
UnitAttr::get(&getContext()));
173+
}
174+
175+
// Customize the bitwidth used for the device side index computations.
170176
LowerToLLVMOptions options(
171177
m.getContext(),
172178
DataLayout(cast<DataLayoutOpInterface>(m.getOperation())));
173-
options.emitCWrappers = true;
174179
if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout)
175180
options.overrideIndexBitwidth(indexBitwidth);
176181

177-
/// MemRef conversion for GPU to NVVM lowering. The GPU dialect uses memory
178-
/// space 5 for private memory attributions, but NVVM represents private
179-
/// memory allocations as local `alloca`s in the default address space. This
180-
/// converter drops the private memory space to support the use case above.
182+
// MemRef conversion for GPU to NVVM lowering. The GPU dialect uses memory
183+
// space 5 for private memory attributions, but NVVM represents private
184+
// memory allocations as local `alloca`s in the default address space. This
185+
// converter drops the private memory space to support the use case above.
181186
LLVMTypeConverter converter(m.getContext(), options);
182187
converter.addConversion([&](MemRefType type) -> Optional<Type> {
183188
if (type.getMemorySpaceAsInt() !=

‎mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -63,11 +63,16 @@ struct LowerGpuOpsToROCDLOpsPass
6363
void runOnOperation() override {
6464
gpu::GPUModuleOp m = getOperation();
6565

66+
// Request C wrapper emission.
67+
for (auto func : m.getOps<func::FuncOp>()) {
68+
func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
69+
UnitAttr::get(&getContext()));
70+
}
71+
6672
/// Customize the bitwidth used for the device side index computations.
6773
LowerToLLVMOptions options(
6874
m.getContext(),
6975
DataLayout(cast<DataLayoutOpInterface>(m.getOperation())));
70-
options.emitCWrappers = true;
7176
if (indexBitwidth != kDeriveIndexBitwidthFromDataLayout)
7277
options.overrideIndexBitwidth(indexBitwidth);
7378
LLVMTypeConverter converter(m.getContext(), options);

‎mlir/lib/Conversion/LinalgToStandard/LinalgToStandard.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@
1111
#include "../PassDetail.h"
1212
#include "mlir/Dialect/Affine/IR/AffineOps.h"
1313
#include "mlir/Dialect/Func/IR/FuncOps.h"
14+
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
1415
#include "mlir/Dialect/Linalg/IR/Linalg.h"
1516
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
1617
#include "mlir/Dialect/MemRef/IR/MemRef.h"
@@ -71,7 +72,8 @@ static FlatSymbolRefAttr getLibraryCallSymbolRef(Operation *op,
7172
// Insert a function attribute that will trigger the emission of the
7273
// corresponding `_mlir_ciface_xxx` interface so that external libraries see
7374
// a normalized ABI. This interface is added during std to llvm conversion.
74-
funcOp->setAttr("llvm.emit_c_interface", UnitAttr::get(op->getContext()));
75+
funcOp->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
76+
UnitAttr::get(op->getContext()));
7577
funcOp.setPrivate();
7678
return fnNameAttr;
7779
}

‎mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,7 @@
2020
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
2121
#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h"
2222
#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h"
23+
#include "mlir/Dialect/Func/IR/FuncOps.h"
2324
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
2425
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
2526
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
@@ -280,9 +281,14 @@ class LowerHostCodeToLLVM
280281
llvm::make_early_inc_range(module.getOps<gpu::GPUModuleOp>()))
281282
gpuModule.erase();
282283

284+
// Request C wrapper emission.
285+
for (auto func : module.getOps<func::FuncOp>()) {
286+
func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
287+
UnitAttr::get(&getContext()));
288+
}
289+
283290
// Specify options to lower to LLVM and pull in the conversion patterns.
284291
LowerToLLVMOptions options(module.getContext());
285-
options.emitCWrappers = true;
286292
auto *context = module.getContext();
287293
RewritePatternSet patterns(context);
288294
LLVMTypeConverter typeConverter(context, options);

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,14 @@
11
add_mlir_dialect_library(MLIRLLVMIRTransforms
22
LegalizeForExport.cpp
33
OptimizeForNVVM.cpp
4+
RequestCWrappers.cpp
45

56
DEPENDS
67
MLIRLLVMPassIncGen
78

89
LINK_LIBS PUBLIC
910
MLIRIR
11+
MLIRFuncDialect
1012
MLIRLLVMDialect
1113
MLIRPass
1214
MLIRTransforms

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,9 @@
1212
#include "mlir/Pass/Pass.h"
1313

1414
namespace mlir {
15+
namespace func {
16+
class FuncOp;
17+
} // namespace func
1518

1619
#define GEN_PASS_CLASSES
1720
#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"

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

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,29 @@
1+
//===- RequestCWrappers.cpp - Annotate funcs with wrap attributes ---------===//
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+
#include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
10+
#include "PassDetail.h"
11+
#include "mlir/Dialect/Func/IR/FuncOps.h"
12+
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
13+
14+
using namespace mlir;
15+
16+
namespace {
17+
class RequestCWrappersPass
18+
: public LLVMRequestCWrappersBase<RequestCWrappersPass> {
19+
public:
20+
void runOnOperation() override {
21+
getOperation()->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
22+
UnitAttr::get(&getContext()));
23+
}
24+
};
25+
} // namespace
26+
27+
std::unique_ptr<Pass> mlir::LLVM::createRequestCWrappersPass() {
28+
return std::make_unique<RequestCWrappersPass>();
29+
}

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -65,7 +65,8 @@ static FlatSymbolRefAttr getFunc(Operation *op, StringRef name,
6565
FunctionType::get(context, operands.getTypes(), resultType));
6666
func.setPrivate();
6767
if (static_cast<bool>(emitCInterface))
68-
func->setAttr("llvm.emit_c_interface", UnitAttr::get(context));
68+
func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
69+
UnitAttr::get(context));
6970
}
7071
return result;
7172
}

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: mlir-opt -convert-memref-to-llvm -convert-func-to-llvm='emit-c-wrappers=1' -reconcile-unrealized-casts %s | FileCheck %s
1+
// RUN: mlir-opt -convert-memref-to-llvm -llvm-request-c-wrappers -convert-func-to-llvm -reconcile-unrealized-casts %s | FileCheck %s
22
// RUN: mlir-opt -convert-memref-to-llvm -convert-func-to-llvm -reconcile-unrealized-casts %s | FileCheck %s --check-prefix=EMIT_C_ATTRIBUTE
33

44
// This tests the default memref calling convention and the emission of C

‎mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-callers.mlir renamed to ‎mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-callers.mlir

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: mlir-opt -convert-func-to-llvm='emit-c-wrappers=1' %s | FileCheck %s
1+
// RUN: mlir-opt -llvm-request-c-wrappers -convert-func-to-llvm %s | FileCheck %s
22

33
// CHECK: llvm.func @res_attrs_with_memref_return() -> (!llvm.struct{{.*}} {test.returnOne})
44
// CHECK-LABEL: llvm.func @_mlir_ciface_res_attrs_with_memref_return

‎mlir/test/Conversion/StandardToLLVM/emit-c-wrappers-for-external-functions.mlir renamed to ‎mlir/test/Conversion/FuncToLLVM/emit-c-wrappers-for-external-functions.mlir

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: mlir-opt -convert-func-to-llvm='emit-c-wrappers=1' %s | FileCheck %s
1+
// RUN: mlir-opt -llvm-request-c-wrappers -convert-func-to-llvm %s | FileCheck %s
22

33
// CHECK: llvm.func @res_attrs_with_memref_return() -> (!llvm.struct{{.*}} {test.returnOne})
44
// CHECK-LABEL: llvm.func @_mlir_ciface_res_attrs_with_memref_return

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44
gpu.module @test_module {
55

66
// CHECK-LABEL: func @gpu_wmma_load_op() ->
7-
// CHECK-SAME: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> {
7+
// CHECK-SAME: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
88
// CHECK32-LABEL: func @gpu_wmma_load_op() ->
99
func.func @gpu_wmma_load_op() -> (!gpu.mma_matrix<16x16xf16, "AOp">) {
1010
%wg = memref.alloca() {alignment = 32} : memref<32x32xf16, 3>
@@ -43,9 +43,9 @@ gpu.module @test_module {
4343
gpu.module @test_module {
4444

4545
// CHECK-LABEL: func @gpu_wmma_store_op
46-
// CHECK-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>) {
46+
// CHECK-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>)
4747
// CHECK32-LABEL: func @gpu_wmma_store_op
48-
// CHECK32-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>) {
48+
// CHECK32-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>)
4949
func.func @gpu_wmma_store_op(%arg0 : !gpu.mma_matrix<16x16xf16, "COp">) -> () {
5050
%sg = memref.alloca(){alignment = 32} : memref<32x32xf16, 3>
5151
%i = arith.constant 16 : index

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

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
// RUN: mlir-opt %s -llvm-request-c-wrappers | FileCheck %s
2+
3+
// CHECK: func.func private @foo() attributes {llvm.emit_c_interface}
4+
func.func private @foo()
5+
6+
// CHECK: func.func @bar() attributes {llvm.emit_c_interface}
7+
func.func @bar() {
8+
return
9+
}

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@
2424
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
2525
#include "mlir/Dialect/GPU/Transforms/Passes.h"
2626
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
27+
#include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
2728
#include "mlir/Dialect/MemRef/IR/MemRef.h"
2829
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
2930
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
@@ -52,8 +53,8 @@ static LogicalResult runMLIRPasses(ModuleOp module) {
5253
modulePM.addPass(spirv::createUpdateVersionCapabilityExtensionPass());
5354
passManager.addPass(createConvertGpuLaunchFuncToVulkanLaunchFuncPass());
5455
LowerToLLVMOptions llvmOptions(module.getContext(), DataLayout(module));
55-
llvmOptions.emitCWrappers = true;
5656
passManager.addPass(createMemRefToLLVMPass());
57+
passManager.nest<func::FuncOp>().addPass(LLVM::createRequestCWrappersPass());
5758
passManager.addPass(createConvertFuncToLLVMPass(llvmOptions));
5859
passManager.addPass(createReconcileUnrealizedCastsPass());
5960
passManager.addPass(createConvertVulkanLaunchFuncToVulkanCallsPass());

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3397,6 +3397,7 @@ cc_library(
33973397
hdrs = glob(["include/mlir/Dialect/LLVMIR/Transforms/*.h"]),
33983398
includes = ["include"],
33993399
deps = [
3400+
":FuncDialect",
34003401
":IR",
34013402
":LLVMDialect",
34023403
":LLVMPassIncGen",
@@ -6484,6 +6485,7 @@ cc_binary(
64846485
":GPUTransforms",
64856486
":LLVMCommonConversion",
64866487
":LLVMDialect",
6488+
":LLVMIRTransforms",
64876489
":LLVMToLLVMIRTranslation",
64886490
":MemRefDialect",
64896491
":MemRefToLLVM",
@@ -7273,6 +7275,7 @@ cc_library(
72737275
":ConversionPassIncGen",
72747276
":FuncDialect",
72757277
":IR",
7278+
":LLVMDialect",
72767279
":LinalgDialect",
72777280
":LinalgTransforms",
72787281
":MemRefDialect",

0 commit comments

Comments
 (0)