[Mlir-commits] [mlir] [RFC][mlir][func] Enforce func.return as sole terminator in func.func regions (PR #184778)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Thu Mar 5 04:08:10 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-spirv
Author: Mehdi Amini (joker-eph)
<details>
<summary>Changes</summary>
Adds a verifier check to FuncOp that requires all block terminators inside a func.func body to be func::ReturnOp instances. Previously any op implementing RegionBranchTerminatorOpInterface (e.g. llvm.return, spirv.ReturnValue, test.return) was silently accepted, which allowed structurally invalid IR.
To preserve partial-conversion semantics, ReturnOpLowering in FuncToLLVM now checks that its parent is already an llvm.func before converting func.return; similarly FuncToSPIRV's ReturnOpPattern bails out when the enclosing func.func cannot be fully converted. Tests across ~50 files are updated accordingly:
- test.return → func.return in Transforms legalizer tests
- spirv.Return/ReturnValue → return in SPIRV IR/Transforms tests
- func.func → llvm.func in LLVM dialect tests that use llvm.return
- llvm.return → return in OpenMP/Transforms/CAPI tests
- CHECK patterns and expected-error annotations updated for conversion tests
I'm not sure we should do this, opening this more as an RFC right now: trying to see if this is worth it or not @<!-- -->jpienaar / @<!-- -->matthias-springer ?
---
Patch is 154.29 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/184778.diff
54 Files Affected:
- (modified) mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp (+5)
- (modified) mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp (+18)
- (modified) mlir/lib/Dialect/Func/IR/FuncOps.cpp (+15-9)
- (modified) mlir/test/CAPI/execution_engine.c (+1-1)
- (modified) mlir/test/Conversion/FuncToLLVM/func-memref-return.mlir (+1-1)
- (modified) mlir/test/Conversion/FuncToLLVM/func-to-llvm.mlir (+2-2)
- (modified) mlir/test/Conversion/MemRefToLLVM/invalid.mlir (+1-1)
- (modified) mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir (+1-1)
- (modified) mlir/test/Conversion/SCFToSPIRV/if.mlir (+1-1)
- (modified) mlir/test/Conversion/SCFToSPIRV/unsupported.mlir (+1-1)
- (modified) mlir/test/Dialect/Bufferization/Transforms/OwnershipBasedBufferDeallocation/dealloc-region-branchop-interface.mlir (+2-2)
- (modified) mlir/test/Dialect/EmitC/invalid_ops.mlir (+1-1)
- (modified) mlir/test/Dialect/Func/invalid.mlir (+6-8)
- (modified) mlir/test/Dialect/LLVMIR/global.mlir (+5-5)
- (modified) mlir/test/Dialect/LLVMIR/inlining.mlir (+2-2)
- (modified) mlir/test/Dialect/LLVMIR/invalid.mlir (+36-36)
- (modified) mlir/test/Dialect/LLVMIR/nvvm.mlir (+27-27)
- (modified) mlir/test/Dialect/LLVMIR/rocdl.mlir (+10-10)
- (modified) mlir/test/Dialect/LLVMIR/roundtrip.mlir (+11-11)
- (modified) mlir/test/Dialect/LLVMIR/terminator.mlir (+2-2)
- (modified) mlir/test/Dialect/OpenMP/invalid-fuse.mlir (+5-5)
- (modified) mlir/test/Dialect/OpenMP/invalid-tile.mlir (+4-4)
- (modified) mlir/test/Dialect/OpenMP/invalid.mlir (+1-1)
- (modified) mlir/test/Dialect/SPIRV/IR/availability.mlir (+2-2)
- (modified) mlir/test/Dialect/SPIRV/IR/bit-ops.mlir (+12-12)
- (modified) mlir/test/Dialect/SPIRV/IR/cast-ops.mlir (+26-26)
- (modified) mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir (+49-52)
- (modified) mlir/test/Dialect/SPIRV/IR/image-ops.mlir (+50-50)
- (modified) mlir/test/Dialect/SPIRV/IR/memory-ops.mlir (+5-5)
- (modified) mlir/test/Dialect/SPIRV/IR/mesh-ops.mlir (+1-1)
- (modified) mlir/test/Dialect/SPIRV/IR/misc-ops.mlir (+5-5)
- (modified) mlir/test/Dialect/SPIRV/IR/primitive-ops.mlir (+2-2)
- (modified) mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir (+33-33)
- (modified) mlir/test/Dialect/SPIRV/Transforms/gl-canonicalize.mlir (+24-24)
- (modified) mlir/test/IR/test-region-branch-op-verifier.mlir (+3-9)
- (modified) mlir/test/Target/LLVMIR/llvmir-invalid.mlir (+1-2)
- (modified) mlir/test/Transforms/print-op-graph.mlir (+3-3)
- (modified) mlir/test/Transforms/remove-dead-values.mlir (+2-2)
- (modified) mlir/test/Transforms/sccp.mlir (+1-1)
- (modified) mlir/test/Transforms/test-dialect-conversion-pdll.mlir (+2-2)
- (modified) mlir/test/Transforms/test-legalize-remapped-value.mlir (+2-2)
- (modified) mlir/test/Transforms/test-legalizer-analysis.mlir (+1)
- (modified) mlir/test/Transforms/test-legalizer-fold-after.mlir (+1-1)
- (modified) mlir/test/Transforms/test-legalizer-fold-before.mlir (+1-1)
- (modified) mlir/test/Transforms/test-legalizer-full-rollback.mlir (+4-4)
- (modified) mlir/test/Transforms/test-legalizer-full.mlir (+8-8)
- (modified) mlir/test/Transforms/test-legalizer-no-fold.mlir (+1-1)
- (modified) mlir/test/Transforms/test-legalizer-no-materializations.mlir (-3)
- (modified) mlir/test/Transforms/test-legalizer-no-rollback.mlir (+1-1)
- (modified) mlir/test/Transforms/test-legalizer-rollback.mlir (+6-8)
- (modified) mlir/test/Transforms/test-legalizer.mlir (+15-33)
- (modified) mlir/test/Transforms/test-merge-blocks.mlir (+3-3)
- (modified) mlir/test/Transforms/test-pattern-selective-replacement.mlir (+2-2)
- (modified) mlir/test/lib/Dialect/Test/TestPatterns.cpp (+30-4)
``````````diff
diff --git a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
index 2220f61ed8a07..2e08a301d5c17 100644
--- a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
+++ b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
@@ -719,7 +719,12 @@ struct ReturnOpLowering : public ConvertOpToLLVMPattern<func::ReturnOp> {
Location loc = op.getLoc();
SmallVector<Value, 4> updatedOperands;
+ // Only convert when the enclosing func.func has already been converted to
+ // llvm.func. If the parent is still a func::FuncOp (e.g., because its
+ // signature has non-convertible types), leave func.return as-is.
auto funcOp = op->getParentOfType<LLVM::LLVMFuncOp>();
+ if (!funcOp)
+ return failure();
bool useBarePtrCallConv =
shouldUseBarePtrCallConv(funcOp, this->getTypeConverter());
diff --git a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp
index 7ceab394a67ff..58e6299b77c78 100644
--- a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp
+++ b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp
@@ -42,6 +42,24 @@ class ReturnOpPattern final : public OpConversionPattern<func::ReturnOp> {
if (returnOp.getNumOperands() > 1)
return failure();
+ // Only convert func.return when the enclosing func.func is being converted
+ // to spirv.func. Check that all argument and result types of the parent
+ // function are convertible; if not, leave this op unconverted so the
+ // func.func stays in the Func dialect with a valid func.return terminator.
+ if (auto funcOp = returnOp->getParentOfType<func::FuncOp>()) {
+ FunctionType fnType = funcOp.getFunctionType();
+ if (fnType.getNumResults() > 1)
+ return failure();
+ for (Type argType : fnType.getInputs()) {
+ if (!getTypeConverter()->convertType(argType))
+ return failure();
+ }
+ if (fnType.getNumResults() == 1) {
+ if (!getTypeConverter()->convertType(fnType.getResult(0)))
+ return failure();
+ }
+ }
+
if (returnOp.getNumOperands() == 1) {
rewriter.replaceOpWithNewOp<spirv::ReturnValueOp>(
returnOp, adaptor.getOperands()[0]);
diff --git a/mlir/lib/Dialect/Func/IR/FuncOps.cpp b/mlir/lib/Dialect/Func/IR/FuncOps.cpp
index d803e99154499..c30e36a7c6281 100644
--- a/mlir/lib/Dialect/Func/IR/FuncOps.cpp
+++ b/mlir/lib/Dialect/Func/IR/FuncOps.cpp
@@ -295,21 +295,27 @@ LogicalResult FuncOp::verifyRegions() {
continue;
// Check func.return or other return-like terminators ops (e.g.
// llvm.return, test.return).
- auto returnOp = dyn_cast<RegionBranchTerminatorOpInterface>(&block.back());
- if (!returnOp)
+ auto regionTerminatorOp =
+ dyn_cast<RegionBranchTerminatorOpInterface>(&block.back());
+ if (!regionTerminatorOp)
continue;
- auto operands =
- returnOp.getMutableSuccessorOperands(RegionSuccessor::parent());
- if (operands.size() != resultTypes.size())
+ auto returnOp = dyn_cast<ReturnOp>(&block.back());
+ if (!returnOp)
+ return regionTerminatorOp->emitOpError(
+ "is not a func.return op: func.func op is expected to have a "
+ "func.return op as the only region terminator");
+
+ if (returnOp->getNumOperands() != resultTypes.size())
return returnOp->emitOpError("has ")
- << operands.size() << " operands, but enclosing function (@"
+ << returnOp->getNumOperands() << " operands, but enclosing function (@"
<< getName() << ") returns " << resultTypes.size();
- for (auto [i, opType] : llvm::enumerate(llvm::zip(operands, resultTypes))) {
+ for (auto [i, opType] :
+ llvm::enumerate(llvm::zip(returnOp->getOperands(), resultTypes))) {
auto [operand, resTy] = opType;
- if (operand.get().getType() != resTy)
+ if (operand.getType() != resTy)
return returnOp->emitError() << "type of return operand " << i << " ("
- << operand.get().getType()
+ << operand.getType()
<< ") doesn't match function result type ("
<< resTy << ") in function @" << getName();
}
diff --git a/mlir/test/CAPI/execution_engine.c b/mlir/test/CAPI/execution_engine.c
index 4df232f3caab3..1898e6dace053 100644
--- a/mlir/test/CAPI/execution_engine.c
+++ b/mlir/test/CAPI/execution_engine.c
@@ -110,7 +110,7 @@ void testOmpCreation(void) {
" } \n"
" omp.terminator \n"
" } \n"
-" llvm.return \n"
+" return \n"
" } \n"
"} \n"
));
diff --git a/mlir/test/Conversion/FuncToLLVM/func-memref-return.mlir b/mlir/test/Conversion/FuncToLLVM/func-memref-return.mlir
index 22ebbf8618bde..78e569835776e 100644
--- a/mlir/test/Conversion/FuncToLLVM/func-memref-return.mlir
+++ b/mlir/test/Conversion/FuncToLLVM/func-memref-return.mlir
@@ -97,7 +97,7 @@ func.func @unconvertible_multiresult(%arg0: memref<?xf32> , %arg1: memref<?xf32>
// BAREPTR-SAME: %{{.*}}: memref<*xi32>)
func.func @unranked_memref(%arg0:memref<*xi32>) {
// BAREPTR: call @printMemrefI32(%arg{{.*}}) : (memref<*xi32>) -> ()
- // BAREPTR-NEXT: llvm.return
+ // BAREPTR-NEXT: return
call @printMemrefI32(%arg0) : (memref<*xi32>) -> ()
return
}
diff --git a/mlir/test/Conversion/FuncToLLVM/func-to-llvm.mlir b/mlir/test/Conversion/FuncToLLVM/func-to-llvm.mlir
index 94dfceadbc449..ef18765f2ce64 100644
--- a/mlir/test/Conversion/FuncToLLVM/func-to-llvm.mlir
+++ b/mlir/test/Conversion/FuncToLLVM/func-to-llvm.mlir
@@ -557,9 +557,9 @@ func.func @index_arg(%arg0: index) -> index {
}
// There is no type conversion rule for tf32, so vector<1xtf32> and, therefore,
-// the func op cannot be converted.
+// the func op and its body remain unchanged.
// CHECK: func.func @non_convertible_arg_type({{.*}}: vector<1xtf32>)
-// CHECK: llvm.return
+// CHECK: return
func.func @non_convertible_arg_type(%arg: vector<1xtf32>) {
return
}
diff --git a/mlir/test/Conversion/MemRefToLLVM/invalid.mlir b/mlir/test/Conversion/MemRefToLLVM/invalid.mlir
index 5462d3278d9e6..4330baef70542 100644
--- a/mlir/test/Conversion/MemRefToLLVM/invalid.mlir
+++ b/mlir/test/Conversion/MemRefToLLVM/invalid.mlir
@@ -4,7 +4,7 @@
llvm.func @malloc(i64)
func.func @redef_reserved() {
%alloc = memref.alloc() : memref<1024x64xf32, 1>
- llvm.return
+ return
}
// -----
diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index c4b8e93b6a9f9..fb11b7cdc8989 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -88,7 +88,7 @@ func.func @cp_async_mbarrier_arrive(%bar_shared: !llvm.ptr<3>, %bar_gen: !llvm.p
nvvm.cp.async.mbarrier.arrive %bar_shared : !llvm.ptr<3>
// CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} {noinc = true}
nvvm.cp.async.mbarrier.arrive %bar_shared {noinc = true} : !llvm.ptr<3>
- llvm.return
+ return
}
// CHECK-LABEL: @tma_load_3d_all
diff --git a/mlir/test/Conversion/SCFToSPIRV/if.mlir b/mlir/test/Conversion/SCFToSPIRV/if.mlir
index 2c18da41dc021..e0f67fce5c9d0 100644
--- a/mlir/test/Conversion/SCFToSPIRV/if.mlir
+++ b/mlir/test/Conversion/SCFToSPIRV/if.mlir
@@ -158,7 +158,7 @@ func.func @simple_if_yield_type_change(%arg2 : memref<10xf32, #spirv.storage_cla
func.func @unsupported_yield_type(%arg0 : memref<8xi32>, %arg1 : memref<8xi32>, %c : i1) {
// CHECK-LABEL: @unsupported_yield_type
// CHECK-NEXT: scf.if
-// CHECK: spirv.Return
+// CHECK: return
%r = scf.if %c -> (memref<8xi32>) {
scf.yield %arg0 : memref<8xi32>
} else {
diff --git a/mlir/test/Conversion/SCFToSPIRV/unsupported.mlir b/mlir/test/Conversion/SCFToSPIRV/unsupported.mlir
index 1a1c24a09aa8c..b2852a8979bf9 100644
--- a/mlir/test/Conversion/SCFToSPIRV/unsupported.mlir
+++ b/mlir/test/Conversion/SCFToSPIRV/unsupported.mlir
@@ -8,7 +8,7 @@
// CHECK-NEXT: spirv.Constant
// CHECK-NEXT: memref.store
// CHECK-NEXT: scf.reduce
-// CHECK: spirv.Return
+// CHECK: return
func.func @func(%arg0: i64) {
%0 = arith.index_cast %arg0 : i64 to index
%alloc = memref.alloc() : memref<16xf32>
diff --git a/mlir/test/Dialect/Bufferization/Transforms/OwnershipBasedBufferDeallocation/dealloc-region-branchop-interface.mlir b/mlir/test/Dialect/Bufferization/Transforms/OwnershipBasedBufferDeallocation/dealloc-region-branchop-interface.mlir
index 423fc4730b137..75a67ceaf3fe8 100644
--- a/mlir/test/Dialect/Bufferization/Transforms/OwnershipBasedBufferDeallocation/dealloc-region-branchop-interface.mlir
+++ b/mlir/test/Dialect/Bufferization/Transforms/OwnershipBasedBufferDeallocation/dealloc-region-branchop-interface.mlir
@@ -525,7 +525,7 @@ func.func @noRegionBranchOpInterface() {
}) : () -> (i32)
"test.return"() : () -> ()
}) : () -> (i32)
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -544,7 +544,7 @@ func.func @noRegionBranchOpInterface() {
}) : () -> (memref<2xi32>)
"test.return"() : () -> ()
}) : () -> (i32)
- "test.return"() : () -> ()
+ func.return
}
// -----
diff --git a/mlir/test/Dialect/EmitC/invalid_ops.mlir b/mlir/test/Dialect/EmitC/invalid_ops.mlir
index 0d878e90cdf0c..471b6701bac51 100644
--- a/mlir/test/Dialect/EmitC/invalid_ops.mlir
+++ b/mlir/test/Dialect/EmitC/invalid_ops.mlir
@@ -428,7 +428,7 @@ emitc.func @return_type_array(%arg : !emitc.array<4xi32>) -> !emitc.array<4xi32>
// -----
func.func @return_inside_func.func(%0: i32) -> (i32) {
- // expected-error at +1 {{'emitc.return' op expects parent op 'emitc.func'}}
+ // expected-error at +1 {{expects parent op 'emitc.func'}}
emitc.return %0 : i32
}
// -----
diff --git a/mlir/test/Dialect/Func/invalid.mlir b/mlir/test/Dialect/Func/invalid.mlir
index 3143bda77ebba..3aa6015b87ad3 100644
--- a/mlir/test/Dialect/Func/invalid.mlir
+++ b/mlir/test/Dialect/Func/invalid.mlir
@@ -194,23 +194,21 @@ func.func @foo() {} // expected-error {{expected non-empty function body}}
// -----
-// test.loop_block_term implements RegionBranchTerminatorOpInterface.
-// getMutableSuccessorOperands(parent) returns only the exit operand (f32).
-// The function returns i32, so the type check must fail.
+// test.loop_block_term implements RegionBranchTerminatorOpInterface but is not
+// a func.return op. func.func requires func.return as the only terminator.
func.func @region_branch_term_type_mismatch(%arg: i32) -> i32 {
%0 = "test.constant"() { value = 5.3 : f32 } : () -> f32
- // expected-error @+1 {{type of return operand 0 ('f32') doesn't match function result type ('i32') in function @region_branch_term_type_mismatch}}
+ // expected-error @+1 {{is not a func.return op}}
test.loop_block_term iter %arg exit %0
}
// -----
-// test.loop_block_term has one exit operand (f32) but the function returns
-// nothing. getMutableSuccessorOperands(parent) returns 1 operand while the
-// function has 0 results, so the count check must fail.
+// test.loop_block_term implements RegionBranchTerminatorOpInterface but is not
+// a func.return op. func.func requires func.return as the only terminator.
func.func @region_branch_term_count_mismatch(%arg: i32) {
%0 = "test.constant"() { value = 5.3 : f32 } : () -> f32
- // expected-error @+1 {{'test.loop_block_term' op has 1 operands, but enclosing function (@region_branch_term_count_mismatch) returns 0}}
+ // expected-error @+1 {{is not a func.return op}}
test.loop_block_term iter %arg exit %0
}
diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir
index 193ab7987a2b6..9c8f7dcfce0a1 100644
--- a/mlir/test/Dialect/LLVMIR/global.mlir
+++ b/mlir/test/Dialect/LLVMIR/global.mlir
@@ -73,7 +73,7 @@ llvm.comdat @__llvm_comdat {
llvm.mlir.global @any() comdat(@__llvm_comdat::@any) {addr_space = 1 : i32} : i64
// CHECK-LABEL: references
-func.func @references() {
+llvm.func @references() {
// CHECK: llvm.mlir.addressof @".string" : !llvm.ptr
%0 = llvm.mlir.addressof @".string" : !llvm.ptr
@@ -164,7 +164,7 @@ func.func @bar() {
// -----
-func.func @foo() {
+llvm.func @foo() {
// The attribute parser will consume the first colon-type, so we put two of
// them to trigger the attribute type mismatch error.
// expected-error @+1 {{invalid kind of attribute specified}}
@@ -177,7 +177,7 @@ func.func @foo() {
func.func @foo() {
// expected-error @+1 {{must reference a global defined by 'llvm.mlir.global'}}
llvm.mlir.addressof @foo : !llvm.ptr
- llvm.return
+ return
}
// -----
@@ -206,7 +206,7 @@ llvm.mlir.global internal @g(43 : i64) : i64 {
// -----
llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
-func.func @mismatch_addr_space_implicit_global() {
+llvm.func @mismatch_addr_space_implicit_global() {
// expected-error @+1 {{pointer address space must match address space of the referenced global}}
llvm.mlir.addressof @g : !llvm.ptr
llvm.return
@@ -216,7 +216,7 @@ func.func @mismatch_addr_space_implicit_global() {
llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
-func.func @mismatch_addr_space() {
+llvm.func @mismatch_addr_space() {
// expected-error @+1 {{pointer address space must match address space of the referenced global}}
llvm.mlir.addressof @g : !llvm.ptr<4>
llvm.return
diff --git a/mlir/test/Dialect/LLVMIR/inlining.mlir b/mlir/test/Dialect/LLVMIR/inlining.mlir
index 70ce7ca20986b..5587f6aea23ed 100644
--- a/mlir/test/Dialect/LLVMIR/inlining.mlir
+++ b/mlir/test/Dialect/LLVMIR/inlining.mlir
@@ -61,7 +61,7 @@ func.func @test_inline(%ptr : !llvm.ptr) -> i32 {
// Check that llvm.return is correctly handled
func.func @func(%arg0 : i32) -> i32 {
- llvm.return %arg0 : i32
+ return %arg0 : i32
}
// CHECK-LABEL: @llvm_ret
// CHECK-NOT: call
@@ -698,7 +698,7 @@ llvm.func @caller(%x : i32) -> i32 {
func.func @func(%arg0 : i32) -> i32 {
llvm.blocktag <id = 1>
- llvm.return %arg0 : i32
+ return %arg0 : i32
}
// CHECK-LABEL: @llvm_ret
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index 5068ddc42e1e5..0ea5268d53947 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -225,7 +225,7 @@ func.func @store_unaligned_atomic(%val : f32, %ptr : !llvm.ptr) {
// -----
-func.func @invalid_call() {
+llvm.func @invalid_call() {
// expected-error at +1 {{'llvm.call' op must have either a `callee` attribute or at least an operand}}
"llvm.call"() {op_bundle_sizes = array<i32>} : () -> ()
llvm.return
@@ -236,14 +236,14 @@ func.func @invalid_call() {
func.func @call_missing_ptr_type(%callee : !llvm.func<i8 (i8)>, %arg : i8) {
// expected-error at +1 {{expected indirect call to have 2 trailing types}}
llvm.call %callee(%arg) : (i8) -> (i8)
- llvm.return
+ return
}
// -----
func.func private @standard_func_callee()
-func.func @call_missing_ptr_type(%arg : i8) {
+llvm.func @call_missing_ptr_type(%arg : i8) {
// expected-error at +2 {{expected '('}}
// expected-error at +1 {{expected direct call to have 1 trailing type}}
llvm.call @standard_func_callee(%arg) : !llvm.ptr, (i8) -> (i8)
@@ -255,12 +255,12 @@ func.func @call_missing_ptr_type(%arg : i8) {
func.func @call_non_pointer_type(%callee : !llvm.func<i8 (i8)>, %arg : i8) {
// expected-error at +1 {{indirect call expects a pointer as callee: '!llvm.func<i8 (i8)>'}}
llvm.call %callee(%arg) : !llvm.func<i8 (i8)>, (i8) -> (i8)
- llvm.return
+ return
}
// -----
-func.func @call_non_function_type(%callee : !llvm.ptr, %arg : i8) {
+llvm.func @call_non_function_type(%callee : !llvm.ptr, %arg : i8) {
// expected-error at +2 {{expected '('}}
// expected-error at +1 {{expected trailing function type}}
llvm.call %callee(%arg) : !llvm.ptr, !llvm.func<i8 (i8)>
@@ -269,7 +269,7 @@ func.func @call_non_function_type(%callee : !llvm.ptr, %arg : i8) {
// -----
-func.func @call_void_result_type(%callee : !llvm.ptr, %arg : i8) {
+llvm.func @call_void_result_type(%callee : !llvm.ptr, %arg : i8) {
// expected-error at +1 {{expected a non-void result type}}
llvm.call %callee(%arg) : !llvm.ptr, (i8) -> (!llvm.void)
llvm.return
@@ -277,7 +277,7 @@ func.func @call_void_result_type(%callee : !llvm.ptr, %arg : i8) {
// -----
-func.func @call_unknown_symbol() {
+llvm.func @call_unknown_symbol() {
// expected-error at +1 {{'llvm.call' op 'missing_callee' does not reference a symbol in the current scope}}
llvm.call @missing_callee() : () -> ()
llvm.return
@@ -287,7 +287,7 @@ func.func @call_unknown_symbol() {
func.func private @standard_func_callee()
-func.func @call_non_llvm() {
+llvm.func @call_non_llvm() {
// expected-error at +1 {{'llvm.call' op 'standard_func_callee' does not reference a valid LLVM function}}
llvm.call @standard_func_callee() : () -> ()
llvm.return
@@ -298,12 +298,12 @@ func.func @call_non_llvm() {
func.func @call_non_llvm_arg(%arg0 : tensor<*xi32>) {
// expected-error at +1 {{'llvm.call' op operand #0 must be variadic of LLVM dialect-compatible type}}
"llvm.call"(%arg0) {operandSegmentSizes = array<i32: 1, 0>, op_bundle_sizes = array<i32>} : (tensor<*xi32>) -> ()
- llvm.return
+ return
}
// -----
-func.func @call_non_llvm_res(%callee : !llvm.ptr) {
+llvm.func @call_non_llvm_res(%callee : !llvm.ptr) {
// expected-error at +1 {{'llvm.call' op result #0 must be LLVM dialect-compatible type}}
llvm.call %callee() : !llvm.ptr, () -> (tensor<*xi32>)
llvm.return
@@ -313,7 +313,7 @@ func.func @call_non_llvm_res(%callee : !llvm.ptr) {
llvm.func @callee_func(i8) -> ()
-func.func @callee_arg_mismatch(%arg0 : i32) {
+llvm.func @callee_arg_mismatch(%arg0 : i32) {
// expected-error at +1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}}
llvm.call @callee_func(%arg0) : (i32) -> ()
llvm.return
@@ -323,7 +323,7 @@ func.func @callee_arg_mismatch(%arg0 : i32) {
llvm.func @callee_func() -> (i8)
-func.func @callee_return_mismatch() {
+llvm.func @callee_return_mismatch() {
// expected-error at +1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}}
%res = llvm.call @callee_func() : () -> (i32)
llvm.return
@@ -331,7 +331,7 @@ func.func @callee_return_mismatch() {
// -----
-func.func @call_too_many_results(%callee : !llvm.ptr) {
+llvm.func @call_too_many_results(%callee : !llvm.ptr) {
// expected-error at +1 {{expected function with 0 or 1 result}}
llvm.call %callee() : !llvm.ptr, () -> (i32, i32)
llvm.return
@@ -684,7 +684,7 @@ func.func @nvvm_invalid_shfl_pred_3(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3
// -----
-func.func @nvvm_invalid_mma_0(%a0 : f16, %a1 : f16,
+llvm.func @nvvm_invalid_mma_0(%a0 : f16, %a1 : f16,
%b0 : vector<2xf16>, %b1 : vector<2xf16>,
%c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,
%c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> {
@@ -696,7 +696,7 @@ func.func @nvvm_invalid_mma_0(%a0 : f16, %a1 : f16,
// -----
-func.func @nvvm_invalid_mma_1(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_invalid_mma_1(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
%b0 : vector<2xf16>, %b1 : vector<2xf16>,
%c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,
%c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)> {
@@ -708,7 +708,7 @@ func.func @nvvm_invalid_mma_1(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
// -----
-func.func @nvvm_invalid_mma_2(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_invalid_mma_2(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
%b0 : vector<2xf16>, %b1 : vector<2xf16>,
%c0 : f32, %c1 ...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/184778
More information about the Mlir-commits
mailing list