Index: mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td =================================================================== --- mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td +++ mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td @@ -217,13 +217,13 @@ 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_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_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_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_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_StackRestoreOp : LLVM_ZeroResultIntrOp<"stackrestore"> { let arguments = (ins LLVM_i8Ptr:$ptr); - let assemblyFormat = "$ptr attr-dict"; + let assemblyFormat = "$ptr attr-dict `:` qualified(type($ptr))"; } // Index: mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td =================================================================== --- mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td +++ mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td @@ -158,11 +158,10 @@ // Type constraints accepting LLVM pointer type to integer of a specific width. class LLVM_IntPtrBase : Type< - LLVM_PointerTo>.predicate, - "LLVM pointer to " # I.summary>, - BuildableType<"::mlir::LLVM::LLVMPointerType::get(" - "::mlir::IntegerType::get($_builder.getContext(), " - # width #"), "# addressSpace #")">; + And<[LLVM_PointerTo>.predicate, + CPred<"$_self.cast<::mlir::LLVM::LLVMPointerType>().getAddressSpace()" + " == " # addressSpace>]>, + "LLVM pointer to " # I.summary>; def LLVM_i8Ptr : LLVM_IntPtrBase<8>; Index: mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td =================================================================== --- mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -265,7 +265,7 @@ } createIntrinsicCall(builder, id, {$dst, $src}); }]; - let assemblyFormat = "$dst `,` $src `,` $size attr-dict"; + let assemblyFormat = "$dst `,` $src `,` $size attr-dict `:` type(operands)"; let hasVerifier = 1; } Index: mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir =================================================================== --- mlir/test/Conversion/AsyncToLLVM/convert-coro-to-llvm.mlir +++ 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 } Index: mlir/test/Dialect/LLVMIR/callgraph.mlir =================================================================== --- mlir/test/Dialect/LLVMIR/callgraph.mlir +++ mlir/test/Dialect/LLVMIR/callgraph.mlir @@ -77,7 +77,7 @@ ^bb1: %10 = llvm.landingpad cleanup (catch %3 : !llvm.ptr>) (catch %6 : !llvm.ptr) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr, i32)> - %11 = llvm.intr.eh.typeid.for %6 : i32 + %11 = llvm.intr.eh.typeid.for %6 : (!llvm.ptr) -> i32 llvm.resume %10 : !llvm.struct<(ptr, i32)> ^bb2: Index: mlir/test/Dialect/LLVMIR/invalid.mlir =================================================================== --- mlir/test/Dialect/LLVMIR/invalid.mlir +++ mlir/test/Dialect/LLVMIR/invalid.mlir @@ -1280,7 +1280,7 @@ func.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { // 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, !llvm.ptr return } @@ -1288,7 +1288,7 @@ func.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { // 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, !llvm.ptr return } Index: mlir/test/Dialect/LLVMIR/nvvm.mlir =================================================================== --- mlir/test/Dialect/LLVMIR/nvvm.mlir +++ mlir/test/Dialect/LLVMIR/nvvm.mlir @@ -290,9 +290,9 @@ // CHECK-LABEL: @cp_async llvm.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { // 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, !llvm.ptr // 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, !llvm.ptr // CHECK: nvvm.cp.async.commit.group nvvm.cp.async.commit.group // CHECK: nvvm.cp.async.wait.group 0 Index: mlir/test/Dialect/LLVMIR/roundtrip.mlir =================================================================== --- mlir/test/Dialect/LLVMIR/roundtrip.mlir +++ mlir/test/Dialect/LLVMIR/roundtrip.mlir @@ -404,11 +404,11 @@ // CHECK: ^[[BB1]]: // CHECK: %[[lp:.*]] = llvm.landingpad cleanup (catch %[[a3]] : !llvm.ptr>) (catch %[[a6]] : !llvm.ptr) (filter %[[a2]] : !llvm.array<1 x i8>) : !llvm.struct<(ptr, i32)> -// CHECK: %{{.*}} = llvm.intr.eh.typeid.for %6 : i32 +// CHECK: %{{.*}} = llvm.intr.eh.typeid.for %6 : (!llvm.ptr) -> i32 // CHECK: llvm.resume %[[lp]] : !llvm.struct<(ptr, i32)> ^bb1: %10 = llvm.landingpad cleanup (catch %3 : !llvm.ptr>) (catch %6 : !llvm.ptr) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr, i32)> - %11 = llvm.intr.eh.typeid.for %6 : i32 + %11 = llvm.intr.eh.typeid.for %6 : (!llvm.ptr) -> i32 llvm.resume %10 : !llvm.struct<(ptr, i32)> // CHECK: ^[[BB2]]: @@ -530,17 +530,17 @@ %2 = llvm.alloca %1 x !llvm.struct<"struct.va_list", (ptr)> {alignment = 8 : i64} : (i32) -> !llvm.ptr)>> %3 = llvm.bitcast %2 : !llvm.ptr)>> to !llvm.ptr // CHECK: llvm.intr.vastart %[[CAST0]] - llvm.intr.vastart %3 + llvm.intr.vastart %3 : !llvm.ptr // CHECK: %[[ALLOCA1:.+]] = llvm.alloca %{{.*}} x !llvm.ptr {alignment = 8 : i64} : (i32) -> !llvm.ptr> // CHECK: %[[CAST1:.+]] = llvm.bitcast %[[ALLOCA1]] : !llvm.ptr> to !llvm.ptr %4 = llvm.alloca %0 x !llvm.ptr {alignment = 8 : i64} : (i32) -> !llvm.ptr> %5 = llvm.bitcast %4 : !llvm.ptr> to !llvm.ptr // CHECK: llvm.intr.vacopy %[[CAST0]] to %[[CAST1]] - llvm.intr.vacopy %3 to %5 + llvm.intr.vacopy %3 to %5 : !llvm.ptr, !llvm.ptr // CHECK: llvm.intr.vaend %[[CAST1]] // CHECK: llvm.intr.vaend %[[CAST0]] - llvm.intr.vaend %5 - llvm.intr.vaend %3 + llvm.intr.vaend %5 : !llvm.ptr + llvm.intr.vaend %3 : !llvm.ptr // CHECK: llvm.return llvm.return } @@ -554,3 +554,40 @@ 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 +} Index: mlir/test/Target/LLVMIR/Import/intrinsic.ll =================================================================== --- mlir/test/Target/LLVMIR/Import/intrinsic.ll +++ mlir/test/Target/LLVMIR/Import/intrinsic.ll @@ -431,16 +431,16 @@ ; 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 %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} %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 %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : %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 %{{.*}}, %{{.*}} %4 = call ptr @llvm.coro.begin(token %3, ptr %1) ret void } @@ -464,14 +464,14 @@ ; CHECK-LABEL: llvm.func @coro_save define void @coro_save(ptr %0) { - ; CHECK: llvm.intr.coro.save %{{.*}} : !llvm.token + ; CHECK: llvm.intr.coro.save %{{.*}} %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 %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} %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 @@ ; 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 %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} %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 %{{.*}}, %{{.*}} %4 = call ptr @llvm.coro.free(token %3, ptr %1) ret void } @@ -503,7 +503,7 @@ ; 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 %{{.*}} %2 = call i32 @llvm.eh.typeid.for(ptr %0) ret void } Index: mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir =================================================================== --- mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir +++ mlir/test/Target/LLVMIR/llvmir-intrinsics.mlir @@ -432,16 +432,16 @@ llvm.func @coro_id(%arg0: i32, %arg1: !llvm.ptr) { // CHECK: call token @llvm.coro.id %null = llvm.mlir.null : !llvm.ptr - llvm.intr.coro.id %arg0, %arg1, %arg1, %null : !llvm.token + llvm.intr.coro.id %arg0, %arg1, %arg1, %null : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token llvm.return } // CHECK-LABEL: @coro_begin llvm.func @coro_begin(%arg0: i32, %arg1: !llvm.ptr) { %null = llvm.mlir.null : !llvm.ptr - %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : !llvm.token + %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token // CHECK: call ptr @llvm.coro.begin - llvm.intr.coro.begin %token, %arg1 : !llvm.ptr + llvm.intr.coro.begin %token, %arg1 : (!llvm.token, !llvm.ptr) -> !llvm.ptr llvm.return } @@ -466,14 +466,14 @@ // CHECK-LABEL: @coro_save llvm.func @coro_save(%arg0: !llvm.ptr) { // CHECK: call token @llvm.coro.save - %0 = llvm.intr.coro.save %arg0 : !llvm.token + %0 = llvm.intr.coro.save %arg0 : (!llvm.ptr) -> !llvm.token llvm.return } // CHECK-LABEL: @coro_suspend llvm.func @coro_suspend(%arg0: i32, %arg1 : i1, %arg2 : !llvm.ptr) { %null = llvm.mlir.null : !llvm.ptr - %token = llvm.intr.coro.id %arg0, %arg2, %arg2, %null : !llvm.token + %token = llvm.intr.coro.id %arg0, %arg2, %arg2, %null : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token // CHECK: call i8 @llvm.coro.suspend %0 = llvm.intr.coro.suspend %token, %arg1 : i8 llvm.return @@ -482,30 +482,30 @@ // CHECK-LABEL: @coro_end llvm.func @coro_end(%arg0: !llvm.ptr, %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, i1) -> i1 llvm.return } // CHECK-LABEL: @coro_free llvm.func @coro_free(%arg0: i32, %arg1 : !llvm.ptr) { %null = llvm.mlir.null : !llvm.ptr - %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : !llvm.token + %token = llvm.intr.coro.id %arg0, %arg1, %arg1, %null : (i32, !llvm.ptr, !llvm.ptr, !llvm.ptr) -> !llvm.token // CHECK: call ptr @llvm.coro.free - %0 = llvm.intr.coro.free %token, %arg1 : !llvm.ptr + %0 = llvm.intr.coro.free %token, %arg1 : (!llvm.token, !llvm.ptr) -> !llvm.ptr llvm.return } // CHECK-LABEL: @coro_resume llvm.func @coro_resume(%arg0: !llvm.ptr) { // CHECK: call void @llvm.coro.resume - llvm.intr.coro.resume %arg0 + llvm.intr.coro.resume %arg0 : !llvm.ptr llvm.return } // CHECK-LABEL: @eh_typeid_for llvm.func @eh_typeid_for(%arg0 : !llvm.ptr) { // CHECK: call i32 @llvm.eh.typeid.for - %0 = llvm.intr.eh.typeid.for %arg0 : i32 + %0 = llvm.intr.eh.typeid.for %arg0 : (!llvm.ptr) -> i32 llvm.return } @@ -519,7 +519,7 @@ // CHECK-LABEL: @stack_restore llvm.func @stack_restore(%arg0: !llvm.ptr) { // CHECK: call void @llvm.stackrestore - llvm.intr.stackrestore %arg0 + llvm.intr.stackrestore %arg0 : !llvm.ptr llvm.return } Index: mlir/test/Target/LLVMIR/llvmir.mlir =================================================================== --- mlir/test/Target/LLVMIR/llvmir.mlir +++ mlir/test/Target/LLVMIR/llvmir.mlir @@ -2120,16 +2120,16 @@ %2 = llvm.alloca %1 x !llvm.struct<"struct.va_list", (ptr)> {alignment = 8 : i64} : (i32) -> !llvm.ptr)>> %3 = llvm.bitcast %2 : !llvm.ptr)>> to !llvm.ptr // CHECK: call void @llvm.va_start(ptr %[[ALLOCA0]]) - llvm.intr.vastart %3 + llvm.intr.vastart %3 : !llvm.ptr // CHECK: %[[ALLOCA1:.+]] = alloca ptr, align 8 %4 = llvm.alloca %0 x !llvm.ptr {alignment = 8 : i64} : (i32) -> !llvm.ptr> %5 = llvm.bitcast %4 : !llvm.ptr> to !llvm.ptr // CHECK: call void @llvm.va_copy(ptr %[[ALLOCA1]], ptr %[[ALLOCA0]]) - llvm.intr.vacopy %3 to %5 + llvm.intr.vacopy %3 to %5 : !llvm.ptr, !llvm.ptr // 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 + llvm.intr.vaend %3 : !llvm.ptr // CHECK: ret void llvm.return } Index: mlir/test/Target/LLVMIR/nvvmir.mlir =================================================================== --- mlir/test/Target/LLVMIR/nvvmir.mlir +++ mlir/test/Target/LLVMIR/nvvmir.mlir @@ -309,13 +309,13 @@ // CHECK-LABEL: @cp_async llvm.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { // 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, !llvm.ptr // 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, !llvm.ptr // 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, !llvm.ptr // 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, !llvm.ptr // 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) Index: mlir/test/mlir-cpu-runner/x86-varargs.mlir =================================================================== --- mlir/test/mlir-cpu-runner/x86-varargs.mlir +++ mlir/test/mlir-cpu-runner/x86-varargs.mlir @@ -38,7 +38,7 @@ %12 = llvm.mlir.constant(1 : i32) : i32 %13 = llvm.alloca %12 x !llvm.array<1 x struct<"struct.va_list", (i32, i32, ptr, ptr)>> {alignment = 8 : i64} : (i32) -> !llvm.ptr, ptr)>>> %14 = llvm.bitcast %13 : !llvm.ptr, ptr)>>> to !llvm.ptr - llvm.intr.vastart %14 + llvm.intr.vastart %14 : !llvm.ptr %15 = llvm.getelementptr %13[%11, %10, 0] : (!llvm.ptr, ptr)>>>, i64, i64) -> !llvm.ptr %16 = llvm.load %15 : !llvm.ptr %17 = llvm.icmp "ult" %16, %8 : i32 @@ -60,7 +60,7 @@ ^bb3(%26: !llvm.ptr): // 2 preds: ^bb1, ^bb2 %27 = llvm.bitcast %26 : !llvm.ptr to !llvm.ptr %28 = llvm.load %27 : !llvm.ptr - llvm.intr.vaend %14 + llvm.intr.vaend %14 : !llvm.ptr llvm.return %28 : i32 }