[Mlir-commits] [mlir] [RFC][mlir][func] Enforce func.return as sole terminator in func.func regions (PR #184778)
Mehdi Amini
llvmlistbot at llvm.org
Thu Mar 5 04:07:31 PST 2026
https://github.com/joker-eph created https://github.com/llvm/llvm-project/pull/184778
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 ?
>From 0d6ccf0eaa6fc333acd6288ef40c9c7d8b307b38 Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Thu, 5 Mar 2026 03:23:06 -0800
Subject: [PATCH] [mlir][func] Enforce func.return as sole terminator in
func.func regions
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
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
---
mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp | 5 +
.../Conversion/FuncToSPIRV/FuncToSPIRV.cpp | 18 ++++
mlir/lib/Dialect/Func/IR/FuncOps.cpp | 24 +++--
mlir/test/CAPI/execution_engine.c | 2 +-
.../FuncToLLVM/func-memref-return.mlir | 2 +-
.../Conversion/FuncToLLVM/func-to-llvm.mlir | 4 +-
.../test/Conversion/MemRefToLLVM/invalid.mlir | 2 +-
.../Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 2 +-
mlir/test/Conversion/SCFToSPIRV/if.mlir | 2 +-
.../Conversion/SCFToSPIRV/unsupported.mlir | 2 +-
.../dealloc-region-branchop-interface.mlir | 4 +-
mlir/test/Dialect/EmitC/invalid_ops.mlir | 2 +-
mlir/test/Dialect/Func/invalid.mlir | 14 ++-
mlir/test/Dialect/LLVMIR/global.mlir | 10 +-
mlir/test/Dialect/LLVMIR/inlining.mlir | 4 +-
mlir/test/Dialect/LLVMIR/invalid.mlir | 72 ++++++-------
mlir/test/Dialect/LLVMIR/nvvm.mlir | 54 +++++-----
mlir/test/Dialect/LLVMIR/rocdl.mlir | 20 ++--
mlir/test/Dialect/LLVMIR/roundtrip.mlir | 22 ++--
mlir/test/Dialect/LLVMIR/terminator.mlir | 4 +-
mlir/test/Dialect/OpenMP/invalid-fuse.mlir | 10 +-
mlir/test/Dialect/OpenMP/invalid-tile.mlir | 8 +-
mlir/test/Dialect/OpenMP/invalid.mlir | 2 +-
mlir/test/Dialect/SPIRV/IR/availability.mlir | 4 +-
mlir/test/Dialect/SPIRV/IR/bit-ops.mlir | 24 ++---
mlir/test/Dialect/SPIRV/IR/cast-ops.mlir | 52 ++++-----
.../Dialect/SPIRV/IR/control-flow-ops.mlir | 101 +++++++++---------
mlir/test/Dialect/SPIRV/IR/image-ops.mlir | 100 ++++++++---------
mlir/test/Dialect/SPIRV/IR/memory-ops.mlir | 10 +-
mlir/test/Dialect/SPIRV/IR/mesh-ops.mlir | 2 +-
mlir/test/Dialect/SPIRV/IR/misc-ops.mlir | 10 +-
mlir/test/Dialect/SPIRV/IR/primitive-ops.mlir | 4 +-
.../SPIRV/Transforms/canonicalize.mlir | 66 ++++++------
.../SPIRV/Transforms/gl-canonicalize.mlir | 48 ++++-----
.../IR/test-region-branch-op-verifier.mlir | 12 +--
mlir/test/Target/LLVMIR/llvmir-invalid.mlir | 3 +-
mlir/test/Transforms/print-op-graph.mlir | 6 +-
mlir/test/Transforms/remove-dead-values.mlir | 4 +-
mlir/test/Transforms/sccp.mlir | 2 +-
.../test-dialect-conversion-pdll.mlir | 4 +-
.../test-legalize-remapped-value.mlir | 4 +-
.../Transforms/test-legalizer-analysis.mlir | 1 +
.../Transforms/test-legalizer-fold-after.mlir | 2 +-
.../test-legalizer-fold-before.mlir | 2 +-
.../test-legalizer-full-rollback.mlir | 8 +-
mlir/test/Transforms/test-legalizer-full.mlir | 16 +--
.../Transforms/test-legalizer-no-fold.mlir | 2 +-
.../test-legalizer-no-materializations.mlir | 3 -
.../test-legalizer-no-rollback.mlir | 2 +-
.../Transforms/test-legalizer-rollback.mlir | 14 ++-
mlir/test/Transforms/test-legalizer.mlir | 48 +++------
mlir/test/Transforms/test-merge-blocks.mlir | 6 +-
.../test-pattern-selective-replacement.mlir | 4 +-
mlir/test/lib/Dialect/Test/TestPatterns.cpp | 34 +++++-
54 files changed, 454 insertions(+), 433 deletions(-)
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 : f32, %c2 : f32, %c3 : f32,
%c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> {
@@ -720,7 +720,7 @@ func.func @nvvm_invalid_mma_2(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
// -----
-func.func @nvvm_invalid_mma_3(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_invalid_mma_3(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
%b0 : vector<2xf16>, %b1 : vector<2xf16>,
%c0 : vector<2xf16>, %c1 : vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> {
// expected-error at +1 {{unimplemented variant for MMA shape <8, 8, 16>}}
@@ -730,7 +730,7 @@ func.func @nvvm_invalid_mma_3(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
// -----
-func.func @nvvm_invalid_mma_8(%a0 : i32, %a1 : i32,
+llvm.func @nvvm_invalid_mma_8(%a0 : i32, %a1 : i32,
%b0 : i32,
%c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// expected-error at +1 {{op requires b1Op attribute}}
@@ -773,7 +773,7 @@ llvm.func @nvvm_mma_m16n8k16_f16_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
// -----
-func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) {
+llvm.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) {
// expected-error at +1 {{op failed to verify that result #0 and operand #1 have the same type}}
%0 = "llvm.atomicrmw"(%f32_ptr, %f32) {bin_op=11, ordering=1} : (!llvm.ptr, f32) -> i32
llvm.return
@@ -781,7 +781,7 @@ func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) {
// -----
-func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr, %i32 : i32) {
+llvm.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{expected LLVM IR floating point type}}
%0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : !llvm.ptr, i32
llvm.return
@@ -789,7 +789,7 @@ func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr, %i32 : i32) {
// -----
-func.func @atomicrmw_scalable_vector(%ptr : !llvm.ptr, %f32_vec : vector<[2]xf32>) {
+llvm.func @atomicrmw_scalable_vector(%ptr : !llvm.ptr, %f32_vec : vector<[2]xf32>) {
// expected-error at +1 {{'val' must be floating point LLVM type or LLVM pointer type or signless integer or LLVM dialect-compatible fixed-length vector type}}
%0 = llvm.atomicrmw fadd %ptr, %f32_vec unordered : !llvm.ptr, vector<[2]xf32>
llvm.return
@@ -797,7 +797,7 @@ func.func @atomicrmw_scalable_vector(%ptr : !llvm.ptr, %f32_vec : vector<[2]xf32
// -----
-func.func @atomicrmw_vector_expected_float(%ptr : !llvm.ptr, %i32_vec : vector<3xi32>) {
+llvm.func @atomicrmw_vector_expected_float(%ptr : !llvm.ptr, %i32_vec : vector<3xi32>) {
// expected-error at +1 {{expected LLVM IR floating point type for vector element}}
%0 = llvm.atomicrmw fadd %ptr, %i32_vec unordered : !llvm.ptr, vector<3xi32>
llvm.return
@@ -805,7 +805,7 @@ func.func @atomicrmw_vector_expected_float(%ptr : !llvm.ptr, %i32_vec : vector<3
// -----
-func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
+llvm.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
// expected-error at +1 {{unexpected LLVM IR type for 'xchg' bin_op}}
%0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : !llvm.ptr, i1
llvm.return
@@ -813,7 +813,7 @@ func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
// -----
-func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {
+llvm.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {
// expected-error at +1 {{expected LLVM IR integer type}}
%0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : !llvm.ptr, f32
llvm.return
@@ -821,7 +821,7 @@ func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {
// -----
-func.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) {
+llvm.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) {
// expected-error at +1 {{op failed to verify that operand #1 and operand #2 have the same type}}
%0 = "llvm.cmpxchg"(%ptr, %i32, %i64) {success_ordering=2,failure_ordering=2} : (!llvm.ptr, i32, i64) -> !llvm.struct<(i32, i1)>
llvm.return
@@ -829,7 +829,7 @@ func.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64
// -----
-func.func @cmpxchg_mismatched_result(%ptr : !llvm.ptr, %i64 : i64) {
+llvm.func @cmpxchg_mismatched_result(%ptr : !llvm.ptr, %i64 : i64) {
// expected-error at +1 {{op failed to verify that result #0 has an LLVM struct type consisting of the type of operand #2 and a bool}}
%0 = "llvm.cmpxchg"(%ptr, %i64, %i64) {success_ordering=2,failure_ordering=2} : (!llvm.ptr, i64, i64) -> !llvm.struct<(i64, i64)>
llvm.return
@@ -837,7 +837,7 @@ func.func @cmpxchg_mismatched_result(%ptr : !llvm.ptr, %i64 : i64) {
// -----
-func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
+llvm.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
// expected-error at +1 {{unexpected LLVM IR type}}
%0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : !llvm.ptr, i1
llvm.return
@@ -845,7 +845,7 @@ func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
// -----
-func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr, %i32 : i32) {
+llvm.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{ordering must be at least 'monotonic'}}
%0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : !llvm.ptr, i32
llvm.return
@@ -853,7 +853,7 @@ func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr, %i32 : i32)
// -----
-func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr, %i32 : i32) {
+llvm.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{ordering must be at least 'monotonic'}}
%0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : !llvm.ptr, i32
llvm.return
@@ -861,7 +861,7 @@ func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr, %i32 : i32)
// -----
-func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr, %i32 : i32) {
+llvm.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{failure ordering cannot be 'release' or 'acq_rel'}}
%0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : !llvm.ptr, i32
llvm.return
@@ -869,7 +869,7 @@ func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr, %i32 : i32) {
// -----
-func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr, %i32 : i32) {
+llvm.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{failure ordering cannot be 'release' or 'acq_rel'}}
%0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr, i32
llvm.return
@@ -998,7 +998,7 @@ module attributes {llvm.data_layout = "#vjkr32"} {
// -----
-func.func @switch_superfluous_comma(%arg0 : i64) {
+llvm.func @switch_superfluous_comma(%arg0 : i64) {
// expected-error at +3 {{custom op 'llvm.switch' expected integer value}}
llvm.switch %arg0 : i32, ^bb1 [
42: ^bb2,
@@ -1011,7 +1011,7 @@ func.func @switch_superfluous_comma(%arg0 : i64) {
// -----
-func.func @switch_wrong_number_of_weights(%arg0 : i32) {
+llvm.func @switch_wrong_number_of_weights(%arg0 : i32) {
// expected-error at +1 {{expects number of branch weights to match number of successors: 3 vs 2}}
llvm.switch %arg0 : i32, ^bb1 [
42: ^bb2(%arg0, %arg0 : i32, i32)
@@ -1026,7 +1026,7 @@ func.func @switch_wrong_number_of_weights(%arg0 : i32) {
// -----
-func.func @switch_case_type_mismatch(%arg0 : i64) {
+llvm.func @switch_case_type_mismatch(%arg0 : i64) {
// expected-error at below {{expects case value type to match condition value type}}
"llvm.switch"(%arg0)[^bb1, ^bb2] <{case_operand_segments = array<i32: 0>, case_values = dense<42> : vector<1xi32>, operandSegmentSizes = array<i32: 1, 0, 0>}> : (i64) -> ()
^bb1: // pred: ^bb0
@@ -1491,7 +1491,7 @@ func.func @invalid_target_ext_atomic(%arg0 : !llvm.ptr) {
// -----
-func.func @invalid_target_ext_constant_unsupported() {
+llvm.func @invalid_target_ext_constant_unsupported() {
// expected-error at +1 {{target extension type does not support zero-initializer}}
%0 = llvm.mlir.zero : !llvm.target<"invalid_constant">
llvm.return
@@ -1499,7 +1499,7 @@ func.func @invalid_target_ext_constant_unsupported() {
// -----
-func.func @invalid_target_ext_constant() {
+llvm.func @invalid_target_ext_constant() {
// expected-error at +1 {{does not support target extension type.}}
%0 = llvm.mlir.constant(0 : index) : !llvm.target<"spirv.Event">
llvm.return
diff --git a/mlir/test/Dialect/LLVMIR/nvvm.mlir b/mlir/test/Dialect/LLVMIR/nvvm.mlir
index 6b7417b4b82bc..2553352c903b1 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-opt %s -split-input-file -verify-diagnostics | FileCheck %s
// CHECK-LABEL: @nvvm_special_regs
-func.func @nvvm_special_regs() -> i32 {
+llvm.func @nvvm_special_regs() -> i32 {
// CHECK: nvvm.read.ptx.sreg.tid.x : i32
%0 = nvvm.read.ptx.sreg.tid.x : i32
// CHECK: nvvm.read.ptx.sreg.tid.y : i32
@@ -30,14 +30,14 @@ func.func @nvvm_special_regs() -> i32 {
}
// CHECK-LABEL: @nvvm_rcp
-func.func @nvvm_rcp(%arg0: f32) -> f32 {
+llvm.func @nvvm_rcp(%arg0: f32) -> f32 {
// CHECK: nvvm.rcp.approx.ftz.f %arg0 : f32
%0 = nvvm.rcp.approx.ftz.f %arg0 : f32
llvm.return %0 : f32
}
// CHECK-LABEL: @llvm_nvvm_barrier0
-func.func @llvm_nvvm_barrier0() {
+llvm.func @llvm_nvvm_barrier0() {
// CHECK: nvvm.barrier0
nvvm.barrier0
llvm.return
@@ -66,7 +66,7 @@ llvm.func @llvm_nvvm_barrier_arrive(%barId : i32, %numberOfThreads : i32) {
}
// CHECK-LABEL: @llvm_nvvm_cluster_arrive
-func.func @llvm_nvvm_cluster_arrive() {
+llvm.func @llvm_nvvm_cluster_arrive() {
// CHECK: nvvm.cluster.arrive
nvvm.cluster.arrive
// CHECK: nvvm.cluster.arrive {aligned}
@@ -75,7 +75,7 @@ func.func @llvm_nvvm_cluster_arrive() {
}
// CHECK-LABEL: @llvm_nvvm_cluster_arrive_relaxed
-func.func @llvm_nvvm_cluster_arrive_relaxed() {
+llvm.func @llvm_nvvm_cluster_arrive_relaxed() {
// CHECK: nvvm.cluster.arrive.relaxed
nvvm.cluster.arrive.relaxed
// CHECK: nvvm.cluster.arrive.relaxed {aligned}
@@ -84,7 +84,7 @@ func.func @llvm_nvvm_cluster_arrive_relaxed() {
}
// CHECK-LABEL: @llvm_nvvm_cluster_wait
-func.func @llvm_nvvm_cluster_wait() {
+llvm.func @llvm_nvvm_cluster_wait() {
// CHECK: nvvm.cluster.wait
nvvm.cluster.wait
// CHECK: nvvm.cluster.wait {aligned}
@@ -93,7 +93,7 @@ func.func @llvm_nvvm_cluster_wait() {
}
// CHECK-LABEL: @nvvm_shfl
-func.func @nvvm_shfl(
+llvm.func @nvvm_shfl(
%arg0 : i32, %arg1 : i32, %arg2 : i32,
%arg3 : i32, %arg4 : f32) -> i32 {
// CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : i32 -> i32
@@ -110,7 +110,7 @@ func.func @nvvm_shfl(
}
// CHECK-LABEL: @nvvm_shfl_pred
-func.func @nvvm_shfl_pred(
+llvm.func @nvvm_shfl_pred(
%arg0 : i32, %arg1 : i32, %arg2 : i32,
%arg3 : i32, %arg4 : f32) -> !llvm.struct<(i32, i1)> {
// CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i1)>
@@ -121,7 +121,7 @@ func.func @nvvm_shfl_pred(
}
// CHECK-LABEL: @nvvm_vote(
-func.func @nvvm_vote(%arg0 : i32, %arg1 : i1) -> i32 {
+llvm.func @nvvm_vote(%arg0 : i32, %arg1 : i1) -> i32 {
// CHECK: nvvm.vote.sync ballot %{{.*}}, %{{.*}} -> i32
%0 = nvvm.vote.sync ballot %arg0, %arg1 -> i32
// CHECK: nvvm.vote.sync all %{{.*}}, %{{.*}} -> i1
@@ -134,14 +134,14 @@ func.func @nvvm_vote(%arg0 : i32, %arg1 : i1) -> i32 {
}
// CHECK-LABEL: @llvm_nvvm_bar_warp_sync
-func.func @llvm_nvvm_bar_warp_sync(%mask : i32) {
+llvm.func @llvm_nvvm_bar_warp_sync(%mask : i32) {
// CHECK: nvvm.bar.warp.sync %{{.*}}
nvvm.bar.warp.sync %mask : i32
llvm.return
}
// CHECK-LABEL: @nvvm_mma_m8n8k4_row_col_f32_f32
-func.func @nvvm_mma_m8n8k4_row_col_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_mma_m8n8k4_row_col_f32_f32(%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, f32)> {
// CHECK: nvvm.mma.sync
@@ -152,7 +152,7 @@ func.func @nvvm_mma_m8n8k4_row_col_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf
}
// CHECK-LABEL: @nvvm_mma_m8n8k4_f16_f16
-func.func @nvvm_mma_m8n8k4_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_mma_m8n8k4_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
%b0 : vector<2xf16>, %b1 : vector<2xf16>,
%c0 : vector<2xf16>, %c1 : vector<2xf16>, %c2 : vector<2xf16>, %c3 : vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> {
// CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}]
@@ -163,7 +163,7 @@ func.func @nvvm_mma_m8n8k4_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
}
// CHECK-LABEL: @nvvm_mma_m16n8k8_bf16_bf16
-func.func @nvvm_mma_m16n8k8_bf16_bf16(%a0 : i32, %a1 : i32, %b0 : i32,
+llvm.func @nvvm_mma_m16n8k8_bf16_bf16(%a0 : i32, %a1 : i32, %b0 : i32,
%c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(f32, f32, f32, f32)> {
// CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>, shape = #nvvm.shape<m = 16, n = 8, k = 8>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>
%0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]
@@ -174,7 +174,7 @@ func.func @nvvm_mma_m16n8k8_bf16_bf16(%a0 : i32, %a1 : i32, %b0 : i32,
}
// CHECK-LABEL: @nvvm_mma_m16n8k16_bf16_bf16
-func.func @nvvm_mma_m16n8k16_bf16_bf16(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,
+llvm.func @nvvm_mma_m16n8k16_bf16_bf16(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,
%b0 : i32, %b1 : i32,
%c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(f32, f32, f32, f32)> {
// CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>
@@ -186,7 +186,7 @@ func.func @nvvm_mma_m16n8k16_bf16_bf16(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i3
}
// CHECK-LABEL: @nvvm_mma_m8n8k16_s8_s8
-func.func @nvvm_mma_m8n8k16_s8_s8(%a0 : i32, %b0 : i32,
+llvm.func @nvvm_mma_m8n8k16_s8_s8(%a0 : i32, %b0 : i32,
%c0 : i32, %c1 : i32) -> !llvm.struct<(i32, i32)> {
// CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<wrapped>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<s8>, shape = #nvvm.shape<m = 8, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32, i32)>
%0 = nvvm.mma.sync A[%a0] B[%b0] C[%c0, %c1]
@@ -198,7 +198,7 @@ func.func @nvvm_mma_m8n8k16_s8_s8(%a0 : i32, %b0 : i32,
}
// CHECK-LABEL: @nvvm_mma_m16n8k8_f16_f16
-func.func @nvvm_mma_m16n8k8_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_mma_m16n8k8_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
%b0 : vector<2xf16>,
%c0 : vector<2xf16>, %c1 : vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> {
// CHECK: nvvm.mma.sync A[%{{.*}}, %{{.*}}] B[%{{.*}}] C[%{{.*}}, %{{.*}}] {{{.*}}} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
@@ -209,7 +209,7 @@ func.func @nvvm_mma_m16n8k8_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
}
// CHECK-LABEL: @nvvm_mma_m16n8k16_f16_f16
-func.func @nvvm_mma_m16n8k16_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_mma_m16n8k16_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
%a2 : vector<2xf16>, %a3 : vector<2xf16>,
%b0 : vector<2xf16>, %b1 : vector<2xf16>,
%c0 : vector<2xf16>, %c1 : vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> {
@@ -221,7 +221,7 @@ func.func @nvvm_mma_m16n8k16_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
}
// CHECK-LABEL: @nvvm_mma_m16n8k16_f32_f32
-func.func @nvvm_mma_m16n8k16_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
+llvm.func @nvvm_mma_m16n8k16_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
%a2 : vector<2xf16>, %a3 : vector<2xf16>,
%b0 : vector<2xf16>, %b1 : vector<2xf16>,
%c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(f32, f32, f32, f32)> {
@@ -233,7 +233,7 @@ func.func @nvvm_mma_m16n8k16_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
}
// CHECK-LABEL: @nvvm_mma_m16n8k4_tf32_f32
-func.func @nvvm_mma_m16n8k4_tf32_f32(%a0 : i32, %a1 : i32,
+llvm.func @nvvm_mma_m16n8k4_tf32_f32(%a0 : i32, %a1 : i32,
%b0 : i32,
%c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(f32, f32, f32, f32)> {
// CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<tf32>, multiplicandBPtxType = #nvvm.mma_type<tf32>, shape = #nvvm.shape<m = 16, n = 8, k = 4>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>
@@ -245,7 +245,7 @@ func.func @nvvm_mma_m16n8k4_tf32_f32(%a0 : i32, %a1 : i32,
}
// CHECK-LABEL: @nvvm_mma_m16n8k16_s8_s8
-func.func @nvvm_mma_m16n8k16_s8_s8(%a0 : i32, %a1 : i32, %b0 : i32,
+llvm.func @nvvm_mma_m16n8k16_s8_s8(%a0 : i32, %a1 : i32, %b0 : i32,
%c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<wrapped>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<s8>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>
%0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]
@@ -257,7 +257,7 @@ func.func @nvvm_mma_m16n8k16_s8_s8(%a0 : i32, %a1 : i32, %b0 : i32,
}
// CHECK-LABEL: @nvvm_mma_m16n8k16_s8_u8
-func.func @nvvm_mma_m16n8k16_s8_u8(%a0 : i32, %a1 : i32,
+llvm.func @nvvm_mma_m16n8k16_s8_u8(%a0 : i32, %a1 : i32,
%b0 : i32,
%c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<satfinite>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<u8>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>
@@ -270,7 +270,7 @@ func.func @nvvm_mma_m16n8k16_s8_u8(%a0 : i32, %a1 : i32,
}
// CHECK-LABEL: @nvvm_mma_m16n8k256_b1_b1
-func.func @nvvm_mma_m16n8k256_b1_b1(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,
+llvm.func @nvvm_mma_m16n8k256_b1_b1(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,
%b0 : i32, %b1 : i32,
%c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}, {{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {b1Op = #nvvm.mma_b1op<xor_popc>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>, shape = #nvvm.shape<m = 16, n = 8, k = 256>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>
@@ -282,7 +282,7 @@ func.func @nvvm_mma_m16n8k256_b1_b1(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,
}
// CHECK-LABEL: @nvvm_mma_m16n8k128_b1_b1
-func.func @nvvm_mma_m16n8k128_b1_b1(%a0 : i32, %a1 : i32,
+llvm.func @nvvm_mma_m16n8k128_b1_b1(%a0 : i32, %a1 : i32,
%b0 : i32,
%c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {b1Op = #nvvm.mma_b1op<xor_popc>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>, shape = #nvvm.shape<m = 16, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>
@@ -295,7 +295,7 @@ func.func @nvvm_mma_m16n8k128_b1_b1(%a0 : i32, %a1 : i32,
}
// CHECK-LABEL: @nvvm_mma_m8n8k128_b1_b1
-func.func @nvvm_mma_m8n8k128_b1_b1(%a0 : i32,
+llvm.func @nvvm_mma_m8n8k128_b1_b1(%a0 : i32,
%b0 : i32,
%c0 : i32, %c1 : i32) -> !llvm.struct<(i32, i32)> {
// CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}] {b1Op = #nvvm.mma_b1op<xor_popc>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>, shape = #nvvm.shape<m = 8, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32, i32)>
@@ -307,7 +307,7 @@ func.func @nvvm_mma_m8n8k128_b1_b1(%a0 : i32,
}
// CHECK-LABEL: @nvvm_mma_m16n8k32_s4_s4
-func.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32,
+llvm.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32,
%b0 : i32,
%c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<wrapped>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s4>, multiplicandBPtxType = #nvvm.mma_type<s4>, shape = #nvvm.shape<m = 16, n = 8, k = 32>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>
@@ -320,7 +320,7 @@ func.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32,
}
// CHECK-LABEL: @nvvm_wmma_load_tf32
-func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
+llvm.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// CHECK: nvvm.wmma.load {{.*}} {eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
%0 = nvvm.wmma.load %arg0, %arg1
{eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
@@ -329,7 +329,7 @@ func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr, %arg1 : i32) -> !llvm.struct<(i
}
// CHECK-LABEL: @nvvm_wmma_mma
-func.func @nvvm_wmma_mma(%0 : i32, %1 : i32, %2 : i32, %3 : i32, %4 : i32, %5 : i32,
+llvm.func @nvvm_wmma_mma(%0 : i32, %1 : i32, %2 : i32, %3 : i32, %4 : i32, %5 : i32,
%6 : i32, %7 : i32, %8 : f32, %9 : f32, %10 : f32,
%11 : f32, %12 : f32, %13 : f32, %14 : f32, %15 : f32)
-> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> {
diff --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir
index 1a810dce706bd..203b5e8df6fe9 100644
--- a/mlir/test/Dialect/LLVMIR/rocdl.mlir
+++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir
@@ -1,6 +1,6 @@
// RUN: mlir-opt %s -split-input-file -verify-diagnostics | FileCheck %s
-func.func @rocdl_special_regs() -> i32 {
+llvm.func @rocdl_special_regs() -> i32 {
// CHECK-LABEL: rocdl_special_regs
// CHECK: rocdl.workitem.id.x : i32
%0 = rocdl.workitem.id.x : i32
@@ -43,21 +43,21 @@ func.func @rocdl_special_regs() -> i32 {
llvm.return %0 : i32
}
-func.func @rocdl.fmed3.scalar(%a: f32, %b: f32, %c: f32) -> f32 {
+llvm.func @rocdl.fmed3.scalar(%a: f32, %b: f32, %c: f32) -> f32 {
// CHECK-LABEL: rocdl.fmed3.scalar
// CHECK: %0 = rocdl.fmed3 %arg0, %arg1, %arg2 : f32
%0 = rocdl.fmed3 %a, %b, %c : f32
llvm.return %0 : f32
}
-func.func @rocdl.fmed3.vector(%a: vector<4xf16>, %b: vector<4xf16>, %c: vector<4xf16>) -> vector<4xf16> {
+llvm.func @rocdl.fmed3.vector(%a: vector<4xf16>, %b: vector<4xf16>, %c: vector<4xf16>) -> vector<4xf16> {
// CHECK-LABEL: rocdl.fmed3.vector
// CHECK: %0 = rocdl.fmed3 %arg0, %arg1, %arg2 : vector<4xf16>
%0 = rocdl.fmed3 %a, %b, %c : vector<4xf16>
llvm.return %0 : vector<4xf16>
}
-func.func @rocdl.math.ops(%a: f32, %b: f16, %c: bf16) {
+llvm.func @rocdl.math.ops(%a: f32, %b: f16, %c: bf16) {
// CHECK-LABEL: rocdl.math.ops
// CHECK: %{{.*}} = rocdl.tanh %{{.*}} f32 -> f32
// CHECK: %{{.*}} = rocdl.tanh %{{.*}} f16 -> f16
@@ -117,37 +117,37 @@ func.func @rocdl.math.ops(%a: f32, %b: f16, %c: bf16) {
llvm.return
}
-func.func @rocdl.barrier() {
+llvm.func @rocdl.barrier() {
// CHECK: rocdl.barrier
rocdl.barrier
llvm.return
}
-func.func @rocdl.sched_barrier() {
+llvm.func @rocdl.sched_barrier() {
// CHECK: rocdl.sched.barrier
rocdl.sched.barrier 0
llvm.return
}
-func.func @rocdl_sched_group_barrier() {
+llvm.func @rocdl_sched_group_barrier() {
// CHECK: rocdl.sched.group.barrier
rocdl.sched.group.barrier 8, 1, 0
llvm.return
}
-func.func @rocdl_iglp_opt() {
+llvm.func @rocdl_iglp_opt() {
// CHECK: rocdl.iglp.opt
rocdl.iglp.opt 0
llvm.return
}
-func.func @rocdl.setprio() {
+llvm.func @rocdl.setprio() {
// CHECK: rocdl.s.setprio
rocdl.s.setprio 0
llvm.return
}
-func.func @rocdl.xdlops(%arg0 : f32, %arg1 : f32,
+llvm.func @rocdl.xdlops(%arg0 : f32, %arg1 : f32,
%arg2 : vector<32xf32>, %arg3 : i32,
%arg4 : vector<16xf32>, %arg5 : vector<4xf32>,
%arg6 : vector<4xf16>, %arg7 : vector<32xi32>,
diff --git a/mlir/test/Dialect/LLVMIR/roundtrip.mlir b/mlir/test/Dialect/LLVMIR/roundtrip.mlir
index c680d0d98ac5f..4e101070a5cdd 100644
--- a/mlir/test/Dialect/LLVMIR/roundtrip.mlir
+++ b/mlir/test/Dialect/LLVMIR/roundtrip.mlir
@@ -7,7 +7,7 @@ llvm.func @baz()
// CHECK-LABEL: func @ops
// CHECK-SAME: (%[[I32:.*]]: i32, %[[FLOAT:.*]]: f32, %[[PTR1:.*]]: !llvm.ptr, %[[PTR2:.*]]: !llvm.ptr, %[[BOOL:.*]]: i1, %[[VPTR1:.*]]: vector<2x!llvm.ptr>)
-func.func @ops(%arg0: i32, %arg1: f32,
+llvm.func @ops(%arg0: i32, %arg1: f32,
%arg2: !llvm.ptr, %arg3: !llvm.ptr,
%arg4: i1, %arg5 : vector<2x!llvm.ptr>) {
// Integer arithmetic binary operations.
@@ -369,7 +369,7 @@ llvm.func @foo(%arg0: i32) -> !llvm.struct<(i32, f64, i32)> {
// CHECK-LABEL: @casts
// CHECK-SAME: (%[[I32:.*]]: i32, %[[I64:.*]]: i64, %[[V4I32:.*]]: vector<4xi32>, %[[V4I64:.*]]: vector<4xi64>, %[[PTR:.*]]: !llvm.ptr)
-func.func @casts(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
+llvm.func @casts(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
%arg3: vector<4xi64>, %arg4: !llvm.ptr) {
// CHECK: = llvm.sext %[[I32]] : i32 to i56
%0 = llvm.sext %arg0 : i32 to i56
@@ -400,7 +400,7 @@ func.func @casts(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
// CHECK-LABEL: @nneg_casts
// CHECK-SAME: (%[[I32:.*]]: i32, %[[I64:.*]]: i64, %[[V4I32:.*]]: vector<4xi32>, %[[V4I64:.*]]: vector<4xi64>, %[[PTR:.*]]: !llvm.ptr)
-func.func @nneg_casts(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
+llvm.func @nneg_casts(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
%arg3: vector<4xi64>, %arg4: !llvm.ptr) {
// CHECK: = llvm.zext nneg %[[I32]] : i32 to i64
%0 = llvm.zext nneg %arg0 : i32 to i64
@@ -413,7 +413,7 @@ func.func @nneg_casts(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
// CHECK-LABEL: @casts_overflow
// CHECK-SAME: (%[[I32:.*]]: i32, %[[I64:.*]]: i64, %[[V4I32:.*]]: vector<4xi32>, %[[V4I64:.*]]: vector<4xi64>, %[[PTR:.*]]: !llvm.ptr)
-func.func @casts_overflow(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
+llvm.func @casts_overflow(%arg0: i32, %arg1: i64, %arg2: vector<4xi32>,
%arg3: vector<4xi64>, %arg4: !llvm.ptr) {
// CHECK: = llvm.trunc %[[I64]] overflow<nsw> : i64 to i56
%0 = llvm.trunc %arg1 overflow<nsw> : i64 to i56
@@ -488,7 +488,7 @@ func.func @vector_deinterleave2(%vec: vector<[8]xf16>) {
}
// CHECK-LABEL: @alloca
-func.func @alloca(%size : i64) {
+llvm.func @alloca(%size : i64) {
// CHECK: llvm.alloca %{{.*}} x i32 : (i64) -> !llvm.ptr
llvm.alloca %size x i32 {alignment = 0} : (i64) -> (!llvm.ptr)
// CHECK: llvm.alloca inalloca %{{.*}} x i32 {alignment = 8 : i64} : (i64) -> !llvm.ptr
@@ -497,21 +497,21 @@ func.func @alloca(%size : i64) {
}
// CHECK-LABEL: @null
-func.func @null() {
+llvm.func @null() {
// CHECK: llvm.mlir.zero : !llvm.ptr
%0 = llvm.mlir.zero : !llvm.ptr
llvm.return
}
// CHECK-LABEL: @zero
-func.func @zero() {
+llvm.func @zero() {
// CHECK: llvm.mlir.zero : i8
%0 = llvm.mlir.zero : i8
llvm.return
}
// CHECK-LABEL: @atomic_load
-func.func @atomic_load(%ptr : !llvm.ptr) {
+llvm.func @atomic_load(%ptr : !llvm.ptr) {
// CHECK: llvm.load %{{.*}} atomic monotonic {alignment = 4 : i64} : !llvm.ptr -> f32
%0 = llvm.load %ptr atomic monotonic {alignment = 4 : i64} : !llvm.ptr -> f32
// CHECK: llvm.load volatile %{{.*}} atomic syncscope("singlethread") monotonic {alignment = 16 : i64} : !llvm.ptr -> f32
@@ -522,7 +522,7 @@ func.func @atomic_load(%ptr : !llvm.ptr) {
}
// CHECK-LABEL: @atomic_store
-func.func @atomic_store(%val : f32, %large_val : i256, %ptr : !llvm.ptr) {
+llvm.func @atomic_store(%val : f32, %large_val : i256, %ptr : !llvm.ptr) {
// CHECK: llvm.store %{{.*}}, %{{.*}} atomic monotonic {alignment = 4 : i64} : f32, !llvm.ptr
llvm.store %val, %ptr atomic monotonic {alignment = 4 : i64} : f32, !llvm.ptr
// CHECK: llvm.store volatile %{{.*}}, %{{.*}} atomic syncscope("singlethread") monotonic {alignment = 16 : i64} : f32, !llvm.ptr
@@ -533,7 +533,7 @@ func.func @atomic_store(%val : f32, %large_val : i256, %ptr : !llvm.ptr) {
}
// CHECK-LABEL: @atomicrmw
-func.func @atomicrmw(%ptr : !llvm.ptr, %f32 : f32, %f16_vec : vector<2xf16>) {
+llvm.func @atomicrmw(%ptr : !llvm.ptr, %f32 : f32, %f16_vec : vector<2xf16>) {
// CHECK: llvm.atomicrmw fadd %{{.*}}, %{{.*}} monotonic : !llvm.ptr, f32
%0 = llvm.atomicrmw fadd %ptr, %f32 monotonic : !llvm.ptr, f32
// CHECK: llvm.atomicrmw volatile fsub %{{.*}}, %{{.*}} syncscope("singlethread") monotonic {alignment = 16 : i64} : !llvm.ptr, f32
@@ -544,7 +544,7 @@ func.func @atomicrmw(%ptr : !llvm.ptr, %f32 : f32, %f16_vec : vector<2xf16>) {
}
// CHECK-LABEL: @cmpxchg
-func.func @cmpxchg(%ptr : !llvm.ptr, %cmp : i32, %new : i32) {
+llvm.func @cmpxchg(%ptr : !llvm.ptr, %cmp : i32, %new : i32) {
// CHECK: llvm.cmpxchg %{{.*}}, %{{.*}}, %{{.*}} acq_rel monotonic : !llvm.ptr, i32
%0 = llvm.cmpxchg %ptr, %cmp, %new acq_rel monotonic : !llvm.ptr, i32
// CHECK: llvm.cmpxchg weak volatile %{{.*}}, %{{.*}}, %{{.*}} syncscope("singlethread") acq_rel monotonic {alignment = 16 : i64} : !llvm.ptr, i32
diff --git a/mlir/test/Dialect/LLVMIR/terminator.mlir b/mlir/test/Dialect/LLVMIR/terminator.mlir
index 86b70735d3efa..8b628b27d6bae 100644
--- a/mlir/test/Dialect/LLVMIR/terminator.mlir
+++ b/mlir/test/Dialect/LLVMIR/terminator.mlir
@@ -3,7 +3,7 @@
// CHECK-LABEL: @return
// CHECK: llvm.return
-func.func @return() {
+llvm.func @return() {
llvm.return
}
@@ -11,7 +11,7 @@ func.func @return() {
// CHECK: llvm.br
// CHECK: llvm.cond_br
// CHECK: llvm.return
-func.func @control_flow(%cond : i1) {
+llvm.func @control_flow(%cond : i1) {
llvm.br ^bb1
^bb1:
llvm.cond_br %cond, ^bb2, ^bb1
diff --git a/mlir/test/Dialect/OpenMP/invalid-fuse.mlir b/mlir/test/Dialect/OpenMP/invalid-fuse.mlir
index ffd1c7300e79b..3eb50703c1ce5 100644
--- a/mlir/test/Dialect/OpenMP/invalid-fuse.mlir
+++ b/mlir/test/Dialect/OpenMP/invalid-fuse.mlir
@@ -48,9 +48,9 @@ func.func @wrong_generatees1(%tc1 : i32, %tc2 : i32) {
%fused1 = omp.new_cli
%fused2 = omp.new_cli
// expected-error at +1 {{'omp.fuse' op in a complete fuse the number of generatees must be exactly 1}}
- omp.fuse (%fused1, %fused2) <-(%canonloop1, %canonloop2)
+ omp.fuse (%fused1, %fused2) <-(%canonloop1, %canonloop2)
- llvm.return
+ return
}
// -----
@@ -70,10 +70,10 @@ func.func @wrong_generatees2(%tc1 : i32, %tc2 : i32, %tc3 : i32) {
}
%fused = omp.new_cli
- // expected-error at +1 {{'omp.fuse' op the number of generatees must be the number of aplyees plus one minus count}}
+ // expected-error at +1 {{'omp.fuse' op the number of generatees must be the number of aplyees plus one minus count}}
omp.fuse (%fused) <-(%canonloop1, %canonloop2, %canonloop3) looprange(first = 1, count = 2)
- llvm.return
+ return
}
// -----
@@ -97,6 +97,6 @@ func.func @wrong_applyees(%tc1 : i32, %tc2 : i32, %tc3 : i32) {
// expected-error at +1 {{'omp.fuse' op the numbers of applyees must be at least first minus one plus count attributes}}
omp.fuse (%fused, %canonloop_fuse) <-(%canonloop1, %canonloop2, %canonloop3) looprange(first = 1, count = 5)
- llvm.return
+ return
}
diff --git a/mlir/test/Dialect/OpenMP/invalid-tile.mlir b/mlir/test/Dialect/OpenMP/invalid-tile.mlir
index e63a062d810ed..42cd3acd1753d 100644
--- a/mlir/test/Dialect/OpenMP/invalid-tile.mlir
+++ b/mlir/test/Dialect/OpenMP/invalid-tile.mlir
@@ -10,7 +10,7 @@ func.func @missing_sizes(%tc : i32, %ts : i32) {
// expected-error at +1 {{'omp.tile' op there must be one tile size for each applyee}}
omp.tile <-(%canonloop)
- llvm.return
+ return
}
// -----
@@ -49,7 +49,7 @@ func.func @insufficient_sizes(%tc : i32, %ts : i32) {
// expected-error at +1 {{'omp.tile' op there must be one tile size for each applyee}}
omp.tile <-(%canonloop1, %canonloop2) sizes(%ts : i32)
- llvm.return
+ return
}
// -----
@@ -97,7 +97,7 @@ func.func @not_perfectly_nested(%tc : i32, %ts : i32) {
// expected-error at +1 {{'omp.tile' op tiled loop nest must be perfectly nested}}
omp.tile <-(%canonloop1, %canonloop2) sizes(%ts, %ts : i32, i32)
- llvm.return
+ return
}
// -----
@@ -115,5 +115,5 @@ func.func @non_nectangular(%tc : i32, %ts : i32) {
// expected-error at +1 {{'omp.tile' op tiled loop nest must be rectangular}}
omp.tile <-(%canonloop1, %canonloop2) sizes(%ts, %ts : i32, i32)
- llvm.return
+ return
}
diff --git a/mlir/test/Dialect/OpenMP/invalid.mlir b/mlir/test/Dialect/OpenMP/invalid.mlir
index bc508d66fbd5f..2c2f610dadc9f 100644
--- a/mlir/test/Dialect/OpenMP/invalid.mlir
+++ b/mlir/test/Dialect/OpenMP/invalid.mlir
@@ -531,7 +531,7 @@ func.func @omp_simd_firstprivate(%lb : index, %ub : index, %step : index) {
omp.yield
}
}
- llvm.return
+ return
}
// -----
diff --git a/mlir/test/Dialect/SPIRV/IR/availability.mlir b/mlir/test/Dialect/SPIRV/IR/availability.mlir
index 4ef242bdc5b16..4fa0c91a6fa6a 100644
--- a/mlir/test/Dialect/SPIRV/IR/availability.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/availability.mlir
@@ -290,7 +290,7 @@ func.func @set_mesh_outputs(%0 : i32, %1 : i32) -> () {
// CHECK: extensions: [ [SPV_EXT_mesh_shader] ]
// CHECK: capabilities: [ [MeshShadingEXT] ]
spirv.EXT.SetMeshOutputs %0, %1 : i32, i32
- spirv.Return
+ return
}
//===----------------------------------------------------------------------===//
@@ -304,7 +304,7 @@ func.func @constant_composite_replicate() -> () {
// CHECK: extensions: [ [SPV_EXT_replicated_composites] ]
// CHECK: capabilities: [ [ReplicatedCompositesEXT] ]
%0 = spirv.EXT.ConstantCompositeReplicate [1 : i32] : vector<2xi32>
- spirv.Return
+ return
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/test/Dialect/SPIRV/IR/bit-ops.mlir b/mlir/test/Dialect/SPIRV/IR/bit-ops.mlir
index 4bdac198a1e8f..b75af1a47b6fb 100644
--- a/mlir/test/Dialect/SPIRV/IR/bit-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/bit-ops.mlir
@@ -9,7 +9,7 @@
func.func @bitcount(%arg: i32) -> i32 {
// CHECK: spirv.BitCount {{%.*}} : i32
%0 = spirv.BitCount %arg : i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -21,7 +21,7 @@ func.func @bitcount(%arg: i32) -> i32 {
func.func @bit_field_insert_vec(%base: vector<3xi32>, %insert: vector<3xi32>, %offset: i32, %count: i16) -> vector<3xi32> {
// CHECK: {{%.*}} = spirv.BitFieldInsert {{%.*}}, {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi32>, i32, i16
%0 = spirv.BitFieldInsert %base, %insert, %offset, %count : vector<3xi32>, i32, i16
- spirv.ReturnValue %0 : vector<3xi32>
+ return %0 : vector<3xi32>
}
// -----
@@ -33,7 +33,7 @@ func.func @bit_field_insert_invalid_insert_type(%base: vector<3xi32>, %insert: v
// insert.
// expected-error @+1 {{type}}
%0 = "spirv.BitFieldInsert" (%base, %insert, %offset, %count) : (vector<3xi32>, vector<2xi32>, i32, i16) -> vector<3xi32>
- spirv.ReturnValue %0 : vector<3xi32>
+ return %0 : vector<3xi32>
}
// -----
@@ -45,7 +45,7 @@ func.func @bit_field_insert_invalid_insert_type(%base: vector<3xi32>, %insert: v
func.func @bit_field_s_extract_vec(%base: vector<3xi32>, %offset: i8, %count: i8) -> vector<3xi32> {
// CHECK: {{%.*}} = spirv.BitFieldSExtract {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi32>, i8, i8
%0 = spirv.BitFieldSExtract %base, %offset, %count : vector<3xi32>, i8, i8
- spirv.ReturnValue %0 : vector<3xi32>
+ return %0 : vector<3xi32>
}
//===----------------------------------------------------------------------===//
@@ -55,7 +55,7 @@ func.func @bit_field_s_extract_vec(%base: vector<3xi32>, %offset: i8, %count: i8
func.func @bit_field_u_extract_vec(%base: vector<3xi32>, %offset: i8, %count: i8) -> vector<3xi32> {
// CHECK: {{%.*}} = spirv.BitFieldUExtract {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi32>, i8, i8
%0 = spirv.BitFieldUExtract %base, %offset, %count : vector<3xi32>, i8, i8
- spirv.ReturnValue %0 : vector<3xi32>
+ return %0 : vector<3xi32>
}
// -----
@@ -63,7 +63,7 @@ func.func @bit_field_u_extract_vec(%base: vector<3xi32>, %offset: i8, %count: i8
func.func @bit_field_u_extract_invalid_result_type(%base: vector<3xi32>, %offset: i32, %count: i16) -> vector<4xi32> {
// expected-error @+1 {{failed to verify that all of {base, result} have same type}}
%0 = "spirv.BitFieldUExtract" (%base, %offset, %count) : (vector<3xi32>, i32, i16) -> vector<4xi32>
- spirv.ReturnValue %0 : vector<4xi32>
+ return %0 : vector<4xi32>
}
// -----
@@ -75,7 +75,7 @@ func.func @bit_field_u_extract_invalid_result_type(%base: vector<3xi32>, %offset
func.func @bitreverse(%arg: i32) -> i32 {
// CHECK: spirv.BitReverse {{%.*}} : i32
%0 = spirv.BitReverse %arg : i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -288,7 +288,7 @@ func.func @bitwise_and_float(%arg0: f16, %arg1: f16) -> f16 {
func.func @not(%arg: i32) -> i32 {
// CHECK: spirv.Not {{%.*}} : i32
%0 = spirv.Not %arg : i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -300,7 +300,7 @@ func.func @not(%arg: i32) -> i32 {
func.func @shift_left_logical(%arg0: i32, %arg1 : i16) -> i32 {
// CHECK: {{%.*}} = spirv.ShiftLeftLogical {{%.*}}, {{%.*}} : i32, i16
%0 = spirv.ShiftLeftLogical %arg0, %arg1: i32, i16
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -308,7 +308,7 @@ func.func @shift_left_logical(%arg0: i32, %arg1 : i16) -> i32 {
func.func @shift_left_logical_invalid_result_type(%arg0: i32, %arg1 : i16) -> i16 {
// expected-error @+1 {{op failed to verify that all of {operand1, result} have same type}}
%0 = "spirv.ShiftLeftLogical" (%arg0, %arg1) : (i32, i16) -> (i16)
- spirv.ReturnValue %0 : i16
+ return %0 : i16
}
// -----
@@ -320,7 +320,7 @@ func.func @shift_left_logical_invalid_result_type(%arg0: i32, %arg1 : i16) -> i1
func.func @shift_right_arithmetic(%arg0: vector<4xi32>, %arg1 : vector<4xi8>) -> vector<4xi32> {
// CHECK: {{%.*}} = spirv.ShiftRightArithmetic {{%.*}}, {{%.*}} : vector<4xi32>, vector<4xi8>
%0 = spirv.ShiftRightArithmetic %arg0, %arg1: vector<4xi32>, vector<4xi8>
- spirv.ReturnValue %0 : vector<4xi32>
+ return %0 : vector<4xi32>
}
// -----
@@ -332,5 +332,5 @@ func.func @shift_right_arithmetic(%arg0: vector<4xi32>, %arg1 : vector<4xi8>) ->
func.func @shift_right_logical(%arg0: vector<2xi32>, %arg1 : vector<2xi8>) -> vector<2xi32> {
// CHECK: {{%.*}} = spirv.ShiftRightLogical {{%.*}}, {{%.*}} : vector<2xi32>, vector<2xi8>
%0 = spirv.ShiftRightLogical %arg0, %arg1: vector<2xi32>, vector<2xi8>
- spirv.ReturnValue %0 : vector<2xi32>
+ return %0 : vector<2xi32>
}
diff --git a/mlir/test/Dialect/SPIRV/IR/cast-ops.mlir b/mlir/test/Dialect/SPIRV/IR/cast-ops.mlir
index 4480a1f3720f2..f8156d2c7511a 100644
--- a/mlir/test/Dialect/SPIRV/IR/cast-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/cast-ops.mlir
@@ -89,7 +89,7 @@ func.func @cast3(%arg0 : i64) {
func.func @convert_f_to_s_scalar(%arg0 : f32) -> i32 {
// CHECK: {{%.*}} = spirv.ConvertFToS {{%.*}} : f32 to i32
%0 = spirv.ConvertFToS %arg0 : f32 to i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -97,7 +97,7 @@ func.func @convert_f_to_s_scalar(%arg0 : f32) -> i32 {
func.func @convert_f64_to_s32_scalar(%arg0 : f64) -> i32 {
// CHECK: {{%.*}} = spirv.ConvertFToS {{%.*}} : f64 to i32
%0 = spirv.ConvertFToS %arg0 : f64 to i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -105,7 +105,7 @@ func.func @convert_f64_to_s32_scalar(%arg0 : f64) -> i32 {
func.func @convert_f_to_s_vector(%arg0 : vector<3xf32>) -> vector<3xi32> {
// CHECK: {{%.*}} = spirv.ConvertFToS {{%.*}} : vector<3xf32> to vector<3xi32>
%0 = spirv.ConvertFToS %arg0 : vector<3xf32> to vector<3xi32>
- spirv.ReturnValue %0 : vector<3xi32>
+ return %0 : vector<3xi32>
}
// -----
@@ -113,7 +113,7 @@ func.func @convert_f_to_s_vector(%arg0 : vector<3xf32>) -> vector<3xi32> {
func.func @convert_bf16_to_s32_scalar(%arg0 : bf16) -> i32 {
// CHECK: {{%.*}} = spirv.ConvertFToS {{%.*}} : bf16 to i32
%0 = spirv.ConvertFToS %arg0 : bf16 to i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -125,7 +125,7 @@ func.func @convert_bf16_to_s32_scalar(%arg0 : bf16) -> i32 {
func.func @convert_f_to_u_scalar(%arg0 : f32) -> i32 {
// CHECK: {{%.*}} = spirv.ConvertFToU {{%.*}} : f32 to i32
%0 = spirv.ConvertFToU %arg0 : f32 to i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -133,7 +133,7 @@ func.func @convert_f_to_u_scalar(%arg0 : f32) -> i32 {
func.func @convert_f64_to_u32_scalar(%arg0 : f64) -> i32 {
// CHECK: {{%.*}} = spirv.ConvertFToU {{%.*}} : f64 to i32
%0 = spirv.ConvertFToU %arg0 : f64 to i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -141,7 +141,7 @@ func.func @convert_f64_to_u32_scalar(%arg0 : f64) -> i32 {
func.func @convert_f_to_u_vector(%arg0 : vector<3xf32>) -> vector<3xi32> {
// CHECK: {{%.*}} = spirv.ConvertFToU {{%.*}} : vector<3xf32> to vector<3xi32>
%0 = spirv.ConvertFToU %arg0 : vector<3xf32> to vector<3xi32>
- spirv.ReturnValue %0 : vector<3xi32>
+ return %0 : vector<3xi32>
}
// -----
@@ -149,7 +149,7 @@ func.func @convert_f_to_u_vector(%arg0 : vector<3xf32>) -> vector<3xi32> {
func.func @convert_f_to_u.coopmatrix(%arg0 : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixB>) {
// CHECK: {{%.*}} = spirv.ConvertFToU {{%.*}} : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixB> to !spirv.coopmatrix<8x16xi32, Subgroup, MatrixB>
%0 = spirv.ConvertFToU %arg0 : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixB> to !spirv.coopmatrix<8x16xi32, Subgroup, MatrixB>
- spirv.Return
+ return
}
// -----
@@ -157,7 +157,7 @@ func.func @convert_f_to_u.coopmatrix(%arg0 : !spirv.coopmatrix<8x16xf32, Subgrou
func.func @convert_bf16_to_u32_scalar(%arg0 : bf16) -> i32 {
// CHECK: {{%.*}} = spirv.ConvertFToU {{%.*}} : bf16 to i32
%0 = spirv.ConvertFToU %arg0 : bf16 to i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// -----
@@ -169,7 +169,7 @@ func.func @convert_bf16_to_u32_scalar(%arg0 : bf16) -> i32 {
func.func @convert_s_to_f_scalar(%arg0 : i32) -> f32 {
// CHECK: {{%.*}} = spirv.ConvertSToF {{%.*}} : i32 to f32
%0 = spirv.ConvertSToF %arg0 : i32 to f32
- spirv.ReturnValue %0 : f32
+ return %0 : f32
}
// -----
@@ -177,7 +177,7 @@ func.func @convert_s_to_f_scalar(%arg0 : i32) -> f32 {
func.func @convert_s64_to_f32_scalar(%arg0 : i64) -> f32 {
// CHECK: {{%.*}} = spirv.ConvertSToF {{%.*}} : i64 to f32
%0 = spirv.ConvertSToF %arg0 : i64 to f32
- spirv.ReturnValue %0 : f32
+ return %0 : f32
}
// -----
@@ -185,7 +185,7 @@ func.func @convert_s64_to_f32_scalar(%arg0 : i64) -> f32 {
func.func @convert_s_to_f_vector(%arg0 : vector<3xi32>) -> vector<3xf32> {
// CHECK: {{%.*}} = spirv.ConvertSToF {{%.*}} : vector<3xi32> to vector<3xf32>
%0 = spirv.ConvertSToF %arg0 : vector<3xi32> to vector<3xf32>
- spirv.ReturnValue %0 : vector<3xf32>
+ return %0 : vector<3xf32>
}
// -----
@@ -193,7 +193,7 @@ func.func @convert_s_to_f_vector(%arg0 : vector<3xi32>) -> vector<3xf32> {
func.func @convert_s32_to_bf16_scalar(%arg0 : i32) -> bf16 {
// CHECK: {{%.*}} = spirv.ConvertSToF {{%.*}} : i32 to bf16
%0 = spirv.ConvertSToF %arg0 : i32 to bf16
- spirv.ReturnValue %0 : bf16
+ return %0 : bf16
}
// -----
@@ -205,7 +205,7 @@ func.func @convert_s32_to_bf16_scalar(%arg0 : i32) -> bf16 {
func.func @convert_u_to_f_scalar(%arg0 : i32) -> f32 {
// CHECK: {{%.*}} = spirv.ConvertUToF {{%.*}} : i32 to f32
%0 = spirv.ConvertUToF %arg0 : i32 to f32
- spirv.ReturnValue %0 : f32
+ return %0 : f32
}
// -----
@@ -213,7 +213,7 @@ func.func @convert_u_to_f_scalar(%arg0 : i32) -> f32 {
func.func @convert_u64_to_f32_scalar(%arg0 : i64) -> f32 {
// CHECK: {{%.*}} = spirv.ConvertUToF {{%.*}} : i64 to f32
%0 = spirv.ConvertUToF %arg0 : i64 to f32
- spirv.ReturnValue %0 : f32
+ return %0 : f32
}
// -----
@@ -221,7 +221,7 @@ func.func @convert_u64_to_f32_scalar(%arg0 : i64) -> f32 {
func.func @convert_u_to_f_vector(%arg0 : vector<3xi32>) -> vector<3xf32> {
// CHECK: {{%.*}} = spirv.ConvertUToF {{%.*}} : vector<3xi32> to vector<3xf32>
%0 = spirv.ConvertUToF %arg0 : vector<3xi32> to vector<3xf32>
- spirv.ReturnValue %0 : vector<3xf32>
+ return %0 : vector<3xf32>
}
// -----
@@ -229,7 +229,7 @@ func.func @convert_u_to_f_vector(%arg0 : vector<3xi32>) -> vector<3xf32> {
func.func @convert_u32_to_bf16_scalar(%arg0 : i32) -> bf16 {
// CHECK: {{%.*}} = spirv.ConvertUToF {{%.*}} : i32 to bf16
%0 = spirv.ConvertUToF %arg0 : i32 to bf16
- spirv.ReturnValue %0 : bf16
+ return %0 : bf16
}
// -----
@@ -241,7 +241,7 @@ func.func @convert_u32_to_bf16_scalar(%arg0 : i32) -> bf16 {
func.func @f_convert_scalar(%arg0 : f32) -> f64 {
// CHECK: {{%.*}} = spirv.FConvert {{%.*}} : f32 to f64
%0 = spirv.FConvert %arg0 : f32 to f64
- spirv.ReturnValue %0 : f64
+ return %0 : f64
}
// -----
@@ -249,7 +249,7 @@ func.func @f_convert_scalar(%arg0 : f32) -> f64 {
func.func @f_convert_vector(%arg0 : vector<3xf32>) -> vector<3xf64> {
// CHECK: {{%.*}} = spirv.FConvert {{%.*}} : vector<3xf32> to vector<3xf64>
%0 = spirv.FConvert %arg0 : vector<3xf32> to vector<3xf64>
- spirv.ReturnValue %0 : vector<3xf64>
+ return %0 : vector<3xf64>
}
// -----
@@ -257,7 +257,7 @@ func.func @f_convert_vector(%arg0 : vector<3xf32>) -> vector<3xf64> {
func.func @f_convert_coop_matrix(%arg0 : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixA>) {
// CHECK: {{%.*}} = spirv.FConvert {{%.*}} : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixA> to !spirv.coopmatrix<8x16xf64, Subgroup, MatrixA>
%0 = spirv.FConvert %arg0 : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixA> to !spirv.coopmatrix<8x16xf64, Subgroup, MatrixA>
- spirv.Return
+ return
}
// -----
@@ -265,7 +265,7 @@ func.func @f_convert_coop_matrix(%arg0 : !spirv.coopmatrix<8x16xf32, Subgroup, M
func.func @f_convert_vector(%arg0 : f32) -> f32 {
// expected-error @+1 {{expected the different bit widths for operand type and result type, but provided 'f32' and 'f32'}}
%0 = spirv.FConvert %arg0 : f32 to f32
- spirv.ReturnValue %0 : f32
+ return %0 : f32
}
// -----
@@ -273,7 +273,7 @@ func.func @f_convert_vector(%arg0 : f32) -> f32 {
func.func @f_convert_bf16_to_f32_scalar(%arg0 : bf16) -> f32 {
// CHECK: {{%.*}} = spirv.FConvert {{%.*}} : bf16 to f32
%0 = spirv.FConvert %arg0 : bf16 to f32
- spirv.ReturnValue %0 : f32
+ return %0 : f32
}
// -----
@@ -281,7 +281,7 @@ func.func @f_convert_bf16_to_f32_scalar(%arg0 : bf16) -> f32 {
func.func @f_convert_f32_to_bf16_vector(%arg0 : vector<3xf32>) -> vector<3xbf16> {
// CHECK: {{%.*}} = spirv.FConvert {{%.*}} : vector<3xf32> to vector<3xbf16>
%0 = spirv.FConvert %arg0 : vector<3xf32> to vector<3xbf16>
- spirv.ReturnValue %0 : vector<3xbf16>
+ return %0 : vector<3xbf16>
}
// -----
@@ -289,7 +289,7 @@ func.func @f_convert_f32_to_bf16_vector(%arg0 : vector<3xf32>) -> vector<3xbf16>
func.func @f_convert_f32_to_bf16_coop_matrix(%arg0 : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixA>) -> !spirv.coopmatrix<8x16xbf16, Subgroup, MatrixA> {
// CHECK: {{%.*}} = spirv.FConvert {{%.*}} : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixA> to !spirv.coopmatrix<8x16xbf16, Subgroup, MatrixA>
%0 = spirv.FConvert %arg0 : !spirv.coopmatrix<8x16xf32, Subgroup, MatrixA> to !spirv.coopmatrix<8x16xbf16, Subgroup, MatrixA>
- spirv.ReturnValue %0 : !spirv.coopmatrix<8x16xbf16, Subgroup, MatrixA>
+ return %0 : !spirv.coopmatrix<8x16xbf16, Subgroup, MatrixA>
}
// -----
@@ -301,7 +301,7 @@ func.func @f_convert_f32_to_bf16_coop_matrix(%arg0 : !spirv.coopmatrix<8x16xf32,
func.func @s_convert_scalar(%arg0 : i32) -> i64 {
// CHECK: {{%.*}} = spirv.SConvert {{%.*}} : i32 to i64
%0 = spirv.SConvert %arg0 : i32 to i64
- spirv.ReturnValue %0 : i64
+ return %0 : i64
}
// -----
@@ -313,7 +313,7 @@ func.func @s_convert_scalar(%arg0 : i32) -> i64 {
func.func @u_convert_scalar(%arg0 : i32) -> i64 {
// CHECK: {{%.*}} = spirv.UConvert {{%.*}} : i32 to i64
%0 = spirv.UConvert %arg0 : i32 to i64
- spirv.ReturnValue %0 : i64
+ return %0 : i64
}
// -----
diff --git a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
index 81dce9822db48..a1a0c487be209 100644
--- a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
@@ -8,7 +8,7 @@ func.func @branch() -> () {
// CHECK: spirv.Branch ^bb1
spirv.Branch ^next
^next:
- spirv.Return
+ return
}
// -----
@@ -18,7 +18,7 @@ func.func @branch_argument() -> () {
// CHECK: spirv.Branch ^bb1(%{{.*}}, %{{.*}} : i32, i32)
spirv.Branch ^next(%zero, %zero: i32, i32)
^next(%arg0: i32, %arg1: i32):
- spirv.Return
+ return
}
// -----
@@ -35,9 +35,9 @@ func.func @wrong_accessor_count() -> () {
// expected-error @+1 {{requires 1 successor but found 2}}
"spirv.Branch"()[^one, ^two] : () -> ()
^one:
- spirv.Return
+ return
^two:
- spirv.Return
+ return
}
// -----
@@ -52,10 +52,10 @@ func.func @cond_branch() -> () {
spirv.BranchConditional %true, ^one, ^two
// CHECK: ^bb1
^one:
- spirv.Return
+ return
// CHECK: ^bb2
^two:
- spirv.Return
+ return
}
// -----
@@ -69,11 +69,11 @@ func.func @cond_branch_argument() -> () {
// CHECK: spirv.BranchConditional %{{.*}}, ^bb3, ^bb4(%{{.*}}, %{{.*}} : i32, i32)
spirv.BranchConditional %true, ^true2, ^false2(%zero, %zero: i32, i32)
^false1:
- spirv.Return
+ return
^true2:
- spirv.Return
+ return
^false2(%arg3: i32, %arg4: i32):
- spirv.Return
+ return
}
// -----
@@ -83,9 +83,9 @@ func.func @cond_branch_with_weights() -> () {
// CHECK: spirv.BranchConditional %{{.*}} [5, 10]
spirv.BranchConditional %true [5, 10], ^one, ^two
^one:
- spirv.Return
+ return
^two:
- spirv.Return
+ return
}
// -----
@@ -94,9 +94,9 @@ func.func @missing_condition() -> () {
// expected-error @+1 {{expected SSA operand}}
spirv.BranchConditional ^one, ^two
^one:
- spirv.Return
+ return
^two:
- spirv.Return
+ return
}
// -----
@@ -107,9 +107,9 @@ func.func @wrong_condition_type() -> () {
// expected-error @+1 {{use of value '%zero' expects different type than prior uses: 'i1' vs 'i32'}}
spirv.BranchConditional %zero, ^one, ^two
^one:
- spirv.Return
+ return
^two:
- spirv.Return
+ return
}
// -----
@@ -119,9 +119,9 @@ func.func @wrong_accessor_count() -> () {
// expected-error @+1 {{requires 2 successors but found 1}}
"spirv.BranchConditional"(%true)[^one] {operandSegmentSizes = array<i32: 1, 0, 0>} : (i1) -> ()
^one:
- spirv.Return
+ return
^two:
- spirv.Return
+ return
}
// -----
@@ -132,9 +132,9 @@ func.func @wrong_number_of_weights() -> () {
"spirv.BranchConditional"(%true)[^one, ^two] {branch_weights = [1 : i32, 2 : i32, 3 : i32],
operandSegmentSizes = array<i32: 1, 0, 0>} : (i1) -> ()
^one:
- spirv.Return
+ return
^two:
- spirv.Return
+ return
}
// -----
@@ -144,9 +144,9 @@ func.func @weights_cannot_both_be_zero() -> () {
// expected-error @+1 {{branch weights cannot both be zero}}
spirv.BranchConditional %true [0, 0], ^one, ^two
^one:
- spirv.Return
+ return
^two:
- spirv.Return
+ return
}
// -----
@@ -195,7 +195,7 @@ spirv.func @callee() "None" {
func.func @caller() {
// CHECK: spirv.FunctionCall
spirv.FunctionCall @callee() : () -> ()
- spirv.Return
+ return
}
// -----
@@ -557,7 +557,7 @@ func.func @loop_yield_result_count_mismatch(%count : i32) -> () {
//===----------------------------------------------------------------------===//
func.func @merge() -> () {
- // expected-error @+1 {{expects parent op to be one of 'spirv.mlir.selection, spirv.mlir.loop'}}
+ // expected-error @+1 {{op expects parent op to be one of 'spirv.mlir.selection, spirv.mlir.loop'}}
spirv.mlir.merge
}
@@ -572,7 +572,7 @@ func.func @only_allowed_in_last_block(%cond : i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -587,7 +587,7 @@ func.func @last_block_no_terminator(%cond : i1) -> () {
spirv.mlir.merge
^merge:
}
- spirv.Return
+ return
}
// -----
@@ -625,7 +625,7 @@ func.func @in_selection(%cond : i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// CHECK-LABEL: func @in_loop
@@ -642,13 +642,12 @@ func.func @in_loop(%cond : i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// CHECK-LABEL: in_other_func_like_op
func.func @in_other_func_like_op() {
- // CHECK: spirv.Return
- spirv.Return
+ return
}
// -----
@@ -694,8 +693,7 @@ spirv.module Logical GLSL450 {
func.func @ret_val() -> (i32) {
%0 = spirv.Constant 42 : i32
- // CHECK: spirv.ReturnValue %{{.*}} : i32
- spirv.ReturnValue %0 : i32
+ return %0 : i32
}
// CHECK-LABEL: func @in_selection
@@ -710,7 +708,7 @@ func.func @in_selection(%cond : i1) -> (i32) {
spirv.mlir.merge
}
%one = spirv.Constant 1 : i32
- spirv.ReturnValue %one : i32
+ return %one : i32
}
// CHECK-LABEL: func @in_loop
@@ -729,13 +727,12 @@ func.func @in_loop(%cond : i1) -> (i32) {
spirv.mlir.merge
}
%one = spirv.Constant 1 : i32
- spirv.ReturnValue %one : i32
+ return %one : i32
}
// CHECK-LABEL: in_other_func_like_op
func.func @in_other_func_like_op(%arg: i32) -> i32 {
- // CHECK: spirv.ReturnValue
- spirv.ReturnValue %arg: i32
+ return %arg: i32
}
// -----
@@ -812,7 +809,7 @@ func.func @selection(%cond: i1) -> () {
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -846,7 +843,7 @@ func.func @selection(%cond: i1) -> () {
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -893,7 +890,7 @@ func.func @selection_switch(%selector: i32) -> () {
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -973,7 +970,7 @@ func.func @selection_yield(%cond: i1) -> () {
// CHECK: spirv.Store "Function" {{%.*}}, {{%.*}}#1 : i32
spirv.Store "Function" %var2, %yield#1 : i32
- spirv.Return
+ return
}
// -----
@@ -999,7 +996,7 @@ func.func @selection_yield_result_type_mismatch(%cond: i1) -> () {
spirv.mlir.merge %merged_1_2, %merged_3_4 : i32, i32
}
- spirv.Return
+ return
}
// -----
@@ -1023,7 +1020,7 @@ func.func @selection_yield_result_count_mismatch(%cond: i1) -> () {
spirv.mlir.merge %merged_1_2 : i32
}
- spirv.Return
+ return
}
// -----
@@ -1035,7 +1032,7 @@ func.func @selection_yield_result_count_mismatch(%cond: i1) -> () {
// CHECK-LABEL: func @unreachable_no_pred
func.func @unreachable_no_pred() {
- spirv.Return
+ return
^next:
// CHECK: spirv.Unreachable
@@ -1044,7 +1041,7 @@ func.func @unreachable_no_pred() {
// CHECK-LABEL: func @unreachable_with_pred
func.func @unreachable_with_pred() {
- spirv.Return
+ return
^parent:
spirv.Branch ^unreachable
@@ -1104,7 +1101,7 @@ func.func @switch(%selector: i32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
func.func @switch_only_default(%selector: i32) -> () {
@@ -1117,7 +1114,7 @@ func.func @switch_only_default(%selector: i32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
func.func @switch_operands(%selector : i32, %operand : i32) {
@@ -1140,7 +1137,7 @@ func.func @switch_operands(%selector : i32, %operand : i32) {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
// -----
@@ -1154,7 +1151,7 @@ func.func @switch_float_selector(%selector: f32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
// -----
@@ -1172,7 +1169,7 @@ func.func @switch_float_selector(%selector: i32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
// -----
@@ -1186,7 +1183,7 @@ func.func @switch_missing_default(%selector: i32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
// -----
@@ -1200,7 +1197,7 @@ func.func @switch_default_no_target(%selector: i32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
// -----
@@ -1218,7 +1215,7 @@ func.func @switch_case_no_target(%selector: i32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
// -----
@@ -1237,6 +1234,6 @@ func.func @switch_missing_operand_type(%selector: i32) -> () {
spirv.Branch ^merge
^merge:
- spirv.Return
+ return
}
diff --git a/mlir/test/Dialect/SPIRV/IR/image-ops.mlir b/mlir/test/Dialect/SPIRV/IR/image-ops.mlir
index 12b5f2ce62a68..934b16309d428 100644
--- a/mlir/test/Dialect/SPIRV/IR/image-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/image-ops.mlir
@@ -7,7 +7,7 @@
func.func @image_dref_gather(%arg0 : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// CHECK: spirv.ImageDrefGather {{.*}}, {{.*}}, {{.*}} : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<4xi32>
%0 = spirv.ImageDrefGather %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<4xi32>
- spirv.Return
+ return
}
// -----
@@ -15,7 +15,7 @@ func.func @image_dref_gather(%arg0 : !spirv.sampled_image<!spirv.image<i32, Dim2
func.func @image_dref_gather_with_single_imageoperands(%arg0 : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// CHECK: spirv.ImageDrefGather {{.*}}, {{.*}}, {{.*}} ["NonPrivateTexel"] : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<4xi32>
%0 = spirv.ImageDrefGather %arg0, %arg1, %arg2 ["NonPrivateTexel"] : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<4xi32>
- spirv.Return
+ return
}
// -----
@@ -23,7 +23,7 @@ func.func @image_dref_gather_with_single_imageoperands(%arg0 : !spirv.sampled_im
func.func @image_dref_gather_with_mismatch_imageoperands(%arg0 : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{the Image Operands should encode what operands follow, as per Image Operands}}
%0 = spirv.ImageDrefGather %arg0, %arg1, %arg2, %arg2, %arg2 : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32, f32, f32 -> vector<4xi32>
- spirv.Return
+ return
}
// -----
@@ -31,7 +31,7 @@ func.func @image_dref_gather_with_mismatch_imageoperands(%arg0 : !spirv.sampled_
func.func @image_dref_gather_error_result_type(%arg0 : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{op result #0 must be vector of 8/16/32/64-bit integer values of length 4 of ranks 1 or vector of 16/32/64-bit float values of length 4 of ranks 1, but got 'vector<3xi32>'}}
%0 = spirv.ImageDrefGather %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<i32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<3xi32>
- spirv.Return
+ return
}
// -----
@@ -39,7 +39,7 @@ func.func @image_dref_gather_error_result_type(%arg0 : !spirv.sampled_image<!spi
func.func @image_dref_gather_error_same_type(%arg0 : !spirv.sampled_image<!spirv.image<i32, Rect, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{the result component type must match the image sampled type}}
%0 = spirv.ImageDrefGather %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<i32, Rect, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -47,7 +47,7 @@ func.func @image_dref_gather_error_same_type(%arg0 : !spirv.sampled_image<!spirv
func.func @image_dref_gather_error_dim(%arg0 : !spirv.sampled_image<!spirv.image<i32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{the Dim operand of the underlying image must be Dim2D or Cube or Rect}}
%0 = spirv.ImageDrefGather %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<i32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<4xi32>
- spirv.Return
+ return
}
// -----
@@ -55,7 +55,7 @@ func.func @image_dref_gather_error_dim(%arg0 : !spirv.sampled_image<!spirv.image
func.func @image_dref_gather_error_ms(%arg0 : !spirv.sampled_image<!spirv.image<i32, Cube, NoDepth, NonArrayed, MultiSampled, NoSampler, Unknown>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{the MS operand of the underlying image type must be SingleSampled}}
%0 = spirv.ImageDrefGather %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<i32, Cube, NoDepth, NonArrayed, MultiSampled, NoSampler, Unknown>>, vector<4xf32>, f32 -> vector<4xi32>
- spirv.Return
+ return
}
// -----
@@ -79,7 +79,7 @@ func.func @image(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth,
func.func @image_query_size(%arg0 : !spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>) -> () {
// CHECK: {{%.*}} = spirv.ImageQuerySize %arg0 : !spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown> -> i32
%0 = spirv.ImageQuerySize %arg0 : !spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown> -> i32
- spirv.Return
+ return
}
// -----
@@ -87,7 +87,7 @@ func.func @image_query_size(%arg0 : !spirv.image<f32, Dim1D, NoDepth, NonArrayed
func.func @image_query_size_error_dim(%arg0 : !spirv.image<f32, SubpassData, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>) -> () {
// expected-error @+1 {{the Dim operand of the image type must be 1D, 2D, 3D, Buffer, Cube, or Rect}}
%0 = spirv.ImageQuerySize %arg0 : !spirv.image<f32, SubpassData, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown> -> i32
- spirv.Return
+ return
}
// -----
@@ -95,7 +95,7 @@ func.func @image_query_size_error_dim(%arg0 : !spirv.image<f32, SubpassData, NoD
func.func @image_query_size_error_dim_sample(%arg0 : !spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Unknown>) -> () {
// expected-error @+1 {{if Dim is 1D, 2D, 3D, or Cube, it must also have either an MS of 1 or a Sampled of 0 or 2}}
%0 = spirv.ImageQuerySize %arg0 : !spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Unknown> -> i32
- spirv.Return
+ return
}
// -----
@@ -103,7 +103,7 @@ func.func @image_query_size_error_dim_sample(%arg0 : !spirv.image<f32, Dim1D, No
func.func @image_query_size_error_result1(%arg0 : !spirv.image<f32, Dim3D, NoDepth, Arrayed, SingleSampled, NoSampler, Unknown>) -> () {
// expected-error @+1 {{expected the result to have 4 component(s), but found 3 component(s)}}
%0 = spirv.ImageQuerySize %arg0 : !spirv.image<f32, Dim3D, NoDepth, Arrayed, SingleSampled, NoSampler, Unknown> -> vector<3xi32>
- spirv.Return
+ return
}
// -----
@@ -111,7 +111,7 @@ func.func @image_query_size_error_result1(%arg0 : !spirv.image<f32, Dim3D, NoDep
func.func @image_query_size_error_result2(%arg0 : !spirv.image<f32, Buffer, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>) -> () {
// expected-error @+1 {{expected the result to have 1 component(s), but found 2 component(s)}}
%0 = spirv.ImageQuerySize %arg0 : !spirv.image<f32, Buffer, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown> -> vector<2xi32>
- spirv.Return
+ return
}
// -----
@@ -123,7 +123,7 @@ func.func @image_query_size_error_result2(%arg0 : !spirv.image<f32, Buffer, NoDe
func.func @image_read(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// CHECK: {{%.*}} = spirv.ImageRead {{%.*}}, {{%.*}} : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba8>, vector<2xsi32> -> vector<4xf32>
%0 = spirv.ImageRead %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba8>, vector<2xsi32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -131,7 +131,7 @@ func.func @image_read(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, Singl
func.func @image_read_type_mismatch(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// expected-error @+1 {{op failed to verify that the result component type must match the image sampled type}}
%0 = spirv.ImageRead %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba8>, vector<2xsi32> -> vector<4xf16>
- spirv.Return
+ return
}
// -----
@@ -139,7 +139,7 @@ func.func @image_read_type_mismatch(%arg0: !spirv.image<f32, Dim2D, NoDepth, Non
func.func @image_read_need_sampler(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// expected-error @+1 {{op failed to verify that the sampled operand of the underlying image must be SamplerUnknown or NoSampler}}
%0 = spirv.ImageRead %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, vector<2xsi32> -> vector<4xf16>
- spirv.Return
+ return
}
// -----
@@ -151,7 +151,7 @@ func.func @image_read_need_sampler(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonA
func.func @image_write(%arg0 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, %arg1 : vector<2xsi32>, %arg2 : vector<4xf32>) -> () {
// CHECK: spirv.ImageWrite {{%.*}}, {{%.*}}, {{%.*}} : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, vector<2xsi32>, vector<4xf32>
spirv.ImageWrite %arg0, %arg1, %arg2 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, vector<2xsi32>, vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -159,7 +159,7 @@ func.func @image_write(%arg0 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, Sin
func.func @image_write_scalar_texel(%arg0 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, %arg1 : vector<2xsi32>, %arg2 : f32) -> () {
// CHECK: spirv.ImageWrite {{%.*}}, {{%.*}}, {{%.*}} : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, vector<2xsi32>, f32
spirv.ImageWrite %arg0, %arg1, %arg2 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, vector<2xsi32>, f32
- spirv.Return
+ return
}
// -----
@@ -167,7 +167,7 @@ func.func @image_write_scalar_texel(%arg0 : !spirv.image<f32, Dim2D, NoDepth, No
func.func @image_write_need_sampler(%arg0 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba16>, %arg1 : vector<2xsi32>, %arg2 : vector<4xf32>) -> () {
// expected-error @+1 {{the sampled operand of the underlying image must be SamplerUnknown or NoSampler}}
spirv.ImageWrite %arg0, %arg1, %arg2 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba16>, vector<2xsi32>, vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -175,7 +175,7 @@ func.func @image_write_need_sampler(%arg0 : !spirv.image<f32, Dim2D, NoDepth, No
func.func @image_write_subpass_data(%arg0 : !spirv.image<f32, SubpassData, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, %arg1 : vector<2xsi32>, %arg2 : vector<4xf32>) -> () {
// expected-error @+1 {{the Dim operand of the underlying image must not be SubpassData}}
spirv.ImageWrite %arg0, %arg1, %arg2 : !spirv.image<f32, SubpassData, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, vector<2xsi32>, vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -183,7 +183,7 @@ func.func @image_write_subpass_data(%arg0 : !spirv.image<f32, SubpassData, NoDep
func.func @image_write_texel_type_mismatch(%arg0 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, %arg1 : vector<2xsi32>, %arg2 : vector<4xi32>) -> () {
// expected-error @+1 {{the texel component type must match the image sampled type}}
spirv.ImageWrite %arg0, %arg1, %arg2 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba16>, vector<2xsi32>, vector<4xi32>
- spirv.Return
+ return
}
// -----
@@ -198,7 +198,7 @@ func.func @image_write_texel_type_mismatch(%arg0 : !spirv.image<f32, Dim2D, NoDe
func.func @sample_explicit_lod(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// CHECK: {{%.*}} = spirv.ImageSampleExplicitLod {{%.*}}, {{%.*}} ["Lod"], {{%.*}} : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32 -> vector<4xf32>
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Lod"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -206,7 +206,7 @@ func.func @sample_explicit_lod(%arg0 : !spirv.sampled_image<!spirv.image<f32, Di
func.func @sample_explicit_lod_multi_sampled(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, MultiSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{the MS operand of the underlying image type must be SingleSampled}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Lod"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, MultiSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -214,7 +214,7 @@ func.func @sample_explicit_lod_multi_sampled(%arg0 : !spirv.sampled_image<!spirv
func.func @sample_explicit_lod_wrong_result_type(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{the result component type must match the image sampled type}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Lod"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32 -> vector<4xsi32>
- spirv.Return
+ return
}
// -----
@@ -222,7 +222,7 @@ func.func @sample_explicit_lod_wrong_result_type(%arg0 : !spirv.sampled_image<!s
func.func @sample_explicit_lod_no_lod(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{either Lod or Grad image operands must be present}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Bias"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -237,7 +237,7 @@ func.func @sample_explicit_lod_no_lod(%arg0 : !spirv.sampled_image<!spirv.image<
func.func @sample_implicit_lod(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>) -> () {
// CHECK: {{%.*}} = spirv.ImageSampleImplicitLod {{%.*}}, {{%.*}} : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32> -> vector<4xf32>
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -245,7 +245,7 @@ func.func @sample_implicit_lod(%arg0 : !spirv.sampled_image<!spirv.image<f32, Di
func.func @sample_implicit_lod_multi_sampled(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, MultiSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>) -> () {
// expected-error @+1 {{the MS operand of the underlying image type must be SingleSampled}}
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, MultiSampled, NeedSampler, Rgba8>>, vector<2xf32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -253,7 +253,7 @@ func.func @sample_implicit_lod_multi_sampled(%arg0 : !spirv.sampled_image<!spirv
func.func @sample_implicit_lod_wrong_result(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>) -> () {
// expected-error @+1 {{the result component type must match the image sampled type}}
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32> -> vector<4xi32>
- spirv.Return
+ return
}
// -----
@@ -268,7 +268,7 @@ func.func @sample_implicit_lod_wrong_result(%arg0 : !spirv.sampled_image<!spirv.
func.func @sample_implicit_proj_dref(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, IsDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// CHECK: {{%.*}} = spirv.ImageSampleProjDrefImplicitLod {{%.*}}, {{%.*}}, {{%.*}} : !spirv.sampled_image<!spirv.image<f32, Dim2D, IsDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<4xf32>, f32 -> f32
%0 = spirv.ImageSampleProjDrefImplicitLod %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, IsDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<4xf32>, f32 -> f32
- spirv.Return
+ return
}
// -----
@@ -276,7 +276,7 @@ func.func @sample_implicit_proj_dref(%arg0 : !spirv.sampled_image<!spirv.image<f
func.func @sample_implicit_proj_dref_multi_sampled(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, IsDepth, NonArrayed, MultiSampled, NeedSampler, Rgba8>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{the MS operand of the underlying image type must be SingleSampled}}
%0 = spirv.ImageSampleProjDrefImplicitLod %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, IsDepth, NonArrayed, MultiSampled, NeedSampler, Rgba8>>, vector<4xf32>, f32 -> f32
- spirv.Return
+ return
}
// -----
@@ -284,7 +284,7 @@ func.func @sample_implicit_proj_dref_multi_sampled(%arg0 : !spirv.sampled_image<
func.func @sample_implicit_proj_dref(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, IsDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<4xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{type of 'result' matches image type of 'sampled_image'}}
%0 = spirv.ImageSampleProjDrefImplicitLod %arg0, %arg1, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, IsDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<4xf32>, f32 -> i32
- spirv.Return
+ return
}
// -----
@@ -296,7 +296,7 @@ func.func @sample_implicit_proj_dref(%arg0 : !spirv.sampled_image<!spirv.image<f
func.func @image_fetch(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// CHECK: {{%.*}} = spirv.ImageFetch {{%.*}}, {{%.*}} : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, vector<2xsi32> -> vector<4xf32>
%0 = spirv.ImageFetch %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, vector<2xsi32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -304,7 +304,7 @@ func.func @image_fetch(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, Sing
func.func @image_fetch_dim_cube(%arg0: !spirv.image<f32, Cube, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// expected-error @+1 {{op failed to verify that the Dim operand of the underlying image must not be Cube}}
%0 = spirv.ImageFetch %arg0, %arg1 : !spirv.image<f32, Cube, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, vector<2xsi32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -312,7 +312,7 @@ func.func @image_fetch_dim_cube(%arg0: !spirv.image<f32, Cube, NoDepth, NonArray
func.func @image_fetch_no_sampler(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// expected-error @+1 {{op failed to verify that the sampled operand of the underlying image must be NeedSampler}}
%0 = spirv.ImageFetch %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NoSampler, Rgba8>, vector<2xsi32> -> vector<4xf16>
- spirv.Return
+ return
}
// -----
@@ -320,7 +320,7 @@ func.func @image_fetch_no_sampler(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonAr
func.func @image_fetch_type_mismatch(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// expected-error @+1 {{op failed to verify that the result component type must match the image sampled type}}
%0 = spirv.ImageFetch %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, vector<2xsi32> -> vector<4xf16>
- spirv.Return
+ return
}
// -----
@@ -328,7 +328,7 @@ func.func @image_fetch_type_mismatch(%arg0: !spirv.image<f32, Dim2D, NoDepth, No
func.func @image_fetch_2d_result(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, %arg1: vector<2xsi32>) -> () {
// expected-error @+1 {{op result #0 must be vector of 16/32/64-bit float values of length 4 of ranks 1 or vector of 8/16/32/64-bit integer values of length 4 of ranks 1, but got 'vector<2xf32>'}}
%0 = spirv.ImageFetch %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, vector<2xsi32> -> vector<2xf32>
- spirv.Return
+ return
}
// -----
@@ -336,7 +336,7 @@ func.func @image_fetch_2d_result(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArr
func.func @image_fetch_float_coords(%arg0: !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, %arg1: vector<2xf32>) -> () {
// expected-error @+1 {{op operand #1 must be 8/16/32/64-bit integer or fixed-length vector of 8/16/32/64-bit integer values of length 2/3/4/8/16 of ranks 1, but got 'vector<2xf32>'}}
%0 = spirv.ImageFetch %arg0, %arg1 : !spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>, vector<2xf32> -> vector<2xf32>
- spirv.Return
+ return
}
//===----------------------------------------------------------------------===//
@@ -346,7 +346,7 @@ func.func @image_fetch_float_coords(%arg0: !spirv.image<f32, Dim2D, NoDepth, Non
func.func @bias_too_many_arguments(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : f32, %arg2 : f32) -> () {
// expected-error @+1 {{too many image operand arguments have been provided}}
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 ["Bias"], %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, f32, f32, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -354,7 +354,7 @@ func.func @bias_too_many_arguments(%arg0 : !spirv.sampled_image<!spirv.image<f32
func.func @bias_too_many_arguments(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : f32, %arg2 : i32) -> () {
// expected-error @+1 {{Bias must be a floating-point type scalar}}
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 ["Bias"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, f32, i32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -362,7 +362,7 @@ func.func @bias_too_many_arguments(%arg0 : !spirv.sampled_image<!spirv.image<f32
func.func @bias_with_rect(%arg0 : !spirv.sampled_image<!spirv.image<f32, Rect, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : f32, %arg2 : f32) -> () {
// expected-error @+1 {{Bias must only be used with an image type that has a dim operand of 1D, 2D, 3D, or Cube}}
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 ["Bias"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Rect, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, f32, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// TODO: We cannot currently test Bias with MS != 0 as all implemented implicit operations already check for that.
@@ -376,7 +376,7 @@ func.func @bias_with_rect(%arg0 : !spirv.sampled_image<!spirv.image<f32, Rect, N
func.func @grad_and_lod_set(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : f32, %arg2 : f32) -> () {
// expected-error @+1 {{it is invalid to set both the Lod and Grad bits}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Lod | Grad"], %arg2, %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, f32, f32, f32, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -384,7 +384,7 @@ func.func @grad_and_lod_set(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim1D
func.func @lod_with_implict_sample(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{Lod is only valid with explicit-lod and fetch instructions}}
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 ["Lod"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -392,7 +392,7 @@ func.func @lod_with_implict_sample(%arg0 : !spirv.sampled_image<!spirv.image<f32
func.func @lod_too_many_arguments(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{too many image operand arguments have been provided}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Lod"], %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -400,7 +400,7 @@ func.func @lod_too_many_arguments(%arg0 : !spirv.sampled_image<!spirv.image<f32,
func.func @lod_with_rect(%arg0 : !spirv.sampled_image<!spirv.image<f32, Rect, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{Lod must only be used with an image type that has a dim operand of 1D, 2D, 3D, or Cube}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Lod"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Rect, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// TODO: We cannot currently test Lod with MS != 0 as all implemented explicit operations already check for that.
@@ -416,7 +416,7 @@ func.func @lod_with_rect(%arg0 : !spirv.sampled_image<!spirv.image<f32, Rect, No
func.func @gard_with_implicit_sample(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : vector<2xf32>) -> () {
// expected-error @+1 {{Grad is only valid with explicit-lod instructions}}
%0 = spirv.ImageSampleImplicitLod %arg0, %arg1 ["Grad"], %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, vector<2xf32>, vector<2xf32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -424,7 +424,7 @@ func.func @gard_with_implicit_sample(%arg0 : !spirv.sampled_image<!spirv.image<f
func.func @gard_not_enough_args(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : vector<2xf32>) -> () {
// expected-error @+1 {{Grad operand requires 2 arguments (scalars or vectors)}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Grad"], %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, vector<2xf32> -> vector<4xf32>
- spirv.Return
+ return
}
// TODO: We cannot currently test Grad with MS != 0 as all implemented explicit operations already check for that.
@@ -434,7 +434,7 @@ func.func @gard_not_enough_args(%arg0 : !spirv.sampled_image<!spirv.image<f32, D
func.func @grad_arg_size_mismatch(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : vector<3xf32>) -> () {
// expected-error @+1 {{number of components of each Grad argument must equal the number of components in coordinate, minus the array layer component, if present}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Grad"], %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, vector<3xf32>, vector<3xf32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -442,7 +442,7 @@ func.func @grad_arg_size_mismatch(%arg0 : !spirv.sampled_image<!spirv.image<f32,
func.func @gard_arg_wrong_type(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : vector<2xsi32>) -> () {
// expected-error @+1 {{Grad arguments must be a vector of floating-point type}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Grad"], %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, vector<2xsi32>, vector<2xsi32> -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -450,7 +450,7 @@ func.func @gard_arg_wrong_type(%arg0 : !spirv.sampled_image<!spirv.image<f32, Di
func.func @gard_arg_size_mismatch_scalar(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : f32) -> () {
// expected-error @+1 {{number of components of each Grad argument must equal the number of components in coordinate, minus the array layer component, if present}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Grad"], %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, f32, f32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -458,7 +458,7 @@ func.func @gard_arg_size_mismatch_scalar(%arg0 : !spirv.sampled_image<!spirv.ima
func.func @gard_int_args(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : i32) -> () {
// expected-error @+1 {{Grad arguments must be a scalar or vector of floating-point type}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Grad"], %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, i32, i32 -> vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -466,5 +466,5 @@ func.func @gard_int_args(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, N
func.func @gard_too_many_args(%arg0 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, %arg1 : vector<2xf32>, %arg2 : vector<2xf32>) -> () {
// expected-error @+1 {{too many image operand arguments have been provided}}
%0 = spirv.ImageSampleExplicitLod %arg0, %arg1 ["Grad"], %arg2, %arg2, %arg2 : !spirv.sampled_image<!spirv.image<f32, Dim2D, NoDepth, NonArrayed, SingleSampled, NeedSampler, Rgba8>>, vector<2xf32>, vector<2xf32>, vector<2xf32>, vector<2xf32> -> vector<4xf32>
- spirv.Return
+ return
}
diff --git a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
index a3b96c698a344..24c3fb3b3ace7 100644
--- a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
@@ -644,7 +644,7 @@ func.func @copy_memory_incompatible_ptrs() {
%1 = spirv.Variable : !spirv.ptr<i32, Function>
// expected-error @+1 {{both operands must be pointers to the same type}}
"spirv.CopyMemory"(%0, %1) {} : (!spirv.ptr<f32, Function>, !spirv.ptr<i32, Function>) -> ()
- spirv.Return
+ return
}
// -----
@@ -654,7 +654,7 @@ func.func @copy_memory_invalid_maa() {
%1 = spirv.Variable : !spirv.ptr<f32, Function>
// expected-error @+1 {{missing alignment value}}
"spirv.CopyMemory"(%0, %1) {memory_access=#spirv.memory_access<Aligned>} : (!spirv.ptr<f32, Function>, !spirv.ptr<f32, Function>) -> ()
- spirv.Return
+ return
}
// -----
@@ -664,7 +664,7 @@ func.func @copy_memory_invalid_source_maa() {
%1 = spirv.Variable : !spirv.ptr<f32, Function>
// expected-error @+1 {{invalid alignment specification with non-aligned memory access specification}}
"spirv.CopyMemory"(%0, %1) {source_memory_access=#spirv.memory_access<Volatile>, memory_access=#spirv.memory_access<Aligned>, source_alignment=8 : i32, alignment=4 : i32} : (!spirv.ptr<f32, Function>, !spirv.ptr<f32, Function>) -> ()
- spirv.Return
+ return
}
// -----
@@ -674,7 +674,7 @@ func.func @copy_memory_invalid_source_maa2() {
%1 = spirv.Variable : !spirv.ptr<f32, Function>
// expected-error @+1 {{missing alignment value}}
"spirv.CopyMemory"(%0, %1) {source_memory_access=#spirv.memory_access<Aligned>, memory_access=#spirv.memory_access<Aligned>, alignment=4 : i32} : (!spirv.ptr<f32, Function>, !spirv.ptr<f32, Function>) -> ()
- spirv.Return
+ return
}
// -----
@@ -695,7 +695,7 @@ func.func @copy_memory_print_maa() {
// CHECK: spirv.CopyMemory "Function" %{{.*}}, "Function" %{{.*}} ["Aligned", 4], ["Aligned", 8] : f32
"spirv.CopyMemory"(%0, %1) {source_memory_access=#spirv.memory_access<Aligned>, memory_access=#spirv.memory_access<Aligned>, source_alignment=8 : i32, alignment=4 : i32} : (!spirv.ptr<f32, Function>, !spirv.ptr<f32, Function>) -> ()
- spirv.Return
+ return
}
// -----
diff --git a/mlir/test/Dialect/SPIRV/IR/mesh-ops.mlir b/mlir/test/Dialect/SPIRV/IR/mesh-ops.mlir
index 436f7d1c9fb15..a87c80347b8dd 100644
--- a/mlir/test/Dialect/SPIRV/IR/mesh-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/mesh-ops.mlir
@@ -30,5 +30,5 @@ func.func @emit_mesh_tasks_wrong_payload(%0 : i32, %1 : !spirv.ptr<i32, Image>)
func.func @set_mesh_outputs(%0 : i32, %1 : i32) {
// CHECK: spirv.EXT.SetMeshOutputs {{%.*}}, {{%.*}} : i32, i32
spirv.EXT.SetMeshOutputs %0, %1 : i32, i32
- spirv.Return
+ return
}
diff --git a/mlir/test/Dialect/SPIRV/IR/misc-ops.mlir b/mlir/test/Dialect/SPIRV/IR/misc-ops.mlir
index 182b661035a61..114439f31cfb5 100644
--- a/mlir/test/Dialect/SPIRV/IR/misc-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/misc-ops.mlir
@@ -9,7 +9,7 @@ func.func @undef() -> () {
%0 = spirv.Undef : f32
// CHECK: %{{.*}} = spirv.Undef : vector<4xf32>
%1 = spirv.Undef : vector<4xf32>
- spirv.Return
+ return
}
// -----
@@ -17,7 +17,7 @@ func.func @undef() -> () {
func.func @undef() -> () {
// expected-error @+1{{expected non-function type}}
%0 = spirv.Undef :
- spirv.Return
+ return
}
// -----
@@ -25,7 +25,7 @@ func.func @undef() -> () {
func.func @undef() -> () {
// expected-error @+1{{expected ':'}}
%0 = spirv.Undef
- spirv.Return
+ return
}
// -----
@@ -33,7 +33,7 @@ func.func @undef() -> () {
func.func @assume_true(%arg : i1) -> () {
// CHECK: spirv.KHR.AssumeTrue %{{.*}}
spirv.KHR.AssumeTrue %arg
- spirv.Return
+ return
}
// -----
@@ -42,5 +42,5 @@ func.func @assume_true(%arg : f32) -> () {
// expected-error @+2{{use of value '%arg' expects different type than prior uses: 'i1' vs 'f32'}}
// expected-note @-2 {{prior use here}}
spirv.KHR.AssumeTrue %arg
- spirv.Return
+ return
}
diff --git a/mlir/test/Dialect/SPIRV/IR/primitive-ops.mlir b/mlir/test/Dialect/SPIRV/IR/primitive-ops.mlir
index 451c3345b4e0d..9d0379b6ca653 100644
--- a/mlir/test/Dialect/SPIRV/IR/primitive-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/primitive-ops.mlir
@@ -7,7 +7,7 @@
func.func @emit_vertex() {
// CHECK: spirv.EmitVertex
spirv.EmitVertex
- spirv.Return
+ return
}
//===----------------------------------------------------------------------===//
@@ -17,5 +17,5 @@ func.func @emit_vertex() {
func.func @end_primitive() {
// CHECK: spirv.EndPrimitive
spirv.EndPrimitive
- spirv.Return
+ return
}
diff --git a/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir b/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
index 722c27586aa61..f2cb6d47608cc 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
@@ -14,7 +14,7 @@ func.func @combine_full_access_chain() -> f32 {
%1 = spirv.AccessChain %0[%c0] : !spirv.ptr<!spirv.struct<(!spirv.array<4x!spirv.array<4xf32>>, !spirv.array<4xi32>)>, Function>, i32 -> !spirv.ptr<!spirv.array<4x!spirv.array<4xf32>>, Function>
%2 = spirv.AccessChain %1[%c0, %c0] : !spirv.ptr<!spirv.array<4x!spirv.array<4xf32>>, Function>, i32, i32 -> !spirv.ptr<f32, Function>
%3 = spirv.Load "Function" %2 : f32
- spirv.ReturnValue %3 : f32
+ return %3 : f32
}
// -----
@@ -33,7 +33,7 @@ func.func @combine_access_chain_multi_use() -> !spirv.array<4xf32> {
%3 = spirv.AccessChain %2[%c0] : !spirv.ptr<!spirv.array<4xf32>, Function>, i32 -> !spirv.ptr<f32, Function>
%4 = spirv.Load "Function" %2 : !spirv.array<4xf32>
%5 = spirv.Load "Function" %3 : f32
- spirv.ReturnValue %4: !spirv.array<4xf32>
+ return %4: !spirv.array<4xf32>
}
// -----
@@ -53,7 +53,7 @@ func.func @dont_combine_access_chain_without_common_base() -> !spirv.array<4xi32
%3 = spirv.AccessChain %1[%c1] : !spirv.ptr<!spirv.struct<(!spirv.array<4x!spirv.array<4xf32>>, !spirv.array<4xi32>)>, Function>, i32 -> !spirv.ptr<!spirv.array<4xi32>, Function>
%4 = spirv.Load "Function" %2 : !spirv.array<4xi32>
%5 = spirv.Load "Function" %3 : !spirv.array<4xi32>
- spirv.ReturnValue %4 : !spirv.array<4xi32>
+ return %4 : !spirv.array<4xi32>
}
// -----
@@ -64,11 +64,11 @@ func.func @dont_combine_access_chain_without_common_base() -> !spirv.array<4xi32
func.func @convert_bitcast_full(%arg0 : vector<2xf32>) -> f64 {
// CHECK: %[[RESULT:.*]] = spirv.Bitcast {{%.*}} : vector<2xf32> to f64
- // CHECK-NEXT: spirv.ReturnValue %[[RESULT]]
+ // CHECK-NEXT: return %[[RESULT]]
%0 = spirv.Bitcast %arg0 : vector<2xf32> to vector<2xi32>
%1 = spirv.Bitcast %0 : vector<2xi32> to i64
%2 = spirv.Bitcast %1 : i64 to f64
- spirv.ReturnValue %2 : f64
+ return %2 : f64
}
// -----
@@ -77,11 +77,11 @@ func.func @convert_bitcast_multi_use(%arg0 : vector<2xf32>, %arg1 : !spirv.ptr<i
// CHECK: %[[RESULT_0:.*]] = spirv.Bitcast {{%.*}} : vector<2xf32> to i64
// CHECK-NEXT: %[[RESULT_1:.*]] = spirv.Bitcast {{%.*}} : vector<2xf32> to f64
// CHECK-NEXT: spirv.Store {{".*"}} {{%.*}}, %[[RESULT_0]]
- // CHECK-NEXT: spirv.ReturnValue %[[RESULT_1]]
+ // CHECK-NEXT: return %[[RESULT_1]]
%0 = spirv.Bitcast %arg0 : vector<2xf32> to i64
%1 = spirv.Bitcast %0 : i64 to f64
spirv.Store "Uniform" %arg1, %0 : i64
- spirv.ReturnValue %1 : f64
+ return %1 : f64
}
// -----
@@ -89,10 +89,10 @@ func.func @convert_bitcast_multi_use(%arg0 : vector<2xf32>, %arg1 : !spirv.ptr<i
// CHECK-LABEL: @convert_bitcast_roundtip
// CHECK-SAME: %[[ARG:.+]]: i64
func.func @convert_bitcast_roundtip(%arg0 : i64) -> i64 {
- // CHECK: spirv.ReturnValue %[[ARG]]
+ // CHECK: return %[[ARG]]
%0 = spirv.Bitcast %arg0 : i64 to f64
%1 = spirv.Bitcast %0 : f64 to i64
- spirv.ReturnValue %1 : i64
+ return %1 : i64
}
// -----
@@ -100,12 +100,12 @@ func.func @convert_bitcast_roundtip(%arg0 : i64) -> i64 {
// CHECK-LABEL: @convert_bitcast_chained_roundtip
// CHECK-SAME: %[[ARG:.+]]: i64
func.func @convert_bitcast_chained_roundtip(%arg0 : i64) -> i64 {
- // CHECK: spirv.ReturnValue %[[ARG]]
+ // CHECK: return %[[ARG]]
%0 = spirv.Bitcast %arg0 : i64 to f64
%1 = spirv.Bitcast %0 : f64 to vector<2xi32>
%2 = spirv.Bitcast %1 : vector<2xi32> to vector<2xf32>
%3 = spirv.Bitcast %2 : vector<2xf32> to i64
- spirv.ReturnValue %3 : i64
+ return %3 : i64
}
// -----
@@ -156,7 +156,7 @@ func.func @extract_from_not_constant() -> i32 {
%1 = spirv.Load "Function" %0 : vector<3xi32>
// CHECK: spirv.CompositeExtract
%2 = spirv.CompositeExtract %1[0 : i32] : vector<3xi32>
- spirv.ReturnValue %2 : i32
+ return %2 : i32
}
// -----
@@ -1222,10 +1222,10 @@ func.func @const_fold_vector_logical_not() -> vector<2xi1> {
func.func @convert_logical_not_to_not_equal(%arg0: vector<3xi64>, %arg1: vector<3xi64>) -> vector<3xi1> {
// CHECK: %[[RESULT:.*]] = spirv.INotEqual {{%.*}}, {{%.*}} : vector<3xi64>
- // CHECK-NEXT: spirv.ReturnValue %[[RESULT]] : vector<3xi1>
+ // CHECK-NEXT: return %[[RESULT]] : vector<3xi1>
%2 = spirv.IEqual %arg0, %arg1 : vector<3xi64>
%3 = spirv.LogicalNot %2 : vector<3xi1>
- spirv.ReturnValue %3 : vector<3xi1>
+ return %3 : vector<3xi1>
}
// -----
@@ -1281,9 +1281,9 @@ func.func @const_fold_vector_logical_equal() -> vector<3xi1> {
// CHECK-SAME: %[[ARG:.+]]: vector<4xi1>
func.func @convert_logical_not_equal_false(%arg: vector<4xi1>) -> vector<4xi1> {
%cst = spirv.Constant dense<false> : vector<4xi1>
- // CHECK: spirv.ReturnValue %[[ARG]] : vector<4xi1>
+ // CHECK: return %[[ARG]] : vector<4xi1>
%0 = spirv.LogicalNotEqual %arg, %cst : vector<4xi1>
- spirv.ReturnValue %0 : vector<4xi1>
+ return %0 : vector<4xi1>
}
// CHECK-LABEL: @logical_not_equal_same
@@ -1327,10 +1327,10 @@ func.func @const_fold_vector_logical_not_equal() -> vector<3xi1> {
func.func @convert_logical_not_to_equal(%arg0: vector<3xi64>, %arg1: vector<3xi64>) -> vector<3xi1> {
// CHECK: %[[RESULT:.*]] = spirv.IEqual {{%.*}}, {{%.*}} : vector<3xi64>
- // CHECK-NEXT: spirv.ReturnValue %[[RESULT]] : vector<3xi1>
+ // CHECK-NEXT: return %[[RESULT]] : vector<3xi1>
%2 = spirv.INotEqual %arg0, %arg1 : vector<3xi64>
%3 = spirv.LogicalNot %2 : vector<3xi1>
- spirv.ReturnValue %3 : vector<3xi1>
+ return %3 : vector<3xi1>
}
// -----
@@ -1339,31 +1339,31 @@ func.func @convert_logical_not_parent_multi_use(%arg0: vector<3xi64>, %arg1: vec
// CHECK: %[[RESULT_0:.*]] = spirv.INotEqual {{%.*}}, {{%.*}} : vector<3xi64>
// CHECK-NEXT: %[[RESULT_1:.*]] = spirv.IEqual {{%.*}}, {{%.*}} : vector<3xi64>
// CHECK-NEXT: spirv.Store "Uniform" {{%.*}}, %[[RESULT_0]]
- // CHECK-NEXT: spirv.ReturnValue %[[RESULT_1]]
+ // CHECK-NEXT: return %[[RESULT_1]]
%0 = spirv.INotEqual %arg0, %arg1 : vector<3xi64>
%1 = spirv.LogicalNot %0 : vector<3xi1>
spirv.Store "Uniform" %arg2, %0 : vector<3xi1>
- spirv.ReturnValue %1 : vector<3xi1>
+ return %1 : vector<3xi1>
}
// -----
func.func @convert_logical_not_to_logical_not_equal(%arg0: vector<3xi1>, %arg1: vector<3xi1>) -> vector<3xi1> {
// CHECK: %[[RESULT:.*]] = spirv.LogicalNotEqual {{%.*}}, {{%.*}} : vector<3xi1>
- // CHECK-NEXT: spirv.ReturnValue %[[RESULT]] : vector<3xi1>
+ // CHECK-NEXT: return %[[RESULT]] : vector<3xi1>
%2 = spirv.LogicalEqual %arg0, %arg1 : vector<3xi1>
%3 = spirv.LogicalNot %2 : vector<3xi1>
- spirv.ReturnValue %3 : vector<3xi1>
+ return %3 : vector<3xi1>
}
// -----
func.func @convert_logical_not_to_logical_equal(%arg0: vector<3xi1>, %arg1: vector<3xi1>) -> vector<3xi1> {
// CHECK: %[[RESULT:.*]] = spirv.LogicalEqual {{%.*}}, {{%.*}} : vector<3xi1>
- // CHECK-NEXT: spirv.ReturnValue %[[RESULT]] : vector<3xi1>
+ // CHECK-NEXT: return %[[RESULT]] : vector<3xi1>
%2 = spirv.LogicalNotEqual %arg0, %arg1 : vector<3xi1>
%3 = spirv.LogicalNot %2 : vector<3xi1>
- spirv.ReturnValue %3 : vector<3xi1>
+ return %3 : vector<3xi1>
}
// -----
@@ -2263,7 +2263,7 @@ func.func @canonicalize_selection_op_scalar_type(%cond: i1) -> () {
// CHECK: %[[SRC_VALUE:.*]] = spirv.Select {{%.*}}, %[[TRUE_VALUE]], %[[FALSE_VALUE]] : i1, i32
// CHECK-NEXT: spirv.Store "Function" %[[DST_VAR]], %[[SRC_VALUE]] ["Aligned", 4] : i32
- // CHECK-NEXT: spirv.Return
+ // CHECK-NEXT: return
spirv.mlir.selection {
spirv.BranchConditional %cond, ^then, ^else
@@ -2278,7 +2278,7 @@ func.func @canonicalize_selection_op_scalar_type(%cond: i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -2294,7 +2294,7 @@ func.func @canonicalize_selection_op_vector_type(%cond: i1) -> () {
// CHECK: %[[SRC_VALUE:.*]] = spirv.Select {{%.*}}, %[[TRUE_VALUE]], %[[FALSE_VALUE]] : i1, vector<3xi32>
// CHECK-NEXT: spirv.Store "Function" %[[DST_VAR]], %[[SRC_VALUE]] ["Aligned", 8] : vector<3xi32>
- // CHECK-NEXT: spirv.Return
+ // CHECK-NEXT: return
spirv.mlir.selection {
spirv.BranchConditional %cond, ^then, ^else
@@ -2309,7 +2309,7 @@ func.func @canonicalize_selection_op_vector_type(%cond: i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -2348,7 +2348,7 @@ func.func @cannot_canonicalize_selection_op_0(%cond: i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -2386,7 +2386,7 @@ func.func @cannot_canonicalize_selection_op_1(%cond: i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -2420,7 +2420,7 @@ func.func @cannot_canonicalize_selection_op_2(%cond: i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -2454,7 +2454,7 @@ func.func @cannot_canonicalize_selection_op_3(%cond: i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
// -----
@@ -2488,5 +2488,5 @@ func.func @cannot_canonicalize_selection_op_4(%cond: i1) -> () {
^merge:
spirv.mlir.merge
}
- spirv.Return
+ return
}
diff --git a/mlir/test/Dialect/SPIRV/Transforms/gl-canonicalize.mlir b/mlir/test/Dialect/SPIRV/Transforms/gl-canonicalize.mlir
index c1447b38f0a48..ace73504814f2 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/gl-canonicalize.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/gl-canonicalize.mlir
@@ -9,8 +9,8 @@ func.func @clamp_fordlessthan(%input: f32, %min: f32, %max: f32) -> f32 {
%1 = spirv.FOrdLessThan %mid, %max : f32
%2 = spirv.Select %1, %mid, %max : i1, f32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : f32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : f32
}
// -----
@@ -24,8 +24,8 @@ func.func @clamp_fordlessthan(%input: f32, %min: f32, %max: f32) -> f32 {
%1 = spirv.FOrdLessThan %max, %input : f32
%2 = spirv.Select %1, %max, %mid : i1, f32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : f32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : f32
}
// -----
@@ -39,8 +39,8 @@ func.func @clamp_fordlessthanequal(%input: f32, %min: f32, %max: f32) -> f32 {
%1 = spirv.FOrdLessThanEqual %mid, %max : f32
%2 = spirv.Select %1, %mid, %max : i1, f32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : f32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : f32
}
// -----
@@ -54,8 +54,8 @@ func.func @clamp_fordlessthanequal(%input: f32, %min: f32, %max: f32) -> f32 {
%1 = spirv.FOrdLessThanEqual %max, %input : f32
%2 = spirv.Select %1, %max, %mid : i1, f32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : f32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : f32
}
// -----
@@ -69,8 +69,8 @@ func.func @clamp_slessthan(%input: si32, %min: si32, %max: si32) -> si32 {
%1 = spirv.SLessThan %mid, %max : si32
%2 = spirv.Select %1, %mid, %max : i1, si32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : si32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : si32
}
// -----
@@ -84,8 +84,8 @@ func.func @clamp_slessthan(%input: si32, %min: si32, %max: si32) -> si32 {
%1 = spirv.SLessThan %max, %input : si32
%2 = spirv.Select %1, %max, %mid : i1, si32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : si32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : si32
}
// -----
@@ -99,8 +99,8 @@ func.func @clamp_slessthanequal(%input: si32, %min: si32, %max: si32) -> si32 {
%1 = spirv.SLessThanEqual %mid, %max : si32
%2 = spirv.Select %1, %mid, %max : i1, si32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : si32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : si32
}
// -----
@@ -114,8 +114,8 @@ func.func @clamp_slessthanequal(%input: si32, %min: si32, %max: si32) -> si32 {
%1 = spirv.SLessThanEqual %max, %input : si32
%2 = spirv.Select %1, %max, %mid : i1, si32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : si32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : si32
}
// -----
@@ -129,8 +129,8 @@ func.func @clamp_ulessthan(%input: i32, %min: i32, %max: i32) -> i32 {
%1 = spirv.ULessThan %mid, %max : i32
%2 = spirv.Select %1, %mid, %max : i1, i32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : i32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : i32
}
// -----
@@ -144,8 +144,8 @@ func.func @clamp_ulessthan(%input: i32, %min: i32, %max: i32) -> i32 {
%1 = spirv.ULessThan %max, %input : i32
%2 = spirv.Select %1, %max, %mid : i1, i32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : i32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : i32
}
// -----
@@ -159,8 +159,8 @@ func.func @clamp_ulessthanequal(%input: i32, %min: i32, %max: i32) -> i32 {
%1 = spirv.ULessThanEqual %mid, %max : i32
%2 = spirv.Select %1, %mid, %max : i1, i32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : i32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : i32
}
// -----
@@ -174,6 +174,6 @@ func.func @clamp_ulessthanequal(%input: i32, %min: i32, %max: i32) -> i32 {
%1 = spirv.ULessThanEqual %max, %input : i32
%2 = spirv.Select %1, %max, %mid : i1, i32
- // CHECK-NEXT: spirv.ReturnValue [[RES]]
- spirv.ReturnValue %2 : i32
+ // CHECK-NEXT: return [[RES]]
+ return %2 : i32
}
diff --git a/mlir/test/IR/test-region-branch-op-verifier.mlir b/mlir/test/IR/test-region-branch-op-verifier.mlir
index 1c7a87c4f0356..dc01266d288b5 100644
--- a/mlir/test/IR/test-region-branch-op-verifier.mlir
+++ b/mlir/test/IR/test-region-branch-op-verifier.mlir
@@ -24,14 +24,8 @@ func.func @test_no_terminator(%arg: index) {
// -----
-// test.loop_block_term has two operands: iter (i32, passed back to the region)
-// and exit (f32, passed to the parent). getMutableSuccessorOperands(parent)
-// returns only the exit operand. The function returns f32, matching the exit
-// operand type, so verification must succeed.
-//
-// A verifier using getNumOperands() instead would incorrectly report "has 2
-// operands, but enclosing function returns 1".
-func.func @func_with_region_branch_terminator(%arg: i32) -> f32 {
+// func.func requires func.return as the only region terminator.
+func.func @func_with_return_terminator(%arg: i32) -> f32 {
%0 = "test.constant"() { value = 5.3 : f32 } : () -> f32
- test.loop_block_term iter %arg exit %0
+ return %0 : f32
}
diff --git a/mlir/test/Target/LLVMIR/llvmir-invalid.mlir b/mlir/test/Target/LLVMIR/llvmir-invalid.mlir
index c263afe553750..6b2fd716c4d59 100644
--- a/mlir/test/Target/LLVMIR/llvmir-invalid.mlir
+++ b/mlir/test/Target/LLVMIR/llvmir-invalid.mlir
@@ -1,7 +1,6 @@
// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s
-// expected-error @below{{cannot be converted to LLVM IR}}
-func.func @foo() {
+llvm.func @foo() {
llvm.return
}
diff --git a/mlir/test/Transforms/print-op-graph.mlir b/mlir/test/Transforms/print-op-graph.mlir
index 6a24d914b65a5..becfa8ee23d0f 100644
--- a/mlir/test/Transforms/print-op-graph.mlir
+++ b/mlir/test/Transforms/print-op-graph.mlir
@@ -17,7 +17,7 @@
// DFG: subgraph {{.*}} {
// DFG: }
// DFG: }
-// DFG: v[[TEST_RET:.*]] [{{.*}}label = "{{.*}}test.return
+// DFG: v[[TEST_RET:.*]] [{{.*}}label = "{{.*}}return
// DFG: v[[ARG0]]:res_arg0:s -> v[[TEST_BR]]:arg_arg0:n
// DFG: v[[CONST10]]:res_c10_i32:s -> v[[TEST_BR]]
// DFG: v[[ANCHOR]] -> v[[TEST_RET]]:arg_1_0:n[ltail = [[CLUSTER_MERGE_BLOCKS]], style = solid];
@@ -42,7 +42,7 @@
// CFG: subgraph {{.*}} {
// CFG: }
// CFG: }
-// CFG: v[[TEST_RET:.*]] [{{.*}}label = "{{.*}}test.return
+// CFG: v[[TEST_RET:.*]] [{{.*}}label = "{{.*}}return
// CFG: v[[C1]] -> v[[C2]]
// CFG: v[[C2]] -> v[[C3]]
// CFG: v[[C3]] -> v[[C4]]
@@ -62,5 +62,5 @@ func.func @merge_blocks(%arg0: i32, %arg1 : i32) -> (i32, i32) {
^bb1(%arg3 : i32, %arg4 : i32, %arg5: i32):
"test.return"(%arg3, %arg4) : (i32, i32) -> ()
}) : () -> (i32, i32)
- "test.return"(%3#0, %3#1) : (i32, i32) -> ()
+ func.return %3#0, %3#1 : i32, i32
}
diff --git a/mlir/test/Transforms/remove-dead-values.mlir b/mlir/test/Transforms/remove-dead-values.mlir
index 19bc6b2fddd66..52b98d82abc2d 100644
--- a/mlir/test/Transforms/remove-dead-values.mlir
+++ b/mlir/test/Transforms/remove-dead-values.mlir
@@ -480,14 +480,14 @@ func.func @kernel(%arg0: memref<18xf32>) {
// CHECK-LABEL: llvm_unreachable
// CHECK-LABEL: @fn_with_llvm_unreachable
// CHECK-LABEL: @main
-// CHECK: llvm.return
+// CHECK: return
module @llvm_unreachable {
func.func private @fn_with_llvm_unreachable(%arg0: tensor<4x4xf32>) -> tensor<4x4xi1> {
llvm.unreachable
}
func.func private @main(%arg0: tensor<4x4xf32>) {
%0 = call @fn_with_llvm_unreachable(%arg0) : (tensor<4x4xf32>) -> tensor<4x4xi1>
- llvm.return
+ func.return
}
}
diff --git a/mlir/test/Transforms/sccp.mlir b/mlir/test/Transforms/sccp.mlir
index c78c8594c0ba5..a2489eb85ed49 100644
--- a/mlir/test/Transforms/sccp.mlir
+++ b/mlir/test/Transforms/sccp.mlir
@@ -253,5 +253,5 @@ func.func @no_crash_with_different_source_type() {
%0 = llvm.mlir.constant(0 : index) : i64
// CHECK: vector.broadcast %[[CST:.*]] : i64 to vector<128xi64>
%1 = vector.broadcast %0 : i64 to vector<128xi64>
- llvm.return
+ func.return
}
diff --git a/mlir/test/Transforms/test-dialect-conversion-pdll.mlir b/mlir/test/Transforms/test-dialect-conversion-pdll.mlir
index 9abc7c83bff8c..afb59883063d3 100644
--- a/mlir/test/Transforms/test-dialect-conversion-pdll.mlir
+++ b/mlir/test/Transforms/test-dialect-conversion-pdll.mlir
@@ -3,9 +3,9 @@
// CHECK-LABEL: @TestSingleConversion
func.func @TestSingleConversion() {
// CHECK: "test.cast"() : () -> f64
- // CHECK-NEXT: "test.return"() : () -> ()
+ // CHECK-NEXT: return
%result = "test.cast"() : () -> (i64)
- "test.return"() : () -> ()
+ func.return
}
// CHECK-LABEL: @TestLingeringConversion
diff --git a/mlir/test/Transforms/test-legalize-remapped-value.mlir b/mlir/test/Transforms/test-legalize-remapped-value.mlir
index d0752405fa109..cae15c8891c8a 100644
--- a/mlir/test/Transforms/test-legalize-remapped-value.mlir
+++ b/mlir/test/Transforms/test-legalize-remapped-value.mlir
@@ -10,7 +10,7 @@
func.func @remap_input_1_to_1(%arg0: i32) {
%0 = "test.one_variadic_out_one_variadic_in1"(%arg0) : (i32) -> i32
%1 = "test.one_variadic_out_one_variadic_in1"(%0) : (i32) -> i32
- "test.return"() : () -> ()
+ func.return
}
// Test the case where an operation is converted before its operands are.
@@ -24,5 +24,5 @@ func.func @remap_unconverted() {
"test.return"(%result) : (f32) -> ()
}) : () -> (f32)
"test.type_consumer"(%region_result) : (f32) -> ()
- "test.return"() : () -> ()
+ func.return
}
diff --git a/mlir/test/Transforms/test-legalizer-analysis.mlir b/mlir/test/Transforms/test-legalizer-analysis.mlir
index 5b070557e2811..8c8e26e60eedf 100644
--- a/mlir/test/Transforms/test-legalizer-analysis.mlir
+++ b/mlir/test/Transforms/test-legalizer-analysis.mlir
@@ -9,6 +9,7 @@ func.func @test(%arg0: f32) {
// expected-remark at +1 {{op 'test.invalid' is legalizable}}
"test.invalid"() : () -> ()
}) : () -> ()
+ // expected-remark at +1 {{op 'func.return' is legalizable}}
return
}
diff --git a/mlir/test/Transforms/test-legalizer-fold-after.mlir b/mlir/test/Transforms/test-legalizer-fold-after.mlir
index 7f80252dc9604..6fd975240d58e 100644
--- a/mlir/test/Transforms/test-legalizer-fold-after.mlir
+++ b/mlir/test/Transforms/test-legalizer-fold-after.mlir
@@ -5,5 +5,5 @@ func.func @fold_legalization() -> i32 {
// CHECK-NOT: op_in_place_self_fold
// CHECK: 97
%1 = "test.op_in_place_self_fold"() : () -> (i32)
- "test.return"(%1) : (i32) -> ()
+ "func.return"(%1) : (i32) -> ()
}
diff --git a/mlir/test/Transforms/test-legalizer-fold-before.mlir b/mlir/test/Transforms/test-legalizer-fold-before.mlir
index fe6e29351a5d7..515a5d0fa83dc 100644
--- a/mlir/test/Transforms/test-legalizer-fold-before.mlir
+++ b/mlir/test/Transforms/test-legalizer-fold-before.mlir
@@ -5,5 +5,5 @@ func.func @fold_legalization() -> i32 {
// CHECK: op_in_place_self_fold
// CHECK-SAME: folded
%1 = "test.op_in_place_self_fold"() : () -> (i32)
- "test.return"(%1) : (i32) -> ()
+ func.return %1 : i32
}
diff --git a/mlir/test/Transforms/test-legalizer-full-rollback.mlir b/mlir/test/Transforms/test-legalizer-full-rollback.mlir
index c61847b55cf01..fb623bc5a93d6 100644
--- a/mlir/test/Transforms/test-legalizer-full-rollback.mlir
+++ b/mlir/test/Transforms/test-legalizer-full-rollback.mlir
@@ -9,7 +9,7 @@
// CHECK: ^[[bb1]](%[[arg1:.*]]: i64):
// CHECK: "test.invalid"(%[[arg1]]) : (i64) -> ()
// CHECK: }) : () -> ()
-// CHECK: "test.return"() : () -> ()
+// CHECK: return
// CHECK: }
// expected-remark at +1 {{applyFullConversion failed}}
@@ -23,7 +23,7 @@ func.func @test_undo_region_inline() {
"test.invalid"(%i1) : (i64) -> ()
}) {} : () -> ()
- "test.return"() : () -> ()
+ func.return
}
}
@@ -40,7 +40,7 @@ func.func @test_undo_region_inline() {
// CHECK: ^[[bb2]](%[[arg2:.*]]: i64):
// CHECK: cf.br ^[[bb1]](%[[arg2]] : i64)
// CHECK: }) {legalizer.erase_old_blocks, legalizer.should_clone} : () -> ()
-// CHECK: "test.return"() : () -> ()
+// CHECK: return
// CHECK: }
// expected-remark at +1 {{applyFullConversion failed}}
@@ -55,6 +55,6 @@ func.func @test_undo_block_erase() {
^bb3(%i2: i64):
cf.br ^bb2(%i2 : i64)
}) {legalizer.should_clone, legalizer.erase_old_blocks} : () -> ()
- "test.return"() : () -> ()
+ func.return
}
}
diff --git a/mlir/test/Transforms/test-legalizer-full.mlir b/mlir/test/Transforms/test-legalizer-full.mlir
index 8da9109a32762..1c63347e3666b 100644
--- a/mlir/test/Transforms/test-legalizer-full.mlir
+++ b/mlir/test/Transforms/test-legalizer-full.mlir
@@ -6,7 +6,7 @@ func.func @multi_level_mapping() {
// CHECK: "test.type_consumer"(%{{.*}}) : (f64) -> ()
%result = "test.type_producer"() : () -> i32
"test.type_consumer"(%result) : (i32) -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -14,12 +14,12 @@ func.func @multi_level_mapping() {
// Test that operations that are erased don't need to be legalized.
// CHECK-LABEL: func @dropped_region_with_illegal_ops
func.func @dropped_region_with_illegal_ops() {
- // CHECK-NEXT: test.return
+ // CHECK-NEXT: return
"test.drop_region_op"() ({
%ignored = "test.illegal_op_f"() : () -> (i32)
"test.return"() : () -> ()
}) : () -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -27,9 +27,9 @@ func.func @dropped_region_with_illegal_ops() {
// CHECK-LABEL: func @replace_non_root_illegal_op
func.func @replace_non_root_illegal_op() {
// CHECK-NEXT: "test.legal_op_b"
- // CHECK-NEXT: test.return
+ // CHECK-NEXT: return
%result = "test.replace_non_root"() : () -> (i32)
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -50,11 +50,11 @@ func.func @recursively_legal_invalid_op() {
func.func @dynamic_func(%arg: i64) attributes {test.recursively_legal} {
// CHECK: "test.illegal_op_f"
%ignored = "test.illegal_op_f"() : () -> (i32)
- "test.return"() : () -> ()
+ func.return
}
}
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -68,7 +68,7 @@ builtin.module {
// expected-error at +1 {{failed to legalize operation 'foo.unknown_op'}}
"foo.unknown_op"() {} : () -> ()
- "test.return"() : () -> ()
+ func.return
}
}
diff --git a/mlir/test/Transforms/test-legalizer-no-fold.mlir b/mlir/test/Transforms/test-legalizer-no-fold.mlir
index 720d17f41943d..3605480c7c9b3 100644
--- a/mlir/test/Transforms/test-legalizer-no-fold.mlir
+++ b/mlir/test/Transforms/test-legalizer-no-fold.mlir
@@ -7,6 +7,6 @@ func.func @remove_foldable_op(%arg0 : i32) -> (i32) {
%0 = "test.op_with_region_fold"(%arg0) ({
"foo.op_with_region_terminator"() : () -> ()
}) : (i32) -> (i32)
- "test.return"(%0) : (i32) -> ()
+ func.return %0 : i32
}
diff --git a/mlir/test/Transforms/test-legalizer-no-materializations.mlir b/mlir/test/Transforms/test-legalizer-no-materializations.mlir
index 82dd7422b22b2..1556632ec5e7f 100644
--- a/mlir/test/Transforms/test-legalizer-no-materializations.mlir
+++ b/mlir/test/Transforms/test-legalizer-no-materializations.mlir
@@ -24,7 +24,6 @@ func.func @test_lookup_without_converter() {
// Make sure that the second "replace_with_valid_consumer" lowering does not
// lookup the materialization that was created for the above op.
"test.replace_with_valid_consumer"(%0) : (i64) -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -43,7 +42,6 @@ func.func @remap_moved_region_args() {
^bb1(%i0: i64, %unused: i16, %i1: i64, %2: f32):
"test.invalid"(%i0, %i1, %2) : (i64, i64, f32) -> ()
}) : () -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -62,6 +60,5 @@ func.func @remap_cloned_region_args() {
^bb1(%i0: i64, %unused: i16, %i1: i64, %2: f32):
"test.invalid"(%i0, %i1, %2) : (i64, i64, f32) -> ()
}) {legalizer.should_clone} : () -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
diff --git a/mlir/test/Transforms/test-legalizer-no-rollback.mlir b/mlir/test/Transforms/test-legalizer-no-rollback.mlir
index 5f421a35d956b..32e34393d5859 100644
--- a/mlir/test/Transforms/test-legalizer-no-rollback.mlir
+++ b/mlir/test/Transforms/test-legalizer-no-rollback.mlir
@@ -19,5 +19,5 @@ func.func @conditional_replacement(%arg0: i42) {
"test.dummy_user_2"(%arg0) {} : (i42) -> ()
// Perform a conditional 1:N replacement.
"test.value_replace"(%arg0, %repl) {conditional} : (i42, i42) -> ()
- "test.return"() : () -> ()
+ func.return
}
diff --git a/mlir/test/Transforms/test-legalizer-rollback.mlir b/mlir/test/Transforms/test-legalizer-rollback.mlir
index f6569201842b7..d64ece7fc50a6 100644
--- a/mlir/test/Transforms/test-legalizer-rollback.mlir
+++ b/mlir/test/Transforms/test-legalizer-rollback.mlir
@@ -39,7 +39,6 @@ func.func @create_illegal_block() {
// expected-remark at +1 {{op 'test.create_illegal_block' is not legalizable}}
"test.create_illegal_block"() : () -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -95,7 +94,6 @@ func.func @undo_block_erase() {
func.func @undo_child_created_before_parent() {
// expected-remark at +1 {{is not legalizable}}
"test.illegal_op_with_region_anchor"() : () -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -106,7 +104,7 @@ builtin.module {
func.func @create_unregistered_op_in_pattern() -> i32 {
// expected-error at +1 {{failed to legalize operation 'test.illegal_op_g'}}
%0 = "test.illegal_op_g"() : () -> (i32)
- "test.return"(%0) : (i32) -> ()
+ func.return %0 : i32
}
}
@@ -121,7 +119,7 @@ func.func @test_move_op_before_rollback() {
%0 = "test.hoist_me"() : () -> (i32)
"test.valid"(%0) : (i32) -> ()
}) : () -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -133,7 +131,7 @@ func.func @test_properties_rollback() {
test.with_properties
a = 32, b = "foo", c = "bar", flag = true, array = [1, 2, 3, 4], array32 = [5, 6]
{modify_inplace}
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -145,7 +143,7 @@ func.func @test_undo_block_move_detached() {
^bb0(%arg0: i64):
"test.return"() : () -> ()
}) : () -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -161,7 +159,7 @@ func.func @test_undo_region_clone() {
// expected-error at +1 {{failed to legalize operation 'test.illegal_op_f'}}
%ignored = "test.illegal_op_f"() : () -> (i32)
- "test.return"() : () -> ()
+ func.return
}
}
@@ -172,7 +170,7 @@ builtin.module {
func.func @create_unregistered_op_in_pattern() -> i32 {
// expected-error at +1 {{failed to legalize operation 'test.illegal_op_g'}}
%0 = "test.illegal_op_g"() : () -> (i32)
- "test.return"(%0) : (i32) -> ()
+ func.return %0 : i32
}
}
diff --git a/mlir/test/Transforms/test-legalizer.mlir b/mlir/test/Transforms/test-legalizer.mlir
index 842d9cfb4a471..676b4fb4662c2 100644
--- a/mlir/test/Transforms/test-legalizer.mlir
+++ b/mlir/test/Transforms/test-legalizer.mlir
@@ -20,7 +20,6 @@
func.func @verifyDirectPattern() -> i32 {
// CHECK-NEXT: "test.legal_op_a"() <{status = "Success"}
%result = "test.illegal_op_a"() : () -> (i32)
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return %result : i32
}
@@ -40,7 +39,6 @@ func.func @verifyDirectPattern() -> i32 {
func.func @verifyLargerBenefit() -> i32 {
// CHECK-NEXT: "test.legal_op_a"() <{status = "Success"}
%result = "test.illegal_op_c"() : () -> (i32)
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return %result : i32
}
@@ -65,7 +63,6 @@ func.func @remap_input_1_to_1(%arg0: i64) {
func.func @remap_call_1_to_1(%arg0: i64) {
// CHECK-NEXT: call @remap_input_1_to_1(%arg0) : (f64) -> ()
call @remap_input_1_to_1(%arg0) : (i64) -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -75,7 +72,7 @@ func.func @remap_call_1_to_1(%arg0: i64) {
// CHECK: notifyBlockInserted into func.func: was unlinked
// Contents of the old block are moved to the new block.
-// CHECK-NEXT: notifyOperationInserted: test.return
+// CHECK-NEXT: notifyOperationInserted: func.return
// The old block is erased.
// CHECK-NEXT: notifyBlockErased
@@ -83,15 +80,9 @@ func.func @remap_call_1_to_1(%arg0: i64) {
// The function op gets a new type attribute.
// CHECK-NEXT: notifyOperationModified: func.func
-// "test.return" is replaced.
-// CHECK-NEXT: notifyOperationInserted: test.return, was unlinked
-// CHECK-NEXT: notifyOperationReplaced: test.return
-// CHECK-NEXT: notifyOperationErased: test.return
-
-// CHECK-LABEL: func @remap_input_1_to_N({{.*}}f16, {{.*}}f16)
-func.func @remap_input_1_to_N(%arg0: f32) -> f32 {
- // CHECK-NEXT: "test.return"{{.*}} : (f16, f16) -> ()
- "test.return"(%arg0) : (f32) -> ()
+// CHECK-LABEL: func @remap_input_1_to_N(%arg0: f16, %arg1: f16)
+func.func @remap_input_1_to_N(%arg0: f32) {
+ func.return
}
// -----
@@ -110,7 +101,7 @@ func.func @remap_materialize_1_to_1(%arg0: i42) {
// CHECK-NEXT: "work"(%[[V]])
// expected-remark at +1 {{op 'work' is not legalizable}}
"work"(%arg0) : (i42) -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -141,7 +132,6 @@ func.func @no_remap_nested() {
// CHECK-NEXT: "test.valid"{{.*}} : (f64, f64)
"test.invalid"(%i0, %i1) : (f64, f64) -> ()
}) : () -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -155,7 +145,6 @@ func.func @remap_drop_region() {
^bb1(%i0: i64, %unused: i16, %i1: i64, %2: f32):
"test.invalid"(%i0, %i1, %2) : (i64, i64, f32) -> ()
}) : () -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -176,7 +165,6 @@ func.func @up_to_date_replacement(%arg: i8) -> i8 {
// CHECK-NEXT: return
%repl_1 = "test.rewrite"(%arg) : (i8) -> i8
%repl_2 = "test.rewrite"(%repl_1) : (i8) -> i8
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return %repl_2 : i8
}
@@ -189,7 +177,6 @@ func.func @remove_foldable_op(%arg0 : i32) -> (i32) {
%0 = "test.op_with_region_fold"(%arg0) ({
"foo.op_with_region_terminator"() : () -> ()
}) : (i32) -> (i32)
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return %0 : i32
}
@@ -202,7 +189,6 @@ func.func @create_block() {
// CHECK: ^{{.*}}(%{{.*}}: i32, %{{.*}}: i32):
"test.create_block"() : () -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -216,7 +202,6 @@ func.func @create_block() {
func.func @bounded_recursion() {
// CHECK: test.recursive_rewrite 0
test.recursive_rewrite 3
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -247,7 +232,7 @@ func.func @replace_block_arg_1_to_n() {
"test.value_replace"(%arg0, %arg1) : (i32, i16) -> ()
"test.return"(%arg0) : (i32) -> ()
}) : () -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -261,9 +246,9 @@ func.func @replace_op_result_1_to_n() -> i32 {
// CHECK-NEXT: %[[cast:.*]] = "test.cast"(%[[repl]], %[[repl]]) : (i16, i16) -> i32
// CHECK-NEXT: "test.value_replace"(%[[cast]], %[[repl]]) {is_legal} : (i32, i16) -> ()
- // CHECK-NEXT: "test.return"(%[[cast]]) : (i32)
+ // CHECK-NEXT: return %[[cast]] : i32
"test.value_replace"(%0, %1) : (i32, i16) -> ()
- "test.return"(%0) : (i32) -> ()
+ func.return %0 : i32
}
// -----
@@ -274,7 +259,6 @@ func.func @replace_op_result_1_to_n() -> i32 {
func.func @blackhole() {
%input = "test.blackhole_producer"() : () -> (i32)
"test.blackhole"(%input) : (i32) -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -296,7 +280,7 @@ func.func @caller() {
// CHECK: "test.some_user"(%[[cast0]], %[[cast1]]) : (f32, i24) -> ()
// expected-remark @below{{'test.some_user' is not legalizable}}
"test.some_user"(%0#0, %0#1) : (f32, i24) -> ()
- "test.return"() : () -> ()
+ func.return
}
}
@@ -319,7 +303,7 @@ func.func @fold_legalization() -> i32 {
// CHECK: op_in_place_self_fold
// CHECK-SAME: folded
%1 = "test.op_in_place_self_fold"() : () -> (i32)
- "test.return"(%1) : (i32) -> ()
+ func.return %1 : i32
}
// -----
@@ -334,7 +318,7 @@ func.func @convert_detached_signature() {
^bb0(%arg0: i64):
"test.return"() : () -> ()
}) : () -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -374,7 +358,7 @@ func.func @test_duplicate_block_arg() {
^bb0(%arg0: i64):
"test.repetitive_1_to_n_consumer"(%arg0) : (i64) -> ()
} : () -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -390,7 +374,7 @@ func.func @test_remap_block_arg() {
^bb0(%arg0: i32):
"test.repetitive_1_to_n_consumer"(%arg0) : (i32) -> ()
} : (i32) -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -426,7 +410,6 @@ func.func @test_lookup_without_converter() {
// Make sure that the second "replace_with_valid_consumer" lowering does not
// lookup the materialization that was created for the above op.
"test.replace_with_valid_consumer"(%0) : (i64) -> ()
- // expected-remark at +1 {{op 'func.return' is not legalizable}}
return
}
@@ -446,9 +429,9 @@ func.func @test_skip_1to1_pattern(%arg0: f32) {
// CHECK-LABEL: @test_working_1to1_pattern(
func.func @test_working_1to1_pattern(%arg0: f16) {
- // CHECK-NEXT: "test.return"() : () -> ()
+ // CHECK-NEXT: return
"test.type_consumer"(%arg0) : (f16) -> ()
- "test.return"() : () -> ()
+ func.return
}
// -----
@@ -487,6 +470,5 @@ func.func @test_preorder_legalization() {
}) : () -> ()
"test.invalid"(%arg0) : (i64) -> ()
}) : () -> ()
- // expected-remark @+1 {{'func.return' is not legalizable}}
return
}
diff --git a/mlir/test/Transforms/test-merge-blocks.mlir b/mlir/test/Transforms/test-merge-blocks.mlir
index ec5af7fb93212..f9897ac6d9305 100644
--- a/mlir/test/Transforms/test-merge-blocks.mlir
+++ b/mlir/test/Transforms/test-merge-blocks.mlir
@@ -5,14 +5,14 @@ func.func @merge_blocks(%arg0: i32, %arg1 : i32) -> (i32, i32) {
// CHECK: "test.merge_blocks"() ({
// CHECK-NEXT: "test.return"
// CHECK-NEXT: })
- // CHECK-NEXT: "test.return"
+ // CHECK-NEXT: return
%0:2 = "test.merge_blocks"() ({
^bb0:
"test.br"(%arg0, %arg1)[^bb1] : (i32, i32) -> ()
^bb1(%arg3 : i32, %arg4 : i32):
"test.return"(%arg3, %arg4) : (i32, i32) -> ()
}) : () -> (i32, i32)
- "test.return"(%0#0, %0#1) : (i32, i32) -> ()
+ func.return %0#0, %0#1 : i32, i32
}
// -----
@@ -55,5 +55,5 @@ func.func @inline_regions() -> ()
}) : () -> ()
"test.finish"() : () -> ()
}) : () -> ()
- "test.return"() : () -> ()
+ func.return
}
diff --git a/mlir/test/Transforms/test-pattern-selective-replacement.mlir b/mlir/test/Transforms/test-pattern-selective-replacement.mlir
index 2bab0b552d311..4b009bf9abdd9 100644
--- a/mlir/test/Transforms/test-pattern-selective-replacement.mlir
+++ b/mlir/test/Transforms/test-pattern-selective-replacement.mlir
@@ -6,10 +6,10 @@
// CHECK-SAME: %[[ARG0:.*]]: i32, %[[ARG1:.*]]: i32
func.func @test1(%arg0: i32, %arg1 : i32) -> (i32, i32) {
// CHECK: arith.addi %[[ARG1]], %[[ARG1]]
- // CHECK-NEXT: "test.return"(%[[ARG0]]
+ // CHECK-NEXT: return %[[ARG0]]
%cast = "test.cast"(%arg0, %arg1) : (i32, i32) -> (i32)
%non_terminator = arith.addi %cast, %cast : i32
- "test.return"(%cast, %non_terminator) : (i32, i32) -> ()
+ func.return %cast, %non_terminator : i32, i32
}
// -----
diff --git a/mlir/test/lib/Dialect/Test/TestPatterns.cpp b/mlir/test/lib/Dialect/Test/TestPatterns.cpp
index 6c44ace831e96..fc237d18df74a 100644
--- a/mlir/test/lib/Dialect/Test/TestPatterns.cpp
+++ b/mlir/test/lib/Dialect/Test/TestPatterns.cpp
@@ -1145,7 +1145,7 @@ struct TestDropAndReplaceInvalidOp : public ConversionPattern {
return success();
}
};
-/// This pattern handles the case of a split return value.
+/// This pattern handles the case of a split return value (test.return variant).
struct TestSplitReturnType : public ConversionPattern {
TestSplitReturnType(MLIRContext *ctx)
: ConversionPattern("test.return", 1, ctx) {}
@@ -1160,6 +1160,25 @@ struct TestSplitReturnType : public ConversionPattern {
}
};
+/// This pattern handles func.return when the operand type is split (1:N
+/// mapping). E.g., func.return %f32 -> func.return %f16_a, %f16_b.
+struct TestSplitFuncReturnType : public OpConversionPattern<func::ReturnOp> {
+ using OpConversionPattern<func::ReturnOp>::OpConversionPattern;
+ LogicalResult
+ matchAndRewrite(func::ReturnOp op, OneToNOpAdaptor adaptor,
+ ConversionPatternRewriter &rewriter) const final {
+ // Only handle the case where a single F32 operand has been split.
+ if (op->getNumOperands() != 1 || !op->getOperand(0).getType().isF32())
+ return failure();
+ // adaptor.getOperands()[0] is the 1:N split of the original F32 value.
+ ValueRange splitOperands = adaptor.getOperands()[0];
+ if (splitOperands.size() != 2)
+ return failure();
+ rewriter.replaceOpWithNewOp<func::ReturnOp>(op, splitOperands);
+ return success();
+ }
+};
+
//===----------------------------------------------------------------------===//
// Multi-Level Type-Conversion Rewrite Testing
struct TestChangeProducerTypeI32ToF32 : public ConversionPattern {
@@ -1583,6 +1602,7 @@ struct TestLegalizePatternDriver
TestUndoPropertiesModification, TestUndoMoveDetachedBlock, TestEraseOp,
TestReplaceWithValidProducer, TestReplaceWithValidConsumer,
TestRepetitive1ToNConsumer>(&getContext());
+ patterns.add<TestSplitFuncReturnType>(converter, &getContext());
patterns.add<TestDropOpSignatureConversion, TestDropAndReplaceInvalidOp,
TestPassthroughInvalidOp, TestMultiple1ToNReplacement,
TestValueReplace, TestReplaceWithValidConsumer,
@@ -1607,6 +1627,11 @@ struct TestLegalizePatternDriver
return llvm::none_of(op.getOperandTypes(),
[](Type type) { return type.isF32(); });
});
+ target.addDynamicallyLegalOp<func::ReturnOp>([](func::ReturnOp op) {
+ // Don't allow F32 operands (mirroring TestReturnOp rule).
+ return llvm::none_of(op.getOperandTypes(),
+ [](Type type) { return type.isF32(); });
+ });
target.addDynamicallyLegalOp<func::FuncOp>([&](func::FuncOp op) {
return converter.isSignatureLegal(op.getFunctionType());
});
@@ -1835,7 +1860,7 @@ struct TestRemappedValue
patterns.add<TestRemapValueInRegion>(typeConverter, &getContext());
mlir::ConversionTarget target(getContext());
- target.addLegalOp<ModuleOp, func::FuncOp, TestReturnOp>();
+ target.addLegalOp<ModuleOp, func::FuncOp, func::ReturnOp, TestReturnOp>();
// Expect the type_producer/type_consumer operations to only operate on f64.
target.addDynamicallyLegalOp<TestTypeProducerOp>(
@@ -2389,8 +2414,9 @@ struct TestMergeBlocksPatternDriver
patterns.add<TestMergeBlock, TestUndoBlocksMerge, TestMergeSingleBlockOps>(
context);
ConversionTarget target(*context);
- target.addLegalOp<func::FuncOp, ModuleOp, TerminatorOp, TestBranchOp,
- TestTypeConsumerOp, TestTypeProducerOp, TestReturnOp>();
+ target.addLegalOp<func::FuncOp, func::ReturnOp, ModuleOp, TerminatorOp,
+ TestBranchOp, TestTypeConsumerOp, TestTypeProducerOp,
+ TestReturnOp>();
target.addIllegalOp<ILLegalOpF>();
/// Expect the op to have a single block after legalization.
More information about the Mlir-commits
mailing list