[Mlir-commits] [mlir] afcceec - [mlir][LLVM] Don't make `LLVM_IntPtrBase` a `BuildableType` to allow the use of opaque pointers
Markus Böck
llvmlistbot at llvm.org
Tue Feb 21 11:46:48 PST 2023
Author: Markus Böck
Date: 2023-02-21T20:46:33+01:00
New Revision: afcceec744097a22037da3f426d1088b360ba164
URL: https://github.com/llvm/llvm-project/commit/afcceec744097a22037da3f426d1088b360ba164
DIFF: https://github.com/llvm/llvm-project/commit/afcceec744097a22037da3f426d1088b360ba164.diff
LOG: [mlir][LLVM] Don't make `LLVM_IntPtrBase` a `BuildableType` to allow the use of opaque pointers
Making the constraint a buildable type makes them incompatible with opaque pointers, at least while we still support typed pointers, since Ops making use of the constraint will then automatically create a typed pointer on parse.
This patch therefore fixes that issue by removing the `BuildableType` mixin. This has a bit of a cascading effect however, as all users of the constraint now need operands of that type to be added to the assembly format, hence a lot of adjustments to the syntax of a lot of (mostly intrinsic) ops.
Few things of note: The syntax as is, is only required while we're supporting both typed and opaque pointers. Once we drop support for typed pointers, we can make it a `BuildableType` again. As a drive by I also fixed the address space not being verified in the constraint. Finally, I added some roundtripping tests, most importantly for ops with `type($specific_operand)` occurences. These are printed incorrectly with typed pointers if not wrapped within a `qualified`.
Differential Revision: https://reviews.llvm.org/D144479
Added:
Modified:
mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir
mlir/test/Dialect/LLVMIR/callgraph.mlir
mlir/test/Dialect/LLVMIR/invalid.mlir
mlir/test/Dialect/LLVMIR/nvvm.mlir
mlir/test/Dialect/LLVMIR/roundtrip.mlir
mlir/test/Target/LLVMIR/Import/intrinsic.ll
mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
mlir/test/Target/LLVMIR/llvmir.mlir
mlir/test/Target/LLVMIR/nvvmir.mlir
mlir/test/mlir-cpu-runner/x86-varargs.mlir
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
index 76e5b34899168..08aca3886ba25 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
@@ -217,13 +217,13 @@ def LLVM_CoroIdOp : LLVM_IntrOp<"coro.id", [], [], [], 1> {
LLVM_i8Ptr:$coroaddr,
LLVM_i8Ptr:$fnaddrs);
let assemblyFormat = "$align `,` $promise `,` $coroaddr `,` $fnaddrs"
- " attr-dict `:` type($res)";
+ " attr-dict `:` functional-type(operands, results)";
}
def LLVM_CoroBeginOp : LLVM_IntrOp<"coro.begin", [], [], [], 1> {
let arguments = (ins LLVM_TokenType:$token,
LLVM_i8Ptr:$mem);
- let assemblyFormat = "$token `,` $mem attr-dict `:` type($res)";
+ let assemblyFormat = "$token `,` $mem attr-dict `:` functional-type(operands, results)";
}
def LLVM_CoroSizeOp : LLVM_IntrOp<"coro.size", [0], [], [], 1> {
@@ -236,7 +236,7 @@ def LLVM_CoroAlignOp : LLVM_IntrOp<"coro.align", [0], [], [], 1> {
def LLVM_CoroSaveOp : LLVM_IntrOp<"coro.save", [], [], [], 1> {
let arguments = (ins LLVM_i8Ptr:$handle);
- let assemblyFormat = "$handle attr-dict `:` type($res)";
+ let assemblyFormat = "$handle attr-dict `:` functional-type(operands, results)";
}
def LLVM_CoroSuspendOp : LLVM_IntrOp<"coro.suspend", [], [], [], 1> {
@@ -248,18 +248,18 @@ def LLVM_CoroSuspendOp : LLVM_IntrOp<"coro.suspend", [], [], [], 1> {
def LLVM_CoroEndOp : LLVM_IntrOp<"coro.end", [], [], [], 1> {
let arguments = (ins LLVM_i8Ptr:$handle,
I1:$unwind);
- let assemblyFormat = "$handle `,` $unwind attr-dict `:` type($res)";
+ let assemblyFormat = "$handle `,` $unwind attr-dict `:` functional-type(operands, results)";
}
def LLVM_CoroFreeOp : LLVM_IntrOp<"coro.free", [], [], [], 1> {
let arguments = (ins LLVM_TokenType:$id,
LLVM_i8Ptr:$handle);
- let assemblyFormat = "$id `,` $handle attr-dict `:` type($res)";
+ let assemblyFormat = "$id `,` $handle attr-dict `:` functional-type(operands, results)";
}
def LLVM_CoroResumeOp : LLVM_IntrOp<"coro.resume", [], [], [], 0> {
let arguments = (ins LLVM_i8Ptr:$handle);
- let assemblyFormat = "$handle attr-dict";
+ let assemblyFormat = "$handle attr-dict `:` qualified(type($handle))";
}
//
@@ -328,19 +328,19 @@ def LLVM_DbgValueOp : LLVM_DbgIntrOp<"dbg.value", "value"> {
def LLVM_VaStartOp : LLVM_ZeroResultIntrOp<"vastart">,
Arguments<(ins LLVM_i8Ptr:$arg_list)> {
- let assemblyFormat = "$arg_list attr-dict";
+ let assemblyFormat = "$arg_list attr-dict `:` qualified(type($arg_list))";
let summary = "Initializes `arg_list` for subsequent variadic argument extractions.";
}
def LLVM_VaCopyOp : LLVM_ZeroResultIntrOp<"vacopy">,
Arguments<(ins LLVM_i8Ptr:$dest_list, LLVM_i8Ptr:$src_list)> {
- let assemblyFormat = "$src_list `to` $dest_list attr-dict";
+ let assemblyFormat = "$src_list `to` $dest_list attr-dict `:` type(operands)";
let summary = "Copies the current argument position from `src_list` to `dest_list`.";
}
def LLVM_VaEndOp : LLVM_ZeroResultIntrOp<"vaend">,
Arguments<(ins LLVM_i8Ptr:$arg_list)> {
- let assemblyFormat = "$arg_list attr-dict";
+ let assemblyFormat = "$arg_list attr-dict `:` qualified(type($arg_list))";
let summary = "Destroys `arg_list`, which has been initialized by `intr.vastart` or `intr.vacopy`.";
}
@@ -350,7 +350,7 @@ def LLVM_VaEndOp : LLVM_ZeroResultIntrOp<"vaend">,
def LLVM_EhTypeidForOp : LLVM_OneResultIntrOp<"eh.typeid.for"> {
let arguments = (ins LLVM_i8Ptr:$type_info);
- let assemblyFormat = "$type_info attr-dict `:` type($res)";
+ let assemblyFormat = "$type_info attr-dict `:` functional-type(operands, results)";
}
//
@@ -363,7 +363,7 @@ def LLVM_StackSaveOp : LLVM_OneResultIntrOp<"stacksave"> {
def LLVM_StackRestoreOp : LLVM_ZeroResultIntrOp<"stackrestore"> {
let arguments = (ins LLVM_i8Ptr:$ptr);
- let assemblyFormat = "$ptr attr-dict";
+ let assemblyFormat = "$ptr attr-dict `:` qualified(type($ptr))";
}
//
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index 6d04cd3d235fb..40455d7ab49d8 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -158,11 +158,10 @@ class LLVM_PointerTo<Type pointee> : Type<
// Type constraints accepting LLVM pointer type to integer of a specific width.
class LLVM_IntPtrBase<int width, int addressSpace = 0> : Type<
- LLVM_PointerTo<I<width>>.predicate,
- "LLVM pointer to " # I<width>.summary>,
- BuildableType<"::mlir::LLVM::LLVMPointerType::get("
- "::mlir::IntegerType::get($_builder.getContext(), "
- # width #"), "# addressSpace #")">;
+ And<[LLVM_PointerTo<I<width>>.predicate,
+ CPred<"$_self.cast<::mlir::LLVM::LLVMPointerType>().getAddressSpace()"
+ " == " # addressSpace>]>,
+ "LLVM pointer to " # I<width>.summary>;
def LLVM_i8Ptr : LLVM_IntPtrBase<8>;
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 32f826d76be51..5dcd5f965ce31 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -265,7 +265,7 @@ def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global">,
}
createIntrinsicCall(builder, id, {$dst, $src});
}];
- let assemblyFormat = "$dst `,` $src `,` $size attr-dict";
+ let assemblyFormat = "$dst `,` $src `,` $size attr-dict `:` type(operands)";
let hasVerifier = 1;
}
diff --git a/mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir b/mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir
index ad5595a0c1db3..fb15a2da2836b 100644
--- a/mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir
+++ b/mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir
@@ -4,7 +4,7 @@
func.func @coro_id() {
// CHECK: %0 = llvm.mlir.constant(0 : i32) : i32
// CHECK: %1 = llvm.mlir.null : !llvm.ptr
- // CHECK: %2 = llvm.intr.coro.id %0, %1, %1, %1 : !llvm.token
+ // CHECK: %2 = llvm.intr.coro.id %0, %1, %1, %1 : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token
%0 = async.coro.id
return
}
diff --git a/mlir/test/Dialect/LLVMIR/callgraph.mlir b/mlir/test/Dialect/LLVMIR/callgraph.mlir
index c61fb2eca16c9..edb5b35d126a5 100644
--- a/mlir/test/Dialect/LLVMIR/callgraph.mlir
+++ b/mlir/test/Dialect/LLVMIR/callgraph.mlir
@@ -77,7 +77,7 @@ module attributes {"test.name" = "Invoke call"} {
^bb1:
%10 = llvm.landingpad cleanup (catch %3 : !llvm.ptr<ptr<i8>>) (catch %6 : !llvm.ptr<i8>) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr<i8>, i32)>
- %11 = llvm.intr.eh.typeid.for %6 : i32
+ %11 = llvm.intr.eh.typeid.for %6 : (!llvm.ptr<i8>) -> i32
llvm.resume %10 : !llvm.struct<(ptr<i8>, i32)>
^bb2:
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index 642639b18deec..d063e9b7a5c2e 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -1280,7 +1280,7 @@ func.func @bitcast(%arg0: vector<2x3xf32>) {
func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
// expected-error @below {{expected byte size to be either 4, 8 or 16.}}
- nvvm.cp.async.shared.global %arg0, %arg1, 32
+ nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
return
}
@@ -1288,7 +1288,7 @@ func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
// expected-error @below {{bypass l1 is only support for 16 bytes copy.}}
- nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1}
+ nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
return
}
diff --git a/mlir/test/Dialect/LLVMIR/nvvm.mlir b/mlir/test/Dialect/LLVMIR/nvvm.mlir
index 2e3b20b1d6611..6596b8503d7a5 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir
@@ -290,9 +290,9 @@ func.func @nvvm_wmma_mma(%0 : i32, %1 : i32, %2 : i32, %3 : i32, %4 : i32, %5 :
// CHECK-LABEL: @cp_async
llvm.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16
- nvvm.cp.async.shared.global %arg0, %arg1, 16
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16 {bypass_l1}
- nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1}
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
// CHECK: nvvm.cp.async.commit.group
nvvm.cp.async.commit.group
// CHECK: nvvm.cp.async.wait.group 0
diff --git a/mlir/test/Dialect/LLVMIR/roundtrip.mlir b/mlir/test/Dialect/LLVMIR/roundtrip.mlir
index 1e49b4e3124ec..f7b340fd66ac8 100644
--- a/mlir/test/Dialect/LLVMIR/roundtrip.mlir
+++ b/mlir/test/Dialect/LLVMIR/roundtrip.mlir
@@ -404,11 +404,11 @@ llvm.func @invokeLandingpad() -> i32 attributes { personality = @__gxx_personali
// CHECK: ^[[BB1]]:
// CHECK: %[[lp:.*]] = llvm.landingpad cleanup (catch %[[a3]] : !llvm.ptr<ptr<i8>>) (catch %[[a6]] : !llvm.ptr<i8>) (filter %[[a2]] : !llvm.array<1 x i8>) : !llvm.struct<(ptr<i8>, i32)>
-// CHECK: %{{.*}} = llvm.intr.eh.typeid.for %6 : i32
+// CHECK: %{{.*}} = llvm.intr.eh.typeid.for %6 : (!llvm.ptr<i8>) -> i32
// CHECK: llvm.resume %[[lp]] : !llvm.struct<(ptr<i8>, i32)>
^bb1:
%10 = llvm.landingpad cleanup (catch %3 : !llvm.ptr<ptr<i8>>) (catch %6 : !llvm.ptr<i8>) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr<i8>, i32)>
- %11 = llvm.intr.eh.typeid.for %6 : i32
+ %11 = llvm.intr.eh.typeid.for %6 : (!llvm.ptr<i8>) -> i32
llvm.resume %10 : !llvm.struct<(ptr<i8>, i32)>
// CHECK: ^[[BB2]]:
@@ -530,17 +530,17 @@ llvm.func @vararg_func(%arg0: i32, ...) {
%2 = llvm.alloca %1 x !llvm.struct<"struct.va_list", (ptr<i8>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>>
%3 = llvm.bitcast %2 : !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>> to !llvm.ptr<i8>
// CHECK: llvm.intr.vastart %[[CAST0]]
- llvm.intr.vastart %3
+ llvm.intr.vastart %3 : !llvm.ptr<i8>
// CHECK: %[[ALLOCA1:.+]] = llvm.alloca %{{.*}} x !llvm.ptr<i8> {alignment = 8 : i64} : (i32) -> !llvm.ptr<ptr<i8>>
// CHECK: %[[CAST1:.+]] = llvm.bitcast %[[ALLOCA1]] : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
%4 = llvm.alloca %0 x !llvm.ptr<i8> {alignment = 8 : i64} : (i32) -> !llvm.ptr<ptr<i8>>
%5 = llvm.bitcast %4 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
// CHECK: llvm.intr.vacopy %[[CAST0]] to %[[CAST1]]
- llvm.intr.vacopy %3 to %5
+ llvm.intr.vacopy %3 to %5 : !llvm.ptr<i8>, !llvm.ptr<i8>
// CHECK: llvm.intr.vaend %[[CAST1]]
// CHECK: llvm.intr.vaend %[[CAST0]]
- llvm.intr.vaend %5
- llvm.intr.vaend %3
+ llvm.intr.vaend %5 : !llvm.ptr<i8>
+ llvm.intr.vaend %3 : !llvm.ptr<i8>
// CHECK: llvm.return
llvm.return
}
@@ -554,3 +554,40 @@ llvm.func @lifetime(%p: !llvm.ptr) {
llvm.intr.lifetime.end 16, %p : !llvm.ptr
llvm.return
}
+
+// CHECK-LABEL: @vararg_func_opaque_pointers
+llvm.func @vararg_func_opaque_pointers(%arg0: i32, ...) {
+ // CHECK: %[[C:.*]] = llvm.mlir.constant(1 : i32)
+ // CHECK: %[[LIST:.*]] = llvm.alloca
+ // CHECK: llvm.intr.vastart %[[LIST]] : !llvm.ptr{{$}}
+ %1 = llvm.mlir.constant(1 : i32) : i32
+ %list = llvm.alloca %1 x !llvm.struct<"struct.va_list_opaque", (ptr)> : (i32) -> !llvm.ptr
+ llvm.intr.vastart %list : !llvm.ptr
+
+ // CHECK: %[[LIST2:.*]] = llvm.alloca
+ // CHECK: llvm.intr.vacopy %[[LIST]] to %[[LIST2]] : !llvm.ptr, !llvm.ptr{{$}}
+ %list2 = llvm.alloca %1 x !llvm.struct<"struct.va_list_opaque", (ptr)> : (i32) -> !llvm.ptr
+ llvm.intr.vacopy %list to %list2 : !llvm.ptr, !llvm.ptr
+
+ // CHECK: llvm.intr.vaend %[[LIST]] : !llvm.ptr{{$}}
+ // CHECK: llvm.intr.vaend %[[LIST2]] : !llvm.ptr{{$}}
+ llvm.intr.vaend %list : !llvm.ptr
+ llvm.intr.vaend %list2 : !llvm.ptr
+ llvm.return
+}
+
+// CHECK-LABEL: @eh_typeid_opaque_pointers
+// CHECK-SAME: %[[ARG0:.*]]: !llvm.ptr
+llvm.func @eh_typeid_opaque_pointers(%arg0: !llvm.ptr) -> i32 {
+ // CHECK: llvm.intr.eh.typeid.for %[[ARG0]] : (!llvm.ptr) -> i32
+ %0 = llvm.intr.eh.typeid.for %arg0 : (!llvm.ptr) -> i32
+ llvm.return %0 : i32
+}
+
+// CHECK-LABEL: @stackrestore_opaque_pointers
+// CHECK-SAME: %[[ARG0:.*]]: !llvm.ptr
+llvm.func @stackrestore_opaque_pointers(%arg0: !llvm.ptr) {
+ // CHECK: llvm.intr.stackrestore %[[ARG0]] : !llvm.ptr
+ llvm.intr.stackrestore %arg0 : !llvm.ptr
+ llvm.return
+}
diff --git a/mlir/test/Target/LLVMIR/Import/intrinsic.ll b/mlir/test/Target/LLVMIR/Import/intrinsic.ll
index 821a0272e8d6a..987fc0c63e006 100644
--- a/mlir/test/Target/LLVMIR/Import/intrinsic.ll
+++ b/mlir/test/Target/LLVMIR/Import/intrinsic.ll
@@ -431,16 +431,16 @@ define void @assume(i1 %true) {
; CHECK-LABEL: llvm.func @coro_id
define void @coro_id(i32 %0, ptr %1) {
- ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : !llvm.token
+ ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token
%3 = call token @llvm.coro.id(i32 %0, ptr %1, ptr %1, ptr null)
ret void
}
; CHECK-LABEL: llvm.func @coro_begin
define void @coro_begin(i32 %0, ptr %1) {
- ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : !llvm.token
+ ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token
%3 = call token @llvm.coro.id(i32 %0, ptr %1, ptr %1, ptr null)
- ; CHECK: llvm.intr.coro.begin %{{.*}}, %{{.*}} : !llvm.ptr
+ ; CHECK: llvm.intr.coro.begin %{{.*}}, %{{.*}} : (!llvm.token, !llvm.ptr) -> !llvm.ptr
%4 = call ptr @llvm.coro.begin(token %3, ptr %1)
ret void
}
@@ -464,14 +464,14 @@ define void @coro_align() {
; CHECK-LABEL: llvm.func @coro_save
define void @coro_save(ptr %0) {
- ; CHECK: llvm.intr.coro.save %{{.*}} : !llvm.token
+ ; CHECK: llvm.intr.coro.save %{{.*}} : (!llvm.ptr) -> !llvm.token
%2 = call token @llvm.coro.save(ptr %0)
ret void
}
; CHECK-LABEL: llvm.func @coro_suspend
define void @coro_suspend(i32 %0, i1 %1, ptr %2) {
- ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : !llvm.token
+ ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token
%4 = call token @llvm.coro.id(i32 %0, ptr %2, ptr %2, ptr null)
; CHECK: llvm.intr.coro.suspend %{{.*}}, %{{.*}} : i8
%5 = call i8 @llvm.coro.suspend(token %4, i1 %1)
@@ -487,9 +487,9 @@ define void @coro_end(ptr %0, i1 %1) {
; CHECK-LABEL: llvm.func @coro_free
define void @coro_free(i32 %0, ptr %1) {
- ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : !llvm.token
+ ; CHECK: llvm.intr.coro.id %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token
%3 = call token @llvm.coro.id(i32 %0, ptr %1, ptr %1, ptr null)
- ; CHECK: llvm.intr.coro.free %{{.*}}, %{{.*}} : !llvm.ptr
+ ; CHECK: llvm.intr.coro.free %{{.*}}, %{{.*}} : (!llvm.token, !llvm.ptr) -> !llvm.ptr
%4 = call ptr @llvm.coro.free(token %3, ptr %1)
ret void
}
@@ -503,7 +503,7 @@ define void @coro_resume(ptr %0) {
; CHECK-LABEL: llvm.func @eh_typeid_for
define void @eh_typeid_for(ptr %0) {
- ; CHECK: llvm.intr.eh.typeid.for %{{.*}} : i32
+ ; CHECK: llvm.intr.eh.typeid.for %{{.*}} : (!llvm.ptr) -> i32
%2 = call i32 @llvm.eh.typeid.for(ptr %0)
ret void
}
diff --git a/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir b/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
index 3dfb0574b9075..681ba0c2e5485 100644
--- a/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
+++ b/mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir
@@ -432,16 +432,16 @@ llvm.func @umul_with_overflow_test(%arg0: i32, %arg1: i32, %arg2: vector<8xi32>,
llvm.func @coro_id(%arg0: i32, %arg1: !llvm.ptr<i8>) {
// CHECK: call token @llvm.coro.id
%null = llvm.mlir.null : !llvm.ptr<i8>
- llvm.intr.coro.id %arg0, %arg1, %arg1, %null : !llvm.token
+ llvm.intr.coro.id %arg0, %arg1, %arg1, %null : (i32, !llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>) -> !llvm.token
llvm.return
}
// CHECK-LABEL: @coro_begin
llvm.func @coro_begin(%arg0: i32, %arg1: !llvm.ptr<i8>) {
%null = llvm.mlir.null : !llvm.ptr<i8>
- %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : !llvm.token
+ %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : (i32, !llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>) -> !llvm.token
// CHECK: call ptr @llvm.coro.begin
- llvm.intr.coro.begin %token, %arg1 : !llvm.ptr<i8>
+ llvm.intr.coro.begin %token, %arg1 : (!llvm.token, !llvm.ptr<i8>) -> !llvm.ptr<i8>
llvm.return
}
@@ -466,14 +466,14 @@ llvm.func @coro_align() {
// CHECK-LABEL: @coro_save
llvm.func @coro_save(%arg0: !llvm.ptr<i8>) {
// CHECK: call token @llvm.coro.save
- %0 = llvm.intr.coro.save %arg0 : !llvm.token
+ %0 = llvm.intr.coro.save %arg0 : (!llvm.ptr<i8>) -> !llvm.token
llvm.return
}
// CHECK-LABEL: @coro_suspend
llvm.func @coro_suspend(%arg0: i32, %arg1 : i1, %arg2 : !llvm.ptr<i8>) {
%null = llvm.mlir.null : !llvm.ptr<i8>
- %token = llvm.intr.coro.id %arg0, %arg2, %arg2, %null : !llvm.token
+ %token = llvm.intr.coro.id %arg0, %arg2, %arg2, %null : (i32, !llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>) -> !llvm.token
// CHECK: call i8 @llvm.coro.suspend
%0 = llvm.intr.coro.suspend %token, %arg1 : i8
llvm.return
@@ -482,30 +482,30 @@ llvm.func @coro_suspend(%arg0: i32, %arg1 : i1, %arg2 : !llvm.ptr<i8>) {
// CHECK-LABEL: @coro_end
llvm.func @coro_end(%arg0: !llvm.ptr<i8>, %arg1 : i1) {
// CHECK: call i1 @llvm.coro.end
- %0 = llvm.intr.coro.end %arg0, %arg1 : i1
+ %0 = llvm.intr.coro.end %arg0, %arg1 : (!llvm.ptr<i8>, i1) -> i1
llvm.return
}
// CHECK-LABEL: @coro_free
llvm.func @coro_free(%arg0: i32, %arg1 : !llvm.ptr<i8>) {
%null = llvm.mlir.null : !llvm.ptr<i8>
- %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : !llvm.token
+ %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : (i32, !llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>) -> !llvm.token
// CHECK: call ptr @llvm.coro.free
- %0 = llvm.intr.coro.free %token, %arg1 : !llvm.ptr<i8>
+ %0 = llvm.intr.coro.free %token, %arg1 : (!llvm.token, !llvm.ptr<i8>) -> !llvm.ptr<i8>
llvm.return
}
// CHECK-LABEL: @coro_resume
llvm.func @coro_resume(%arg0: !llvm.ptr<i8>) {
// CHECK: call void @llvm.coro.resume
- llvm.intr.coro.resume %arg0
+ llvm.intr.coro.resume %arg0 : !llvm.ptr<i8>
llvm.return
}
// CHECK-LABEL: @eh_typeid_for
llvm.func @eh_typeid_for(%arg0 : !llvm.ptr<i8>) {
// CHECK: call i32 @llvm.eh.typeid.for
- %0 = llvm.intr.eh.typeid.for %arg0 : i32
+ %0 = llvm.intr.eh.typeid.for %arg0 : (!llvm.ptr<i8>) -> i32
llvm.return
}
@@ -519,7 +519,7 @@ llvm.func @stack_save() {
// CHECK-LABEL: @stack_restore
llvm.func @stack_restore(%arg0: !llvm.ptr<i8>) {
// CHECK: call void @llvm.stackrestore
- llvm.intr.stackrestore %arg0
+ llvm.intr.stackrestore %arg0 : !llvm.ptr<i8>
llvm.return
}
diff --git a/mlir/test/Target/LLVMIR/llvmir.mlir b/mlir/test/Target/LLVMIR/llvmir.mlir
index a58dfde99b053..f897429e563e1 100644
--- a/mlir/test/Target/LLVMIR/llvmir.mlir
+++ b/mlir/test/Target/LLVMIR/llvmir.mlir
@@ -2120,16 +2120,16 @@ llvm.func @vararg_function(%arg0: i32, ...) {
%2 = llvm.alloca %1 x !llvm.struct<"struct.va_list", (ptr<i8>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>>
%3 = llvm.bitcast %2 : !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>> to !llvm.ptr<i8>
// CHECK: call void @llvm.va_start(ptr %[[ALLOCA0]])
- llvm.intr.vastart %3
+ llvm.intr.vastart %3 : !llvm.ptr<i8>
// CHECK: %[[ALLOCA1:.+]] = alloca ptr, align 8
%4 = llvm.alloca %0 x !llvm.ptr<i8> {alignment = 8 : i64} : (i32) -> !llvm.ptr<ptr<i8>>
%5 = llvm.bitcast %4 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
// CHECK: call void @llvm.va_copy(ptr %[[ALLOCA1]], ptr %[[ALLOCA0]])
- llvm.intr.vacopy %3 to %5
+ llvm.intr.vacopy %3 to %5 : !llvm.ptr<i8>, !llvm.ptr<i8>
// CHECK: call void @llvm.va_end(ptr %[[ALLOCA1]])
// CHECK: call void @llvm.va_end(ptr %[[ALLOCA0]])
- llvm.intr.vaend %5
- llvm.intr.vaend %3
+ llvm.intr.vaend %5 : !llvm.ptr<i8>
+ llvm.intr.vaend %3 : !llvm.ptr<i8>
// CHECK: ret void
llvm.return
}
diff --git a/mlir/test/Target/LLVMIR/nvvmir.mlir b/mlir/test/Target/LLVMIR/nvvmir.mlir
index 4ead9dc5317ee..522cce57a88a6 100644
--- a/mlir/test/Target/LLVMIR/nvvmir.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir.mlir
@@ -309,13 +309,13 @@ llvm.func @nvvm_wmma_mma(%0 : i32, %1 : i32, %2 : i32, %3 : i32, %4 : i32, %5 :
// CHECK-LABEL: @cp_async
llvm.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
// CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})
- nvvm.cp.async.shared.global %arg0, %arg1, 4
+ nvvm.cp.async.shared.global %arg0, %arg1, 4 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
// CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})
- nvvm.cp.async.shared.global %arg0, %arg1, 8
+ nvvm.cp.async.shared.global %arg0, %arg1, 8 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
// CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})
- nvvm.cp.async.shared.global %arg0, %arg1, 16
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
// CHECK: call void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})
- nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1}
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
// CHECK: call void @llvm.nvvm.cp.async.commit.group()
nvvm.cp.async.commit.group
// CHECK: call void @llvm.nvvm.cp.async.wait.group(i32 0)
diff --git a/mlir/test/mlir-cpu-runner/x86-varargs.mlir b/mlir/test/mlir-cpu-runner/x86-varargs.mlir
index d3a2297b2c42d..44024113c2b7b 100644
--- a/mlir/test/mlir-cpu-runner/x86-varargs.mlir
+++ b/mlir/test/mlir-cpu-runner/x86-varargs.mlir
@@ -38,7 +38,7 @@ llvm.func @foo(%arg0: i32, ...) -> i32 {
%12 = llvm.mlir.constant(1 : i32) : i32
%13 = llvm.alloca %12 x !llvm.array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>> {alignment = 8 : i64} : (i32) -> !llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>
%14 = llvm.bitcast %13 : !llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>> to !llvm.ptr<i8>
- llvm.intr.vastart %14
+ llvm.intr.vastart %14 : !llvm.ptr<i8>
%15 = llvm.getelementptr %13[%11, %10, 0] : (!llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>, i64, i64) -> !llvm.ptr<i32>
%16 = llvm.load %15 : !llvm.ptr<i32>
%17 = llvm.icmp "ult" %16, %8 : i32
@@ -60,7 +60,7 @@ llvm.func @foo(%arg0: i32, ...) -> i32 {
^bb3(%26: !llvm.ptr<i8>): // 2 preds: ^bb1, ^bb2
%27 = llvm.bitcast %26 : !llvm.ptr<i8> to !llvm.ptr<i32>
%28 = llvm.load %27 : !llvm.ptr<i32>
- llvm.intr.vaend %14
+ llvm.intr.vaend %14 : !llvm.ptr<i8>
llvm.return %28 : i32
}
More information about the Mlir-commits
mailing list