[mlir][LLVM] Don't make `LLVM_IntPtrBase` a `BuildableType` to allow the use of opaque pointers

[zero9178]

Summary

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.

[zero9178]

Rebase

[ftynse]

mlir/test/Target/LLVMIR/Import/intrinsic.ll

434

Should the trailing type be updated rather than dropped?

[zero9178]

Address review comments

mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td

 //
 
 def LLVM_CoroIdOp : LLVM_IntrOp<"coro.id", [], [], [], 1> {
   let arguments = (ins I32:$align,
                        LLVM_i8Ptr:$promise,
                        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> {
   let assemblyFormat = "attr-dict `:` type($res)";
 }
 
 def LLVM_CoroAlignOp : LLVM_IntrOp<"coro.align", [0], [], [], 1> {
   let assemblyFormat = "attr-dict `:` type($res)";
 }
 
 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> {
   let arguments = (ins LLVM_TokenType:$save,
                        I1:$final);
   let assemblyFormat = "$save `,` $final attr-dict `:` type($res)";
 }
 
 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))";
 }
 
 //
 // Debug function intrinsics.
 //
 
 class LLVM_DbgIntrOp<string name, string argName> : LLVM_IntrOp<name, [], [], [], 0> {
   let llvmBuilder = [{

 }
 
 //
 // Variadic function intrinsics.
 //
 
 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`.";
 }
 
 //
 // Exception handling intrinsics.
 //
 
 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)";
 }
 
 //
 // Stack save/restore intrinsics.
 //
 
 def LLVM_StackSaveOp : LLVM_OneResultIntrOp<"stacksave"> {
   let assemblyFormat = "attr-dict `:` type($res)";
 }
 
 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))";
 }
 
 //
 // Vector Reductions.
 //
 
 // LLVM vector reduction over a single vector.
 class LLVM_VecReductionBase<string mnem, Type element>

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

And<[LLVM_AnyPointer.predicate,
            SubstLeaves<
              "$_self",
              "$_self.cast<::mlir::LLVM::LLVMPointerType>().getElementType()",
              pointee.predicate>]>]>,
   "LLVM pointer to " # pointee.summary, "::mlir::LLVM::LLVMPointerType">;
 
 // 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>;
 
 // Type constraint accepting any LLVM structure type.
 def LLVM_AnyStruct : Type<CPred<"$_self.isa<::mlir::LLVM::LLVMStructType>()">,
                          "LLVM structure type">;
 
 // Type constraint accepting opaque LLVM structure type.

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

string llvmBuilder = [{
           else
             id = llvm::Intrinsic::nvvm_cp_async_ca_shared_global_16;
           break;
         default:
           llvm_unreachable("unsupported async copy size");
       }
       createIntrinsicCall(builder, id, {$dst, $src});
   }];
-  let assemblyFormat = "$dst `,` $src `,` $size attr-dict";
+  let assemblyFormat = "$dst `,` $src `,` $size attr-dict `:` type(operands)";
   let hasVerifier = 1;
 }
 
 def NVVM_CpAsyncCommitGroupOp : NVVM_Op<"cp.async.commit.group"> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_cp_async_commit_group);
   }];
   let assemblyFormat = "attr-dict";

mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir

 // RUN: mlir-opt %s -convert-async-to-llvm='use-opaque-pointers=1' | FileCheck %s
 
 // CHECK-LABEL: @coro_id
 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
 }
 
 // CHECK-LABEL: @coro_begin
 func.func @coro_begin() {
   // CHECK: %[[ID:.*]] = llvm.intr.coro.id
   %0 = async.coro.id

mlir/test/Dialect/LLVMIR/callgraph.mlir

llvm.func @invokeLandingpad() -> i32 attributes { personality = @__gxx_personality_v0 } {
     %5 = llvm.mlir.addressof @_ZTIi : !llvm.ptr<ptr<i8>>
     %6 = llvm.bitcast %5 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
     %7 = llvm.mlir.constant(1 : i32) : i32
     %8 = llvm.alloca %7 x i8 : (i32) -> !llvm.ptr<i8>
     %9 = llvm.invoke @foo(%7) to ^bb2 unwind ^bb1 : (i32) -> !llvm.struct<(i32, f64, 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)>
 
   ^bb2:
     llvm.return %7 : i32
 
   ^bb3:
     llvm.invoke @bar(%8, %6, %4) to ^bb2 unwind ^bb1 : (!llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>) -> ()
 
   ^bb4:
     llvm.return %0 : i32
   }
 }

mlir/test/Dialect/LLVMIR/invalid.mlir

func.func @bitcast(%arg0: vector<2x3xf32>) {
   llvm.bitcast %arg0 : vector<2x3xf32> to vector<2x3xi32>
   return
 }
 
 // -----
 
 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
 }
 
 // -----
 
 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
 }
 
 // -----
 
 func.func @gep_struct_variable(%arg0: !llvm.ptr<struct<(i32)>>, %arg1: i32, %arg2: i32) {
   // expected-error @below {{op expected index 1 indexing a struct to be constant}}
   llvm.getelementptr %arg0[%arg1, %arg1] : (!llvm.ptr<struct<(i32)>>, i32, i32) -> !llvm.ptr<i32>

mlir/test/Dialect/LLVMIR/nvvm.mlir

%r = nvvm.wmma.mma %0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15
     : (i32, i32, i32, i32, i32, i32, i32, i32, f32, f32, f32, f32, f32, f32, f32, f32)
     -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
   llvm.return %r : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
 }
 
 // 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
   nvvm.cp.async.wait.group 0
   llvm.return
 }
 
 // CHECK-LABEL: llvm.func @ld_matrix

mlir/test/Dialect/LLVMIR/roundtrip.mlir

// CHECK: %{{.*}} = llvm.invoke @foo(%[[a7]]) to ^[[BB2:.*]] unwind ^[[BB1:.*]] : (i32) -> !llvm.struct<(i32, f64, i32)>
   %5 = llvm.mlir.addressof @_ZTIi : !llvm.ptr<ptr<i8>>
   %6 = llvm.bitcast %5 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
   %7 = llvm.mlir.constant(1 : i32) : i32
   %8 = llvm.alloca %7 x i8 : (i32) -> !llvm.ptr<i8>
   %9 = llvm.invoke @foo(%7) to ^bb2 unwind ^bb1 : (i32) -> !llvm.struct<(i32, f64, i32)>
 
 // 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]]:
 // CHECK:   llvm.return %[[a7]] : i32
 ^bb2:
   llvm.return %7 : i32
 
 // CHECK: ^[[BB3:.*]]:
llvm.func @vararg_func(%arg0: i32, ...) {
   // CHECK: %{{.*}} = llvm.mlir.constant(1 : i32) : i32
   %0 = llvm.mlir.constant(1 : i32) : i32
   %1 = llvm.mlir.constant(1 : i32) : i32
   // CHECK: %[[ALLOCA0:.+]] = llvm.alloca %{{.*}} x !llvm.struct<"struct.va_list", (ptr<i8>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>>
   // CHECK: %[[CAST0:.+]] = llvm.bitcast %[[ALLOCA0]] : !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>> to !llvm.ptr<i8>
   %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
 }
 
 // CHECK-LABEL: @lifetime
 // CHECK-SAME: %[[P:.*]]: !llvm.ptr
 llvm.func @lifetime(%p: !llvm.ptr) {
   // CHECK: llvm.intr.lifetime.start 16, %[[P]]
   llvm.intr.lifetime.start 16, %p : !llvm.ptr
   // CHECK: llvm.intr.lifetime.end 16, %[[P]]
   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
+}

mlir/test/Target/LLVMIR/Import/intrinsic.ll

 define void @assume(i1 %true) {
   ; CHECK:  "llvm.intr.assume"(%[[TRUE]]) : (i1) -> ()
   call void @llvm.assume(i1 %true)
   ret void
 }
 
 ; 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
 }
 
 ; CHECK-LABEL:  llvm.func @coro_size()
 define void @coro_size() {
   ; CHECK: llvm.intr.coro.size : i64
   %1 = call i64 @llvm.coro.size.i64()
   ; CHECK: llvm.intr.coro.size : i32
   %2 = call i32 @llvm.coro.size.i32()
   ret void
 }
 ; CHECK-LABEL:  llvm.func @coro_align()
 define void @coro_align() {
   ; CHECK: llvm.intr.coro.align : i64
   %1 = call i64 @llvm.coro.align.i64()
   ; CHECK: llvm.intr.coro.align : i32
   %2 = call i32 @llvm.coro.align.i32()
   ret void
 }
 
 ; 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)
   ret void
 }
 
 ; CHECK-LABEL:  llvm.func @coro_end
 define void @coro_end(ptr %0, i1 %1) {
   ; CHECK:  llvm.intr.coro.end
   call i1 @llvm.coro.end(ptr %0, i1 %1)
   ret void
 }
 
 ; 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
 }
 
 ; CHECK-LABEL:  llvm.func @coro_resume
 define void @coro_resume(ptr %0) {
   ; CHECK: llvm.intr.coro.resume %{{.*}}
   call void @llvm.coro.resume(ptr %0)
   ret void
 }
 
 ; 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
 }
 
 ; CHECK-LABEL:  llvm.func @stack_save() {
 define void @stack_save() {
   ; CHECK: llvm.intr.stacksave : !llvm.ptr
   %1 = call ptr @llvm.stacksave()

mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir

llvm.func @umul_with_overflow_test(%arg0: i32, %arg1: i32, %arg2: vector<8xi32>, %arg3: vector<8xi32>) {
   "llvm.intr.umul.with.overflow"(%arg2, %arg3) : (vector<8xi32>, vector<8xi32>) -> !llvm.struct<(vector<8xi32>, vector<8xi1>)>
   llvm.return
 }
 
 // CHECK-LABEL: @coro_id
 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
 }
 
 // CHECK-LABEL: @coro_size
 llvm.func @coro_size() {
   // CHECK: call i64 @llvm.coro.size.i64
   %0 = llvm.intr.coro.size : i64
   // CHECK: call i32 @llvm.coro.size.i32
   %1 = llvm.intr.coro.size : i32
   llvm.return
 }
 
 // CHECK-LABEL: @coro_align
 llvm.func @coro_align() {
   // CHECK: call i64 @llvm.coro.align.i64
   %0 = llvm.intr.coro.align : i64
   // CHECK: call i32 @llvm.coro.align.i32
   %1 = llvm.intr.coro.align : i32
   llvm.return
 }
 
 // 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
 }
 
 // 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
 }
 
 // CHECK-LABEL: @stack_save
 llvm.func @stack_save() {
   // CHECK: call ptr @llvm.stacksave
   %0 = llvm.intr.stacksave : !llvm.ptr<i8>
   llvm.return
 }
 
 // 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
 }
 
 // CHECK-LABEL: @vector_predication_intrinsics
 llvm.func @vector_predication_intrinsics(%A: vector<8xi32>, %B: vector<8xi32>,
                                          %C: vector<8xf32>, %D: vector<8xf32>,
                                          %E: vector<8xi64>, %F: vector<8xf64>,
                                          %G: !llvm.vec<8 x !llvm.ptr<i32>>,

mlir/test/Target/LLVMIR/llvmir.mlir

 // CHECK: define void @vararg_function(i32 %{{.*}}, ...)
 llvm.func @vararg_function(%arg0: i32, ...) {
   %0 = llvm.mlir.constant(1 : i32) : i32
   %1 = llvm.mlir.constant(1 : i32) : i32
   // CHECK: %[[ALLOCA0:.+]] = alloca %struct.va_list, align 8
   %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
 }
 
 // -----
 
 // CHECK: declare void @readonly_function([[PTR:.+]] readonly)
 llvm.func @readonly_function(%arg0: !llvm.ptr<f32> {llvm.readonly})

mlir/test/Target/LLVMIR/nvvmir.mlir

%r = nvvm.wmma.mma %0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15
     : (i32, i32, i32, i32, i32, i32, i32, i32, f32, f32, f32, f32, f32, f32, f32, f32)
     -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
   llvm.return
 }
 
 // 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)
   nvvm.cp.async.wait.group 0
   llvm.return
 }
 
 // CHECK-LABEL: @ld_matrix

mlir/test/mlir-cpu-runner/x86-varargs.mlir

llvm.func @foo(%arg0: i32, ...) -> i32 {
   %7 = llvm.mlir.constant(0 : i64) : i64
   %8 = llvm.mlir.constant(41 : i32) : i32
   %9 = llvm.mlir.constant(0 : i32) : i32
   %10 = llvm.mlir.constant(0 : i64) : i64
   %11 = llvm.mlir.constant(0 : i64) : i64
   %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
   llvm.cond_br %17, ^bb1, ^bb2
 ^bb1:  // pred: ^bb0
   %18 = llvm.getelementptr %13[%7, %6, 3] : (!llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>, i64, i64) -> !llvm.ptr<ptr<i8>>
   %19 = llvm.load %18 : !llvm.ptr<ptr<i8>>
   %20 = llvm.zext %16 : i32 to i64
   %21 = llvm.getelementptr %19[%20] : (!llvm.ptr<i8>, i64) -> !llvm.ptr<i8>
   %22 = llvm.add %16, %4  : i32
   llvm.store %22, %15 : !llvm.ptr<i32>
   llvm.br ^bb3(%21 : !llvm.ptr<i8>)
 ^bb2:  // pred: ^bb0
   %23 = llvm.getelementptr %13[%3, %2, 2] : (!llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>, i64, i64) -> !llvm.ptr<ptr<i8>>
   %24 = llvm.load %23 : !llvm.ptr<ptr<i8>>
   %25 = llvm.getelementptr %24[%0] : (!llvm.ptr<i8>, i64) -> !llvm.ptr<i8>
   llvm.store %25, %23 : !llvm.ptr<ptr<i8>>
   llvm.br ^bb3(%24 : !llvm.ptr<i8>)
 ^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
 }
 
 // CHECK: 2