[mlir][llvm] Switch remaining LLVM dialect tests to opaque pointers.

[gysit]

Summary

The revision switches the remaining LLVM dialect tests to use opaque
pointers. Selected tests are copied to a postfixed test file for the
time being.

A number of tests disappear once we fully switch to opaque pointers.
In particular, all tests that check verify a pointer element type
matches another type as well as tests of recursive types.

Part of https://discourse.llvm.org/t/rfc-switching-the-llvm-dialect-and-dialect-lowerings-to-opaque-pointers/68179

[Dinistro]

LGTM! Thanks for pushing on this.

[zero9178]

LGTM!

mlir/test/Dialect/LLVMIR/callgraph.mlir

 
 // CHECK: Testing : "Invoke call"
 module attributes {"test.name" = "Invoke call"} {
   // CHECK-LABEL: ---- CallGraph ----
   // CHECK: - Node : 'llvm.func' {{.*}} sym_name = "invokeLandingpad"
   // CHECK-DAG: -- Call-Edge : <Unknown-Callee-Node>
 
   // CHECK: -- SCCs --
-  llvm.mlir.global external constant @_ZTIi() : !llvm.ptr<i8>
+  llvm.mlir.global external constant @_ZTIi() : !llvm.ptr
   llvm.func @foo(%arg0: i32) -> !llvm.struct<(i32, f64, i32)>
-  llvm.func @bar(!llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>)
+  llvm.func @bar(!llvm.ptr, !llvm.ptr, !llvm.ptr)
   llvm.func @__gxx_personality_v0(...) -> i32
 
   llvm.func @invokeLandingpad() -> i32 attributes { personality = @__gxx_personality_v0 } {
     %0 = llvm.mlir.constant(0 : i32) : i32
     %1 = llvm.mlir.constant(3 : i32) : i32
     %2 = llvm.mlir.constant("\01") : !llvm.array<1 x i8>
-    %3 = llvm.mlir.null : !llvm.ptr<ptr<i8>>
-    %4 = llvm.mlir.null : !llvm.ptr<i8>
-    %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)>
+    %3 = llvm.mlir.null : !llvm.ptr
+    %4 = llvm.mlir.null : !llvm.ptr
+    %5 = llvm.mlir.addressof @_ZTIi : !llvm.ptr
+    %6 = llvm.mlir.constant(1 : i32) : i32
+    %7 = llvm.alloca %6 x i8 : (i32) -> !llvm.ptr
+    %8 = llvm.invoke @foo(%6) 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 : (!llvm.ptr<i8>) -> i32
-    llvm.resume %10 : !llvm.struct<(ptr<i8>, i32)>
+    %10 = llvm.landingpad cleanup (catch %3 : !llvm.ptr) (catch %5 : !llvm.ptr) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr, i32)>
+    %11 = llvm.intr.eh.typeid.for %5 : (!llvm.ptr) -> i32
+    llvm.resume %10 : !llvm.struct<(ptr, i32)>
 
   ^bb2:
-    llvm.return %7 : i32
+    llvm.return %6 : i32
 
   ^bb3:
-    llvm.invoke @bar(%8, %6, %4) to ^bb2 unwind ^bb1 : (!llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>) -> ()
+    llvm.invoke @bar(%7, %5, %4) to ^bb2 unwind ^bb1 : (!llvm.ptr, !llvm.ptr, !llvm.ptr) -> ()
 
   ^bb4:
     llvm.return %0 : i32
   }
 }

mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir

 // RUN: mlir-opt --pass-pipeline='builtin.module(llvm.func(canonicalize{test-convergence}))' %s -split-input-file | FileCheck %s
 
-// CHECK-LABEL: fold_extractvalue
-llvm.func @fold_extractvalue() -> i32 {
-  //  CHECK-DAG: %[[C0:.*]] = arith.constant 0 : i32
-  %c0 = arith.constant 0 : i32
-  //  CHECK-DAG: %[[C1:.*]] = arith.constant 1 : i32
-  %c1 = arith.constant 1 : i32
-
-  %0 = llvm.mlir.undef : !llvm.struct<(i32, i32)>
-
-  // CHECK-NOT: insertvalue
-  %1 = llvm.insertvalue %c0, %0[0] : !llvm.struct<(i32, i32)>
-  %2 = llvm.insertvalue %c1, %1[1] : !llvm.struct<(i32, i32)>
-
-  // CHECK-NOT: extractvalue
-  %3 = llvm.extractvalue %2[0] : !llvm.struct<(i32, i32)>
-  %4 = llvm.extractvalue %2[1] : !llvm.struct<(i32, i32)>
-
-  //     CHECK: llvm.add %[[C0]], %[[C1]]
-  %5 = llvm.add %3, %4 : i32
-  llvm.return %5 : i32
-}
-
-// -----
-
-// CHECK-LABEL: no_fold_extractvalue
-llvm.func @no_fold_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 {
-  %f0 = arith.constant 0.0 : f32
-  %0 = llvm.mlir.undef : !llvm.array<4 x !llvm.array<4 x f32>>
-
-  // CHECK: insertvalue
-  // CHECK: insertvalue
-  // CHECK: extractvalue
-  %1 = llvm.insertvalue %f0, %0[0, 0] : !llvm.array<4 x !llvm.array<4 x f32>>
-  %2 = llvm.insertvalue %arr, %1[0] : !llvm.array<4 x !llvm.array<4 x f32>>
-  %3 = llvm.extractvalue %2[0, 0] : !llvm.array<4 x !llvm.array<4 x f32>>
-
-  llvm.return %3 : f32
-
-}
-// -----
-
-// CHECK-LABEL: fold_unrelated_extractvalue
-llvm.func @fold_unrelated_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 {
-  %f0 = arith.constant 0.0 : f32
-  // CHECK-NOT: insertvalue
-  // CHECK: extractvalue
-  %2 = llvm.insertvalue %f0, %arr[0] : !llvm.array<4 x f32>
-  %3 = llvm.extractvalue %2[1] : !llvm.array<4 x f32>
-  llvm.return %3 : f32
-}
-
-// -----
-
 // CHECK-LABEL: fold_bitcast
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: llvm.return %[[a0]]
 llvm.func @fold_bitcast(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
   %c = llvm.bitcast %x : !llvm.ptr<i8> to !llvm.ptr<i8>
   llvm.return %c : !llvm.ptr<i8>
 }
 

 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: llvm.return %[[a0]]
 llvm.func @fold_gep(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
   %c0 = arith.constant 0 : i32
   %c = llvm.getelementptr %x[%c0] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
   llvm.return %c : !llvm.ptr<i8>
 }
 
-// CHECK-LABEL: fold_gep_neg
-// CHECK-SAME: %[[a0:arg[0-9]+]]
-// CHECK-NEXT: %[[RES:.*]] = llvm.getelementptr inbounds %[[a0]][0, 1]
-// CHECK-NEXT: llvm.return %[[RES]]
-llvm.func @fold_gep_neg(%x : !llvm.ptr) -> !llvm.ptr {
-  %c0 = arith.constant 0 : i32
-  %0 = llvm.getelementptr inbounds %x[%c0, 1] : (!llvm.ptr, i32) -> !llvm.ptr, !llvm.struct<(i32, i32)>
-  llvm.return %0 : !llvm.ptr
-}
+// -----
 
 // CHECK-LABEL: fold_gep_canon
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: %[[RES:.*]] = llvm.getelementptr %[[a0]][2]
 // CHECK-NEXT: llvm.return %[[RES]]
 llvm.func @fold_gep_canon(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
   %c2 = arith.constant 2 : i32
   %c = llvm.getelementptr %x[%c2] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
   llvm.return %c : !llvm.ptr<i8>
 }
 
-
-// -----
-
-// Check that LLVM constants participate in cross-dialect constant folding. The
-// resulting constant is created in the arith dialect because the last folded
-// operation belongs to it.
-// CHECK-LABEL: llvm_constant
-llvm.func @llvm_constant() -> i32 {
-  // CHECK-NOT: llvm.mlir.constant
-  %0 = llvm.mlir.constant(40 : i32) : i32
-  %1 = llvm.mlir.constant(42 : i32) : i32
-  // CHECK: %[[RES:.*]] = arith.constant 82 : i32
-  // CHECK-NOT: arith.addi
-  %2 = arith.addi %0, %1 : i32
-  // CHECK: return %[[RES]]
-  llvm.return %2 : i32
-}
-
 // -----
 
 // CHECK-LABEL: load_dce
 // CHECK-NEXT: llvm.return
 llvm.func @load_dce(%x : !llvm.ptr<i8>) {
   %0 = llvm.load %x : !llvm.ptr<i8>
   llvm.return
 }

mlir/test/Dialect/LLVMIR/canonicalize.mlir

llvm.func @no_fold_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 {
   // CHECK: insertvalue
   // CHECK: insertvalue
   // CHECK: extractvalue
   %1 = llvm.insertvalue %f0, %0[0, 0] : !llvm.array<4 x !llvm.array<4 x f32>>
   %2 = llvm.insertvalue %arr, %1[0] : !llvm.array<4 x !llvm.array<4 x f32>>
   %3 = llvm.extractvalue %2[0, 0] : !llvm.array<4 x !llvm.array<4 x f32>>
 
   llvm.return %3 : f32
-
 }
+
 // -----
 
 // CHECK-LABEL: fold_unrelated_extractvalue
 llvm.func @fold_unrelated_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 {
   %f0 = arith.constant 0.0 : f32
   // CHECK-NOT: insertvalue
   // CHECK: extractvalue
   %2 = llvm.insertvalue %f0, %arr[0] : !llvm.array<4 x f32>
   %3 = llvm.extractvalue %2[1] : !llvm.array<4 x f32>
   llvm.return %3 : f32
 }
 
 // -----
 
 // CHECK-LABEL: fold_bitcast
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_bitcast(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
-  %c = llvm.bitcast %x : !llvm.ptr<i8> to !llvm.ptr<i8>
-  llvm.return %c : !llvm.ptr<i8>
+llvm.func @fold_bitcast(%x : !llvm.ptr) -> !llvm.ptr {
+  %c = llvm.bitcast %x : !llvm.ptr to !llvm.ptr
+  llvm.return %c : !llvm.ptr
 }
 
 // CHECK-LABEL: fold_bitcast2
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_bitcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
-  %c = llvm.bitcast %x : !llvm.ptr<i8> to !llvm.ptr<i32>
-  %d = llvm.bitcast %c : !llvm.ptr<i32> to !llvm.ptr<i8>
-  llvm.return %d : !llvm.ptr<i8>
+llvm.func @fold_bitcast2(%x : i32) -> i32 {
+  %c = llvm.bitcast %x : i32 to f32
+  %d = llvm.bitcast %c : f32 to i32
+  llvm.return %d : i32
 }
 
 // -----
 
 // CHECK-LABEL: fold_addrcast
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_addrcast(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
-  %c = llvm.addrspacecast %x : !llvm.ptr<i8> to !llvm.ptr<i8>
-  llvm.return %c : !llvm.ptr<i8>
+llvm.func @fold_addrcast(%x : !llvm.ptr) -> !llvm.ptr {
+  %c = llvm.addrspacecast %x : !llvm.ptr to !llvm.ptr
+  llvm.return %c : !llvm.ptr
 }
 
 // CHECK-LABEL: fold_addrcast2
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_addrcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
-  %c = llvm.addrspacecast %x : !llvm.ptr<i8> to !llvm.ptr<i32, 5>
-  %d = llvm.addrspacecast %c : !llvm.ptr<i32, 5> to !llvm.ptr<i8>
-  llvm.return %d : !llvm.ptr<i8>
+llvm.func @fold_addrcast2(%x : !llvm.ptr) -> !llvm.ptr {
+  %c = llvm.addrspacecast %x : !llvm.ptr to !llvm.ptr<5>
+  %d = llvm.addrspacecast %c : !llvm.ptr<5> to !llvm.ptr
+  llvm.return %d : !llvm.ptr
 }
 
 // -----
 
 // CHECK-LABEL: fold_gep
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_gep(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+llvm.func @fold_gep(%x : !llvm.ptr) -> !llvm.ptr {
   %c0 = arith.constant 0 : i32
-  %c = llvm.getelementptr %x[%c0] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
-  llvm.return %c : !llvm.ptr<i8>
+  %c = llvm.getelementptr %x[%c0] : (!llvm.ptr, i32) -> !llvm.ptr, i8
+  llvm.return %c : !llvm.ptr
 }
 
 // CHECK-LABEL: fold_gep_neg
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: %[[RES:.*]] = llvm.getelementptr inbounds %[[a0]][0, 1]
 // CHECK-NEXT: llvm.return %[[RES]]
 llvm.func @fold_gep_neg(%x : !llvm.ptr) -> !llvm.ptr {
   %c0 = arith.constant 0 : i32
   %0 = llvm.getelementptr inbounds %x[%c0, 1] : (!llvm.ptr, i32) -> !llvm.ptr, !llvm.struct<(i32, i32)>
   llvm.return %0 : !llvm.ptr
 }
 
 // CHECK-LABEL: fold_gep_canon
 // CHECK-SAME: %[[a0:arg[0-9]+]]
 // CHECK-NEXT: %[[RES:.*]] = llvm.getelementptr %[[a0]][2]
 // CHECK-NEXT: llvm.return %[[RES]]
-llvm.func @fold_gep_canon(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+llvm.func @fold_gep_canon(%x : !llvm.ptr) -> !llvm.ptr {
   %c2 = arith.constant 2 : i32
-  %c = llvm.getelementptr %x[%c2] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
-  llvm.return %c : !llvm.ptr<i8>
+  %c = llvm.getelementptr %x[%c2] : (!llvm.ptr, i32) -> !llvm.ptr, i8
+  llvm.return %c : !llvm.ptr
 }
 
-
 // -----
 
 // Check that LLVM constants participate in cross-dialect constant folding. The
 // resulting constant is created in the arith dialect because the last folded
 // operation belongs to it.
 // CHECK-LABEL: llvm_constant
 llvm.func @llvm_constant() -> i32 {
   // CHECK-NOT: llvm.mlir.constant
   %0 = llvm.mlir.constant(40 : i32) : i32
   %1 = llvm.mlir.constant(42 : i32) : i32
   // CHECK: %[[RES:.*]] = arith.constant 82 : i32
   // CHECK-NOT: arith.addi
   %2 = arith.addi %0, %1 : i32
   // CHECK: return %[[RES]]
   llvm.return %2 : i32
 }
 
 // -----
 
 // CHECK-LABEL: load_dce
 // CHECK-NEXT: llvm.return
-llvm.func @load_dce(%x : !llvm.ptr<i8>) {
-  %0 = llvm.load %x : !llvm.ptr<i8>
+llvm.func @load_dce(%x : !llvm.ptr) {
+  %0 = llvm.load %x : !llvm.ptr -> i8
   llvm.return
 }
 
-llvm.mlir.global external @fp() : !llvm.ptr<i8>
+llvm.mlir.global external @fp() : !llvm.ptr
 
 // CHECK-LABEL: addr_dce
 // CHECK-NEXT: llvm.return
-llvm.func @addr_dce(%x : !llvm.ptr<i8>) {
-  %0 = llvm.mlir.addressof @fp : !llvm.ptr<ptr<i8>>
+llvm.func @addr_dce(%x : !llvm.ptr) {
+  %0 = llvm.mlir.addressof @fp : !llvm.ptr
   llvm.return
 }
 
 // CHECK-LABEL: alloca_dce
 // CHECK-NEXT: llvm.return
 llvm.func @alloca_dce() {
   %c1_i64 = arith.constant 1 : i64
-  %0 = llvm.alloca %c1_i64 x i32 : (i64) -> !llvm.ptr<i32>
+  %0 = llvm.alloca %c1_i64 x i32 : (i64) -> !llvm.ptr
   llvm.return
 }

mlir/test/Dialect/LLVMIR/debuginfo.mlir

#var2 = #llvm.di_local_variable<
   // Omit the optional parameters.
   scope = #block2, name = "arg2"
 >
 
 // CHECK: llvm.func @addr(%[[ARG:.*]]: i64)
 llvm.func @addr(%arg: i64) {
   // CHECK: %[[ALLOC:.*]] = llvm.alloca
   %allocCount = llvm.mlir.constant(1 : i32) : i32
-  %alloc = llvm.alloca %allocCount x i64 : (i32) -> !llvm.ptr<i64>
+  %alloc = llvm.alloca %allocCount x i64 : (i32) -> !llvm.ptr
 
   // CHECK: llvm.intr.dbg.declare #[[VAR0]] = %[[ALLOC]]
-  llvm.intr.dbg.declare #var0 = %alloc : !llvm.ptr<i64>
+  llvm.intr.dbg.declare #var0 = %alloc : !llvm.ptr
   llvm.return
 }
 
 // CHECK: llvm.func @value(%[[ARG1:.*]]: i32, %[[ARG2:.*]]: i32)
 llvm.func @value(%arg1: i32, %arg2: i32) {
   // CHECK: llvm.intr.dbg.value #[[VAR1]] = %[[ARG1]]
   llvm.intr.dbg.value #var1 = %arg1 : i32
   // CHECK: llvm.intr.dbg.value #[[VAR2]] = %[[ARG2]]
   llvm.intr.dbg.value #var2 = %arg2 : i32
   llvm.return
 }

mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir

 // RUN: mlir-opt %s | FileCheck %s
 
 module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<i1, dense<8> : vector<2xi32>>, #dlti.dl_entry<i8, dense<8> : vector<2xi32>>, #dlti.dl_entry<i16, dense<16> : vector<2xi32>>, #dlti.dl_entry<i32, dense<32> : vector<2xi32>>, #dlti.dl_entry<i64, dense<[32, 64]> : vector<2xi32>>, #dlti.dl_entry<f16, dense<16> : vector<2xi32>>, #dlti.dl_entry<f64, dense<64> : vector<2xi32>>, #dlti.dl_entry<f128, dense<128> : vector<2xi32>>>} {
-  // CHECK: llvm.func @foo(%[[ARG0:.+]]: !llvm.ptr<struct<"my_struct", {{.+}}>>, %[[ARG1:.+]]: i32)
-  llvm.func @foo(%arg0: !llvm.ptr<struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>>, %arg1: i32) {
+  // CHECK: llvm.func @foo(%[[ARG0:.+]]: !llvm.ptr, %[[ARG1:.+]]: i32)
+  llvm.func @foo(%arg0: !llvm.ptr, %arg1: i32) {
     // CHECK: %[[C0:.+]] = llvm.mlir.constant(0 : i32)
     %0 = llvm.mlir.constant(0 : i32) : i32
     // CHECK: llvm.getelementptr %[[ARG0]][%[[C0]], 1, %[[ARG1]]]
-    %1 = "llvm.getelementptr"(%arg0, %0, %arg1) {rawConstantIndices = array<i32: -2147483648, 1, -2147483648>} : (!llvm.ptr<struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>>, i32, i32) -> !llvm.ptr<i32>
+    %1 = "llvm.getelementptr"(%arg0, %0, %arg1) {elem_type = !llvm.struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>, rawConstantIndices = array<i32: -2147483648, 1, -2147483648>} : (!llvm.ptr, i32, i32) -> !llvm.ptr
     llvm.return
   }
 }

mlir/test/Dialect/LLVMIR/func.mlir

^bb0(%arg0: i64):
     // GENERIC: llvm.return
     llvm.return %arg0 : i64
 
   // GENERIC: function_type = !llvm.func<i64 (i64)>
   // GENERIC-SAME: sym_name = "baz"
   // GENERIC-SAME: () -> ()
   }) {sym_name = "baz", function_type = !llvm.func<i64 (i64)>} : () -> ()
 
-  // CHECK: llvm.func @qux(!llvm.ptr<i64> {llvm.noalias}, i64)
+  // CHECK: llvm.func @qux(!llvm.ptr {llvm.noalias}, i64)
   // CHECK: attributes {xxx = {yyy = 42 : i64}}
   "llvm.func"() ({
-  }) {sym_name = "qux", function_type = !llvm.func<void (ptr<i64>, i64)>,
+  }) {sym_name = "qux", function_type = !llvm.func<void (ptr, i64)>,
       arg_attrs = [{llvm.noalias}, {}], xxx = {yyy = 42}} : () -> ()
 
   // CHECK: llvm.func @roundtrip1()
   llvm.func @roundtrip1()
 
   // CHECK: llvm.func @roundtrip2(i64, f32) -> f64
   llvm.func @roundtrip2(i64, f32) -> f64
 
module {
   // CHECK: llvm.func @roundtrip7() {
   llvm.func @roundtrip7() attributes {} {
     llvm.return
   }
 
   // CHECK: llvm.func @roundtrip8() -> i32
   llvm.func @roundtrip8() -> i32 attributes {}
 
-  // CHECK: llvm.func @roundtrip9(!llvm.ptr<i32> {llvm.noalias})
-  llvm.func @roundtrip9(!llvm.ptr<i32> {llvm.noalias})
+  // CHECK: llvm.func @roundtrip9(!llvm.ptr {llvm.noalias})
+  llvm.func @roundtrip9(!llvm.ptr {llvm.noalias})
 
-  // CHECK: llvm.func @roundtrip10(!llvm.ptr<i32> {llvm.noalias})
-  llvm.func @roundtrip10(%arg0: !llvm.ptr<i32> {llvm.noalias})
+  // CHECK: llvm.func @roundtrip10(!llvm.ptr {llvm.noalias})
+  llvm.func @roundtrip10(%arg0: !llvm.ptr {llvm.noalias})
 
-  // CHECK: llvm.func @roundtrip11(%{{.*}}: !llvm.ptr<i32> {llvm.noalias}) {
-  llvm.func @roundtrip11(%arg0: !llvm.ptr<i32> {llvm.noalias}) {
+  // CHECK: llvm.func @roundtrip11(%{{.*}}: !llvm.ptr {llvm.noalias}) {
+  llvm.func @roundtrip11(%arg0: !llvm.ptr {llvm.noalias}) {
     llvm.return
   }
 
-  // CHECK: llvm.func @roundtrip12(%{{.*}}: !llvm.ptr<i32> {llvm.noalias})
+  // CHECK: llvm.func @roundtrip12(%{{.*}}: !llvm.ptr {llvm.noalias})
   // CHECK: attributes {foo = 42 : i32}
-  llvm.func @roundtrip12(%arg0: !llvm.ptr<i32> {llvm.noalias})
+  llvm.func @roundtrip12(%arg0: !llvm.ptr {llvm.noalias})
   attributes {foo = 42 : i32} {
     llvm.return
   }
 
-  // CHECK: llvm.func @byvalattr(%{{.*}}: !llvm.ptr<i32> {llvm.byval = i32})
-  llvm.func @byvalattr(%arg0: !llvm.ptr<i32> {llvm.byval = i32}) {
+  // CHECK: llvm.func @byvalattr(%{{.*}}: !llvm.ptr {llvm.byval = i32})
+  llvm.func @byvalattr(%arg0: !llvm.ptr {llvm.byval = i32}) {
     llvm.return
   }
 
-  // CHECK: llvm.func @sretattr(%{{.*}}: !llvm.ptr<i32> {llvm.sret = i32})
-  // LOCINFO: llvm.func @sretattr(%{{.*}}: !llvm.ptr<i32> {llvm.sret = i32} loc("some_source_loc"))
-  llvm.func @sretattr(%arg0: !llvm.ptr<i32> {llvm.sret = i32} loc("some_source_loc")) {
+  // CHECK: llvm.func @sretattr(%{{.*}}: !llvm.ptr {llvm.sret = i32})
+  // LOCINFO: llvm.func @sretattr(%{{.*}}: !llvm.ptr {llvm.sret = i32} loc("some_source_loc"))
+  llvm.func @sretattr(%arg0: !llvm.ptr {llvm.sret = i32} loc("some_source_loc")) {
     llvm.return
   }
 
-  // CHECK: llvm.func @nestattr(%{{.*}}: !llvm.ptr<i32> {llvm.nest})
-  llvm.func @nestattr(%arg0: !llvm.ptr<i32> {llvm.nest}) {
+  // CHECK: llvm.func @nestattr(%{{.*}}: !llvm.ptr {llvm.nest})
+  llvm.func @nestattr(%arg0: !llvm.ptr {llvm.nest}) {
     llvm.return
   }
 
-  // CHECK: llvm.func @llvm_noalias_decl(!llvm.ptr<f32> {llvm.noalias})
-  llvm.func @llvm_noalias_decl(!llvm.ptr<f32> {llvm.noalias})
-  // CHECK: llvm.func @byrefattr_decl(!llvm.ptr<i32> {llvm.byref = i32})
-  llvm.func @byrefattr_decl(!llvm.ptr<i32> {llvm.byref = i32})
-  // CHECK: llvm.func @byvalattr_decl(!llvm.ptr<i32> {llvm.byval = i32})
-  llvm.func @byvalattr_decl(!llvm.ptr<i32> {llvm.byval = i32})
-  // CHECK: llvm.func @sretattr_decl(!llvm.ptr<i32> {llvm.sret = i32})
-  llvm.func @sretattr_decl(!llvm.ptr<i32> {llvm.sret = i32})
-  // CHECK: llvm.func @nestattr_decl(!llvm.ptr<i32> {llvm.nest})
-  llvm.func @nestattr_decl(!llvm.ptr<i32> {llvm.nest})
+  // CHECK: llvm.func @llvm_noalias_decl(!llvm.ptr {llvm.noalias})
+  llvm.func @llvm_noalias_decl(!llvm.ptr {llvm.noalias})
+  // CHECK: llvm.func @byrefattr_decl(!llvm.ptr {llvm.byref = i32})
+  llvm.func @byrefattr_decl(!llvm.ptr {llvm.byref = i32})
+  // CHECK: llvm.func @byvalattr_decl(!llvm.ptr {llvm.byval = i32})
+  llvm.func @byvalattr_decl(!llvm.ptr {llvm.byval = i32})
+  // CHECK: llvm.func @sretattr_decl(!llvm.ptr {llvm.sret = i32})
+  llvm.func @sretattr_decl(!llvm.ptr {llvm.sret = i32})
+  // CHECK: llvm.func @nestattr_decl(!llvm.ptr {llvm.nest})
+  llvm.func @nestattr_decl(!llvm.ptr {llvm.nest})
   // CHECK: llvm.func @noundefattr_decl(i32 {llvm.noundef})
   llvm.func @noundefattr_decl(i32 {llvm.noundef})
-  // CHECK: llvm.func @llvm_align_decl(!llvm.ptr<f32> {llvm.align = 4 : i64})
-  llvm.func @llvm_align_decl(!llvm.ptr<f32> {llvm.align = 4})
-  // CHECK: llvm.func @inallocaattr_decl(!llvm.ptr<i32> {llvm.inalloca = i32})
-  llvm.func @inallocaattr_decl(!llvm.ptr<i32> {llvm.inalloca = i32})
+  // CHECK: llvm.func @llvm_align_decl(!llvm.ptr {llvm.align = 4 : i64})
+  llvm.func @llvm_align_decl(!llvm.ptr {llvm.align = 4})
+  // CHECK: llvm.func @inallocaattr_decl(!llvm.ptr {llvm.inalloca = i32})
+  llvm.func @inallocaattr_decl(!llvm.ptr {llvm.inalloca = i32})
 
 
   // CHECK: llvm.func @variadic(...)
   llvm.func @variadic(...)
 
   // CHECK: llvm.func @variadic_args(i32, i32, ...)
   llvm.func @variadic_args(i32, i32, ...)
 

mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir

+// RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
+
+// CHECK: llvm.mlir.global internal @global(42 : i64) {addr_space = 0 : i32} : i64
+llvm.mlir.global internal @global(42 : i64) : i64
+
+// CHECK: llvm.mlir.global internal constant @".string"("foobar")
+llvm.mlir.global internal constant @".string"("foobar") : !llvm.array<6 x i8>
+
+func.func @references() {
+  // CHECK: llvm.mlir.addressof @global : !llvm.ptr<i64>
+  %0 = llvm.mlir.addressof @global : !llvm.ptr<i64>
+
+  // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
+  %1 = llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
+
+  llvm.return
+}
+
+// -----
+
+llvm.mlir.global internal @foo(0: i32) : i32
+
+func.func @bar() {
+  // expected-error @+1 {{the type must be a pointer to the type of the referenced global}}
+  llvm.mlir.addressof @foo : !llvm.ptr<i64>
+  llvm.return
+}
+
+// -----
+
+llvm.func @foo()
+
+llvm.func @bar() {
+  // expected-error @+1 {{the type must be a pointer to the type of the referenced function}}
+  llvm.mlir.addressof @foo : !llvm.ptr<i8>
+  llvm.return
+}
+
+// -----
+
+llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
+func.func @mismatch_addr_space() {
+  // expected-error @+1 {{pointer address space must match address space of the referenced global}}
+  llvm.mlir.addressof @g : !llvm.ptr<i64, 4>
+  llvm.return
+}

mlir/test/Dialect/LLVMIR/global.mlir

 llvm.mlir.global external @has_thr_local(42 : i64) {thr_local} : i64
 // CHECK: llvm.mlir.global external @has_dso_local(42 : i64) {addr_space = 0 : i32, dso_local} : i64
 llvm.mlir.global external @has_dso_local(42 : i64) {dso_local} : i64
 // CHECK: llvm.mlir.global external @has_addr_space(32 : i64) {addr_space = 3 : i32} : i64
 llvm.mlir.global external @has_addr_space(32 : i64) {addr_space = 3: i32} : i64
 
 // CHECK-LABEL: references
 func.func @references() {
-  // CHECK: llvm.mlir.addressof @global : !llvm.ptr<i64>
-  %0 = llvm.mlir.addressof @global : !llvm.ptr<i64>
-
-  // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
-  %1 = llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
+  // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr
+  %0 = llvm.mlir.addressof @".string" : !llvm.ptr
 
   // CHECK: llvm.mlir.addressof @global : !llvm.ptr
-  %2 = llvm.mlir.addressof @global : !llvm.ptr
+  %1 = llvm.mlir.addressof @global : !llvm.ptr
 
   // CHECK: llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
-  %3 = llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
+  %2 = llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
 
   llvm.return
 }
 
 // CHECK: llvm.mlir.global private local_unnamed_addr constant @local(42 : i64) {addr_space = 0 : i32} : i64
 llvm.mlir.global private local_unnamed_addr constant @local(42 : i64) : i64
 
 // CHECK: llvm.mlir.global private unnamed_addr constant @foo(42 : i64) {addr_space = 0 : i32} : i64

 }
 
 // -----
 
 func.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}}
-  llvm.mlir.addressof "foo" : i64 : !llvm.ptr<func<void ()>>
+  llvm.mlir.addressof "foo" : i64 : !llvm.ptr
   llvm.return
 }
 
 // -----
 
 func.func @foo() {
   // expected-error @+1 {{must reference a global defined by 'llvm.mlir.global'}}
-  llvm.mlir.addressof @foo : !llvm.ptr<func<void ()>>
-  llvm.return
-}
-
-// -----
-
-llvm.mlir.global internal @foo(0: i32) : i32
-
-func.func @bar() {
-  // expected-error @+1 {{the type must be a pointer to the type of the referenced global}}
-  llvm.mlir.addressof @foo : !llvm.ptr<i64>
-  llvm.return
-}
-
-// -----
-
-llvm.func @foo()
-
-llvm.func @bar() {
-  // expected-error @+1 {{the type must be a pointer to the type of the referenced function}}
-  llvm.mlir.addressof @foo : !llvm.ptr<i8>
+  llvm.mlir.addressof @foo : !llvm.ptr
   llvm.return
 }
 
 // -----
 
 // expected-error @+2 {{block with no terminator}}
 llvm.mlir.global internal @g() : i64 {
   %c = llvm.mlir.constant(42 : i64) : i64
llvm.mlir.global internal @g(43 : i64) : i64 {
   llvm.return %c : i64
 }
 
 // -----
 
 llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
 func.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<i64>
+  llvm.mlir.addressof @g : !llvm.ptr
   llvm.return
 }
 
 // -----
 
 llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
-func.func @mismatch_addr_space() {
-  // expected-error @+1 {{pointer address space must match address space of the referenced global}}
-  llvm.mlir.addressof @g : !llvm.ptr<i64, 4>
-  llvm.return
-}
-// -----
 
-llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
-
-func.func @mismatch_addr_space_opaque() {
+func.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
 }
 
 // -----
 
 llvm.func @ctor() {

mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir

+// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -verify-diagnostics
+
+func.func @alloca_ptr_type_attr_non_opaque_ptr(%sz : i64) {
+  // expected-error@below {{unexpected 'elem_type' attribute when non-opaque pointer type is used}}
+  "llvm.alloca"(%sz) { elem_type = i32 } : (i64) -> !llvm.ptr<i32>
+}
+
+// -----
+
+func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr<f32>) {
+  // expected-error@+1 {{2 operands present, but expected 0}}
+  llvm.getelementptr %base[%pos] : () -> (!llvm.ptr<f32>)
+}
+
+// -----
+
+func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr<f32>) {
+  // expected-error@+1 {{op requires one result}}
+  llvm.getelementptr %base[%pos] : (!llvm.ptr<f32>, i64) -> ()
+}
+
+// -----
+
+func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr<f32>) {
+  // expected-error@+1 {{invalid kind of type specified}}
+  llvm.getelementptr %base[%pos] : !llvm.ptr<f32>
+}
+
+// -----
+
+func.func @gep_too_few_dynamic(%base : !llvm.ptr<f32>) {
+  // expected-error@+1 {{expected as many dynamic indices as specified in 'rawConstantIndices'}}
+  %1 = "llvm.getelementptr"(%base) {rawConstantIndices = array<i32: -2147483648>} : (!llvm.ptr<f32>) -> !llvm.ptr<f32>
+}
+
+// -----
+
+func.func @call_variadic(%callee : !llvm.ptr<func<i8 (i8, ...)>>, %arg : i8) {
+  // expected-error@+1 {{indirect calls to variadic functions are not supported}}
+  llvm.call %callee(%arg) : !llvm.ptr<func<i8 (i8, ...)>>, (i8) -> (i8)
+  llvm.return
+}
+
+// -----
+
+func.func @indirect_callee_arg_mismatch(%arg0 : i32, %callee : !llvm.ptr<func<void(i8)>>) {
+  // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}}
+  "llvm.call"(%callee, %arg0) : (!llvm.ptr<func<void(i8)>>, i32) -> ()
+  llvm.return
+}
+
+// -----
+
+func.func @indirect_callee_return_mismatch(%callee : !llvm.ptr<func<i8()>>) {
+  // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}}
+  "llvm.call"(%callee) : (!llvm.ptr<func<i8()>>) -> (i32)
+  llvm.return
+}
+
+// -----
+
+func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %i32 : i32) {
+  // expected-error@+1 {{expected LLVM IR element type for operand #0 to match type for operand #1}}
+  %0 = "llvm.atomicrmw"(%f32_ptr, %i32) {bin_op=11, ordering=1} : (!llvm.ptr<f32>, i32) -> i32
+  llvm.return
+}
+
+// -----
+
+func.func @cmpxchg_expected_ptr(%f32 : f32) {
+  // expected-error@+1 {{op operand #0 must be LLVM pointer to integer or LLVM pointer type}}
+  %0 = "llvm.cmpxchg"(%f32, %f32, %f32) {success_ordering=2,failure_ordering=2} : (f32, f32, f32) -> !llvm.struct<(f32, i1)>
+  llvm.return
+}
+
+// -----
+
+func.func @cmpxchg_mismatched_operands(%i64_ptr : !llvm.ptr<i64>, %i32 : i32) {
+  // expected-error@+1 {{expected LLVM IR element type for operand #0 to match type for all other operands}}
+  %0 = "llvm.cmpxchg"(%i64_ptr, %i32, %i32) {success_ordering=2,failure_ordering=2} : (!llvm.ptr<i64>, i32, i32) -> !llvm.struct<(i32, i1)>
+  llvm.return
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @bad_landingpad(%arg0: !llvm.ptr<ptr<i8>>) -> i32 attributes { personality = @__gxx_personality_v0} {
+  %0 = llvm.mlir.constant(3 : i32) : i32
+  %1 = llvm.mlir.constant(2 : i32) : i32
+  %2 = llvm.invoke @foo(%1) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1:  // pred: ^bb0
+  llvm.return %1 : i32
+^bb2:  // pred: ^bb0
+  // expected-error@+1 {{clause #0 is not a known constant - null, addressof, bitcast}}
+  %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr<ptr<i8>>) : !llvm.struct<(ptr<i8>, i32)>
+  llvm.return %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {
+  %0 = llvm.mlir.constant(1 : i32) : i32
+  %1 = llvm.alloca %0 x !llvm.ptr<i8> : (i32) -> !llvm.ptr<ptr<i8>>
+  // expected-note@+1 {{global addresses expected as operand to bitcast used in clauses for landingpad}}
+  %2 = llvm.bitcast %1 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
+  %3 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+  llvm.return %0 : i32
+^bb2: // pred: ^bb0
+  // expected-error@+1 {{constant clauses expected}}
+  %5 = llvm.landingpad (catch %2 : !llvm.ptr<i8>) : !llvm.struct<(ptr<i8>, i32)>
+  llvm.return %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {
+  %0 = llvm.mlir.constant(1 : i32) : i32
+  %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+  llvm.return %0 : i32
+^bb2: // pred: ^bb0
+  // expected-error@+1 {{landingpad instruction expects at least one clause or cleanup attribute}}
+  %2 = llvm.landingpad : !llvm.struct<(ptr<i8>, i32)>
+  llvm.return %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0 } {
+  %0 = llvm.mlir.constant(1 : i32) : i32
+  %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+  llvm.return %0 : i32
+^bb2: // pred: ^bb0
+  %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
+  // expected-error@+1 {{'llvm.resume' op expects landingpad value as operand}}
+  llvm.resume %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 {
+  %0 = llvm.mlir.constant(1 : i32) : i32
+  %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+  llvm.return %0 : i32
+^bb2: // pred: ^bb0
+  // expected-error@+1 {{llvm.landingpad needs to be in a function with a personality}}
+  %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
+  llvm.resume %2 : !llvm.struct<(ptr<i8>, i32)>
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32) {
+  // expected-error@+1 {{'nvvm.wmma.load' op expected source pointer in memory space 0, 1, 3}}
+  %0 = nvvm.wmma.load %arg0, %arg1
+    {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+    : (!llvm.ptr<5>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
+  // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
+  %0 = nvvm.wmma.load %arg0, %arg1
+  {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+  : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
+  // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
+ %0 = nvvm.wmma.load %arg0, %arg1
+ {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<b>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<3>, %arg1: i32) {
+  // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 4 elements of type 'vector<2xf16>'}}
+ %0 = nvvm.wmma.load %arg0, %arg1
+   {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<c>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+   : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
+
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32,
+                            %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,
+                            %arg4: vector<2 x f16>, %arg5: vector<2 xf16>) {
+  // expected-error@+1 {{'nvvm.wmma.store' op expected operands to be a source pointer in memory space 0, 1, 3}}
+  nvvm.wmma.store %arg0, %arg1, %arg2, %arg3, %arg4, %arg5
+    {eltype = #nvvm.mma_type<f16>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+    : !llvm.ptr<5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32>) {
+  // expected-error@+1 {{'nvvm.ldmatrix' op expected source pointer in memory space 3}}
+  %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32>) -> i32
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+  // expected-error@+1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4}}
+  %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+  // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is i32}}
+  %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32)>
+  llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+  // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is a structure of 4 elements of type i32}}
+  %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+  llvm.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 : !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} : !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>
+  return
+}
+
+// -----
+
+func.func @gep_out_of_bounds(%ptr: !llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, %idx: i64) {
+  // expected-error @below {{index 2 indexing a struct is out of bounds}}
+  llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, i64) -> !llvm.ptr<i32>
+  return
+}

mlir/test/Dialect/LLVMIR/invalid.mlir

func.func @alloca_missing_input_result_type(%size : i64) {
   // expected-error@+1 {{expected trailing function type with one argument and one result}}
   llvm.alloca %size x i32 : () -> ()
 }
 
 // -----
 
 func.func @alloca_missing_input_type() {
   // expected-error@+1 {{expected trailing function type with one argument and one result}}
-  llvm.alloca %size x i32 : () -> (!llvm.ptr<i32>)
+  llvm.alloca %size x i32 : () -> (!llvm.ptr)
 }
 
 // -----
 
 func.func @alloca_missing_result_type() {
   // expected-error@+1 {{expected trailing function type with one argument and one result}}
   llvm.alloca %size x i32 : (i64) -> ()
 }
 
 // -----
 
 func.func @alloca_non_function_type() {
   // expected-error@+1 {{expected trailing function type with one argument and one result}}
-  llvm.alloca %size x i32 : !llvm.ptr<i32>
+  llvm.alloca %size x i32 : !llvm.ptr
 }
 
 // -----
 
 func.func @alloca_non_integer_alignment() {
   // expected-error@+1 {{expected integer alignment}}
-  llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr<i32>
+  llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr
 }
 
 // -----
 
 func.func @alloca_opaque_ptr_no_type(%sz : i64) {
   // expected-error@below {{expected 'elem_type' attribute if opaque pointer type is used}}
   "llvm.alloca"(%sz) : (i64) -> !llvm.ptr
 }
 
 // -----
 
-func.func @alloca_ptr_type_attr_non_opaque_ptr(%sz : i64) {
-  // expected-error@below {{unexpected 'elem_type' attribute when non-opaque pointer type is used}}
-  "llvm.alloca"(%sz) { elem_type = i32 } : (i64) -> !llvm.ptr<i32>
-}
-
-// -----
-
-func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr) {
   // expected-error@+1 {{2 operands present, but expected 0}}
   llvm.getelementptr %base[%pos] : () -> ()
 }
 
 // -----
 
-func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) {
   // expected-error@+1 {{2 operands present, but expected 0}}
-  llvm.getelementptr %base[%pos] : () -> (!llvm.ptr<f32>)
+  llvm.getelementptr %base[%pos] : () -> (!llvm.ptr)
 }
 
 // -----
 
-func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) {
   // expected-error@+1 {{op requires one result}}
-  llvm.getelementptr %base[%pos] : (!llvm.ptr<f32>, i64) -> ()
+  llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> ()
 }
 
 // -----
 
-func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr) {
   // expected-error@+1 {{invalid kind of type specified}}
-  llvm.getelementptr %base[%pos] : !llvm.ptr<f32>
+  llvm.getelementptr %base[%pos] : !llvm.ptr
 }
 
 // -----
 
-func.func @gep_too_few_dynamic(%base : !llvm.ptr<f32>) {
+func.func @gep_too_few_dynamic(%base : !llvm.ptr) {
   // expected-error@+1 {{expected as many dynamic indices as specified in 'rawConstantIndices'}}
-  %1 = "llvm.getelementptr"(%base) {rawConstantIndices = array<i32: -2147483648>} : (!llvm.ptr<f32>) -> !llvm.ptr<f32>
+  %1 = "llvm.getelementptr"(%base) {elem_type = f32, rawConstantIndices = array<i32: -2147483648>} : (!llvm.ptr) -> !llvm.ptr
 }
 
 // -----
 
 func.func @load_non_llvm_type(%foo : memref<f32>) {
   // expected-error@+1 {{expected LLVM pointer type}}
   llvm.load %foo : memref<f32>
 }

 func.func @call_unknown_symbol() {
   // expected-error@+1 {{'llvm.call' op 'missing_callee' does not reference a symbol in the current scope}}
   llvm.call @missing_callee() : () -> ()
   llvm.return
 }
 
 // -----
 
-func.func @call_variadic(%callee : !llvm.ptr<func<i8 (i8, ...)>>, %arg : i8) {
-  // expected-error@+1 {{indirect calls to variadic functions are not supported}}
-  llvm.call %callee(%arg) : !llvm.ptr<func<i8 (i8, ...)>>, (i8) -> (i8)
-  llvm.return
-}
-
-// -----
-
 func.func private @standard_func_callee()
 
 func.func @call_non_llvm() {
   // expected-error@+1 {{'llvm.call' op 'standard_func_callee' does not reference a valid LLVM function}}
   llvm.call @standard_func_callee() : () -> ()
   llvm.return
 }
 

 func.func @callee_arg_mismatch(%arg0 : i32) {
   // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}}
   llvm.call @callee_func(%arg0) : (i32) -> ()
   llvm.return
 }
 
 // -----
 
-func.func @indirect_callee_arg_mismatch(%arg0 : i32, %callee : !llvm.ptr<func<void(i8)>>) {
-  // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}}
-  "llvm.call"(%callee, %arg0) : (!llvm.ptr<func<void(i8)>>, i32) -> ()
-  llvm.return
-}
-
-// -----
-
 llvm.func @callee_func() -> (i8)
 
 func.func @callee_return_mismatch() {
   // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}}
   %res = llvm.call @callee_func() : () -> (i32)
   llvm.return
 }
 
 // -----
 
-func.func @indirect_callee_return_mismatch(%callee : !llvm.ptr<func<i8()>>) {
-  // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}}
-  "llvm.call"(%callee) : (!llvm.ptr<func<i8()>>) -> (i32)
-  llvm.return
-}
-
-// -----
-
 func.func @call_too_many_results(%callee : !llvm.ptr) {
   // expected-error@+1 {{expected function with 0 or 1 result}}
   llvm.call %callee() : !llvm.ptr, () -> (i32, i32)
   llvm.return
 }
 
 // -----
 
llvm.func @func_result_mismatch(%arg0: f32) -> i32 {
   // expected-note@above {{when returning from function}}
   llvm.return %arg0 : f32
 }
 
 // -----
 
 func.func @constant_wrong_type() {
   // expected-error@+1 {{only supports integer, float, string or elements attributes}}
-  llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr<func<void ()>>
+  llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr
 }
 
 // -----
 
 func.func @constant_wrong_type_string() {
   // expected-error@below {{expected array type of 3 i8 elements for the string constant}}
-  llvm.mlir.constant("foo") : !llvm.ptr<i8>
+  llvm.mlir.constant("foo") : !llvm.ptr
 }
 
 // -----
 
 llvm.func @array_attribute_one_element() -> !llvm.struct<(f64, f64)> {
   // expected-error @+1 {{expected array attribute with two elements, representing a complex constant}}
   %0 = llvm.mlir.constant([1.0 : f64]) : !llvm.struct<(f64, f64)>
   llvm.return %0 : !llvm.struct<(f64, f64)>

 func.func @atomicrmw_expected_ptr(%f32 : f32) {
   // expected-error@+1 {{operand #0 must be LLVM pointer to floating point LLVM type or integer}}
   %0 = "llvm.atomicrmw"(%f32, %f32) {bin_op=11, ordering=1} : (f32, f32) -> f32
   llvm.return
 }
 
 // -----
 
-func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %i32 : i32) {
-  // expected-error@+1 {{expected LLVM IR element type for operand #0 to match type for operand #1}}
-  %0 = "llvm.atomicrmw"(%f32_ptr, %i32) {bin_op=11, ordering=1} : (!llvm.ptr<f32>, i32) -> i32
-  llvm.return
-}
-
-// -----
-
-func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) {
+func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) {
   // expected-error@+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>, f32) -> i32
+  %0 = "llvm.atomicrmw"(%f32_ptr, %f32) {bin_op=11, ordering=1} : (!llvm.ptr, f32) -> i32
   llvm.return
 }
 
 // -----
 
-func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr, %i32 : i32) {
   // expected-error@+1 {{expected LLVM IR floating point type}}
-  %0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : !llvm.ptr<i32>, i32
+  %0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : !llvm.ptr, i32
   llvm.return
 }
 
 // -----
 
-func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr<i1>, %i1 : i1) {
+func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
   // expected-error@+1 {{unexpected LLVM IR type for 'xchg' bin_op}}
-  %0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : !llvm.ptr<i1>, i1
+  %0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : !llvm.ptr, i1
   llvm.return
 }
 
 // -----
 
-func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) {
+func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {
   // expected-error@+1 {{expected LLVM IR integer type}}
-  %0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : !llvm.ptr<f32>, f32
+  %0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : !llvm.ptr, f32
   llvm.return
 }
 
 // -----
 
-func.func @cmpxchg_expected_ptr(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) {
+func.func @cmpxchg_expected_ptr(%f32 : f32) {
   // expected-error@+1 {{op operand #0 must be LLVM pointer to integer or LLVM pointer type}}
   %0 = "llvm.cmpxchg"(%f32, %f32, %f32) {success_ordering=2,failure_ordering=2} : (f32, f32, f32) -> !llvm.struct<(f32, i1)>
   llvm.return
 }
 
 // -----
 
-func.func @cmpxchg_mismatched_operands(%i64_ptr : !llvm.ptr<i64>, %i32 : i32) {
-  // expected-error@+1 {{expected LLVM IR element type for operand #0 to match type for all other operands}}
-  %0 = "llvm.cmpxchg"(%i64_ptr, %i32, %i32) {success_ordering=2,failure_ordering=2} : (!llvm.ptr<i64>, i32, i32) -> !llvm.struct<(i32, i1)>
-  llvm.return
-}
-
-// -----
-
 func.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) {
   // expected-error@+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
 }
 
 // -----
 
 func.func @cmpxchg_mismatched_result(%ptr : !llvm.ptr, %i64 : i64) {
   // expected-error@+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
 }
 
 // -----
 
-func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr<i1>, %i1 : i1) {
+func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
   // expected-error@+1 {{unexpected LLVM IR type}}
-  %0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : !llvm.ptr<i1>, i1
+  %0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : !llvm.ptr, i1
   llvm.return
 }
 
 // -----
 
-func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr, %i32 : i32) {
   // expected-error@+1 {{ordering must be at least 'monotonic'}}
-  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : !llvm.ptr<i32>, i32
+  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : !llvm.ptr, i32
   llvm.return
 }
 
 // -----
 
-func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr, %i32 : i32) {
   // expected-error@+1 {{ordering must be at least 'monotonic'}}
-  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : !llvm.ptr<i32>, i32
+  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : !llvm.ptr, i32
   llvm.return
 }
 
 // -----
 
-func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr, %i32 : i32) {
   // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}}
-  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : !llvm.ptr<i32>, i32
+  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : !llvm.ptr, i32
   llvm.return
 }
 
 // -----
 
-func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr, %i32 : i32) {
   // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}}
-  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr<i32>, i32
+  %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr, i32
   llvm.return
 }
 
 // -----
 
 llvm.func @foo(i32) -> i32
 llvm.func @__gxx_personality_v0(...) -> i32
 
-llvm.func @bad_landingpad(%arg0: !llvm.ptr<ptr<i8>>) -> i32 attributes { personality = @__gxx_personality_v0} {
+llvm.func @bad_landingpad(%arg0: !llvm.ptr) -> i32 attributes { personality = @__gxx_personality_v0} {
   %0 = llvm.mlir.constant(3 : i32) : i32
   %1 = llvm.mlir.constant(2 : i32) : i32
   %2 = llvm.invoke @foo(%1) to ^bb1 unwind ^bb2 : (i32) -> i32
 ^bb1:  // pred: ^bb0
   llvm.return %1 : i32
 ^bb2:  // pred: ^bb0
   // expected-error@+1 {{clause #0 is not a known constant - null, addressof, bitcast}}
-  %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr<ptr<i8>>) : !llvm.struct<(ptr<i8>, i32)>
+  %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr) : !llvm.struct<(ptr, i32)>
   llvm.return %0 : i32
 }
 
 // -----
 
 llvm.func @foo(i32) -> i32
 llvm.func @__gxx_personality_v0(...) -> i32
 
 llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {
   %0 = llvm.mlir.constant(1 : i32) : i32
-  %1 = llvm.alloca %0 x !llvm.ptr<i8> : (i32) -> !llvm.ptr<ptr<i8>>
+  %1 = llvm.alloca %0 x !llvm.ptr : (i32) -> !llvm.ptr
   // expected-note@+1 {{global addresses expected as operand to bitcast used in clauses for landingpad}}
-  %2 = llvm.bitcast %1 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
+  %2 = llvm.bitcast %1 : !llvm.ptr to !llvm.ptr
   %3 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
 ^bb1: // pred: ^bb0
   llvm.return %0 : i32
 ^bb2: // pred: ^bb0
   // expected-error@+1 {{constant clauses expected}}
-  %5 = llvm.landingpad (catch %2 : !llvm.ptr<i8>) : !llvm.struct<(ptr<i8>, i32)>
+  %5 = llvm.landingpad (catch %2 : !llvm.ptr) : !llvm.struct<(ptr, i32)>
   llvm.return %0 : i32
 }
 
 // -----
 
 llvm.func @foo(i32) -> i32
 llvm.func @__gxx_personality_v0(...) -> i32
 
 llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {
   %0 = llvm.mlir.constant(1 : i32) : i32
   %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
 ^bb1: // pred: ^bb0
   llvm.return %0 : i32
 ^bb2: // pred: ^bb0
   // expected-error@+1 {{landingpad instruction expects at least one clause or cleanup attribute}}
-  %2 = llvm.landingpad : !llvm.struct<(ptr<i8>, i32)>
+  %2 = llvm.landingpad : !llvm.struct<(ptr, i32)>
   llvm.return %0 : i32
 }
 
 // -----
 
 llvm.func @foo(i32) -> i32
 llvm.func @__gxx_personality_v0(...) -> i32
 
 llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0 } {
   %0 = llvm.mlir.constant(1 : i32) : i32
   %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
 ^bb1: // pred: ^bb0
   llvm.return %0 : i32
 ^bb2: // pred: ^bb0
-  %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
+  %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>
   // expected-error@+1 {{'llvm.resume' op expects landingpad value as operand}}
   llvm.resume %0 : i32
 }
 
 // -----
 
 llvm.func @foo(i32) -> i32
 
 llvm.func @caller(%arg0: i32) -> i32 {
   %0 = llvm.mlir.constant(1 : i32) : i32
   %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
 ^bb1: // pred: ^bb0
   llvm.return %0 : i32
 ^bb2: // pred: ^bb0
   // expected-error@+1 {{llvm.landingpad needs to be in a function with a personality}}
-  %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
-  llvm.resume %2 : !llvm.struct<(ptr<i8>, i32)>
+  %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>
+  llvm.resume %2 : !llvm.struct<(ptr, i32)>
 }
 
 // -----
 
 func.func @invalid_ordering_in_fence() {
   // expected-error @+1 {{can be given only acquire, release, acq_rel, and seq_cst orderings}}
   llvm.fence syncscope("agent") monotonic
 }
module {
   }
   llvm.metadata @other_metadata {
     llvm.alias_scope_domain @domain
   }
 }
 
 // -----
 
-llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<i32, 5>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32) {
   // expected-error@+1 {{'nvvm.wmma.load' op expected source pointer in memory space 0, 1, 3}}
   %0 = nvvm.wmma.load %arg0, %arg1
     {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
-    : (!llvm.ptr<i32, 5>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+    : (!llvm.ptr<5>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
   // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
   %0 = nvvm.wmma.load %arg0, %arg1
   {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
-  : (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+  : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
   // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
  %0 = nvvm.wmma.load %arg0, %arg1
  {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<b>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
- : (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<3>, %arg1: i32) {
   // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 4 elements of type 'vector<2xf16>'}}
  %0 = nvvm.wmma.load %arg0, %arg1
    {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<c>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
-   : (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
+   : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
 
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<i32, 5>, %arg1: i32,
+llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32,
                             %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,
                             %arg4: vector<2 x f16>, %arg5: vector<2 xf16>) {
   // expected-error@+1 {{'nvvm.wmma.store' op expected operands to be a source pointer in memory space 0, 1, 3}}
   nvvm.wmma.store %arg0, %arg1, %arg2, %arg3, %arg4, %arg5
     {eltype = #nvvm.mma_type<f16>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
-    : !llvm.ptr<i32, 5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
+    : !llvm.ptr<5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
   llvm.return
 }
 
 // -----
 
 llvm.func @gpu_wmma_mma_op_invalid_operands(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,
                         %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,
                         %arg4: vector<2 x f16>, %arg5: vector<2 x f16>,
llvm.func @gpu_wmma_mma_op_invalid_result(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,
   %0 = nvvm.wmma.mma %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19, %arg20, %arg21, %arg22, %arg23
     {eltypeA = #nvvm.mma_type<f16>, eltypeB = #nvvm.mma_type<f32>, k = 16 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
     : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, vector<2xf16>)>
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr) {
   // expected-error@+1 {{'nvvm.ldmatrix' op expected source pointer in memory space 3}}
-  %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32>) -> i32
+  %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr) -> i32
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) {
   // expected-error@+1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4}}
-  %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
+  %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> i32
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) {
   // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is i32}}
-  %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32)>
+  %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32)>
   llvm.return
 }
 
 // -----
 
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) {
   // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is a structure of 4 elements of type i32}}
-  %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+  %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>
   llvm.return
 }
 
 // -----
 
 llvm.func @caller() {
   // expected-error @below {{expected function call to produce a value}}
   llvm.call @callee() : () -> ()

 func.func @bitcast(%arg0: vector<2x3xf32>) {
   // expected-error @below {{op operand #0 must be LLVM-compatible non-aggregate type}}
   llvm.bitcast %arg0 : vector<2x3xf32> to vector<2x3xi32>
   return
 }
 
 // -----
 
-func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+func.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {
   // expected-error @below {{expected byte size to be either 4, 8 or 16.}}
-  nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+  nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr<3>, !llvm.ptr<1>
   return
 }
 
 // -----
 
-func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+func.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {
   // expected-error @below {{bypass l1 is only support for 16 bytes copy.}}
-  nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+  nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1} : !llvm.ptr<3>, !llvm.ptr<1>
   return
 }
 
 // -----
 
-func.func @gep_struct_variable(%arg0: !llvm.ptr<struct<(i32)>>, %arg1: i32, %arg2: i32) {
+func.func @gep_struct_variable(%arg0: !llvm.ptr, %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>
+  llvm.getelementptr %arg0[%arg1, %arg1] : (!llvm.ptr, i32, i32) -> !llvm.ptr, !llvm.struct<(i32)>
   return
 }
 
 // -----
 
-func.func @gep_out_of_bounds(%ptr: !llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, %idx: i64) {
+func.func @gep_out_of_bounds(%ptr: !llvm.ptr, %idx: i64) {
   // expected-error @below {{index 2 indexing a struct is out of bounds}}
-  llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, i64) -> !llvm.ptr<i32>
+  llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(i32, struct<(i32, f32)>)>
   return
 }
 
 // -----
 
 func.func @non_splat_shuffle_on_scalable_vector(%arg0: vector<[4]xf32>) {
   // expected-error@below {{expected a splat operation for scalable vectors}}
   %0 = llvm.shufflevector %arg0, %arg0 [0, 0, 0, 1] : vector<[4]xf32>
   return
 }
 
 // -----
 
 llvm.mlir.global internal @side_effecting_global() : !llvm.struct<(i8)> {
   %0 = llvm.mlir.constant(1 : i64) : i64
   // expected-error@below {{ops with side effects not allowed in global initializers}}
-  %1 = llvm.alloca %0 x !llvm.struct<(i8)> : (i64) -> !llvm.ptr<struct<(i8)>>
-  %2 = llvm.load %1 : !llvm.ptr<struct<(i8)>>
+  %1 = llvm.alloca %0 x !llvm.struct<(i8)> : (i64) -> !llvm.ptr
+  %2 = llvm.load %1 : !llvm.ptr -> !llvm.struct<(i8)>
   llvm.return %2 : !llvm.struct<(i8)>
 }
 
 // -----
 
 func.func @insert_vector_invalid_source_vector_size(%arg0 : vector<16385xi8>, %arg1 : vector<[16]xi8>) {
   // expected-error@+1 {{op failed to verify that vectors are not bigger than 2^17 bits.}}
   %0 = llvm.intr.vector.insert %arg0, %arg1[0] : vector<16385xi8> into vector<[16]xi8>

mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir

+// RUN: mlir-opt --test-data-layout-query --split-input-file --verify-diagnostics %s | FileCheck %s
+
+module {
+  // CHECK: @no_spec
+  func.func @no_spec() {
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 0
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<i8>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 0
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<i32>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 0
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<bf16>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 0
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 0
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 0
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 0
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<5>
+    return
+  }
+}
+
+// -----
+
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+  #dlti.dl_entry<!llvm.ptr<i8>, dense<[32, 32, 64]> : vector<3xi32>>,
+  #dlti.dl_entry<!llvm.ptr<i8, 5>, dense<[64, 64, 64]> : vector<3xi32>>,
+  #dlti.dl_entry<!llvm.ptr<4>, dense<[32, 64, 64]> : vector<3xi32>>,
+  #dlti.dl_entry<"dlti.alloca_memory_space", 5 : ui32>
+>} {
+  // CHECK: @spec
+  func.func @spec() {
+    // CHECK: alignment = 4
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 32
+    // CHECK: preferred = 8
+    // CHECK: size = 4
+    "test.data_layout_query"() : () -> !llvm.ptr<i8>
+    // CHECK: alignment = 4
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 32
+    // CHECK: preferred = 8
+    // CHECK: size = 4
+    "test.data_layout_query"() : () -> !llvm.ptr<i32>
+    // CHECK: alignment = 4
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 32
+    // CHECK: preferred = 8
+    // CHECK: size = 4
+    "test.data_layout_query"() : () -> !llvm.ptr<bf16>
+    // CHECK: alignment = 4
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 32
+    // CHECK: preferred = 8
+    // CHECK: size = 4
+    "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
+    // CHECK: alignment = 4
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 32
+    // CHECK: preferred = 8
+    // CHECK: size = 4
+    "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 64
+    // CHECK: preferred = 8
+    // CHECK: size = 8
+    "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+    // CHECK: alignment = 4
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 32
+    // CHECK: preferred = 8
+    // CHECK: size = 4
+    "test.data_layout_query"() : () -> !llvm.ptr<3>
+    // CHECK: alignment = 8
+    // CHECK: alloca_memory_space = 5
+    // CHECK: bitsize = 32
+    // CHECK: preferred = 8
+    // CHECK: size = 4
+  "test.data_layout_query"() : () -> !llvm.ptr<4>
+    return
+  }
+}
+
+// -----
+
+// expected-error@below {{unexpected layout attribute for pointer to 'i32'}}
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+  #dlti.dl_entry<!llvm.ptr<i32>, dense<[64, 64, 64]> : vector<3xi32>>
+>} {
+  func.func @pointer() {
+    return
+  }
+}
+
+// -----
+
+// expected-error@below {{expected layout attribute for '!llvm.ptr<i8>' to be a dense integer elements attribute with 3 or 4 elements}}
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+  #dlti.dl_entry<!llvm.ptr<i8>, dense<[64.0, 64.0, 64.0]> : vector<3xf32>>
+>} {
+  func.func @pointer() {
+    return
+  }
+}
+
+// -----
+
+// expected-error@below {{preferred alignment is expected to be at least as large as ABI alignment}}
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+  #dlti.dl_entry<!llvm.ptr<i8>, dense<[64, 64, 32]> : vector<3xi32>>
+>} {
+  func.func @pointer() {
+    return
+  }
+}

mlir/test/Dialect/LLVMIR/layout.mlir

 // RUN: mlir-opt --test-data-layout-query --split-input-file --verify-diagnostics %s | FileCheck %s
 
 module {
   // CHECK: @no_spec
   func.func @no_spec() {
+    "test.data_layout_query"() : () -> !llvm.ptr
     // CHECK: alignment = 8
     // CHECK: alloca_memory_space = 0
     // CHECK: bitsize = 64
     // CHECK: preferred = 8
     // CHECK: size = 8
-    "test.data_layout_query"() : () -> !llvm.ptr<i8>
-    // CHECK: alignment = 8
-    // CHECK: alloca_memory_space = 0
-    // CHECK: bitsize = 64
-    // CHECK: preferred = 8
-    // CHECK: size = 8
-    "test.data_layout_query"() : () -> !llvm.ptr<i32>
-    // CHECK: alignment = 8
-    // CHECK: alloca_memory_space = 0
-    // CHECK: bitsize = 64
-    // CHECK: preferred = 8
-    // CHECK: size = 8
-    "test.data_layout_query"() : () -> !llvm.ptr<bf16>
-    // CHECK: alignment = 8
-    // CHECK: alloca_memory_space = 0
-    // CHECK: bitsize = 64
-    // CHECK: preferred = 8
-    // CHECK: size = 8
-    "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
-    // CHECK: alignment = 8
-    // CHECK: alloca_memory_space = 0
-    // CHECK: bitsize = 64
-    // CHECK: preferred = 8
-    // CHECK: size = 8
-    "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
-    // CHECK: alignment = 8
-    // CHECK: alloca_memory_space = 0
-    // CHECK: bitsize = 64
-    // CHECK: preferred = 8
-    // CHECK: size = 8
-    "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+    "test.data_layout_query"() : () -> !llvm.ptr<3>
     // CHECK: alignment = 8
     // CHECK: alloca_memory_space = 0
     // CHECK: bitsize = 64
     // CHECK: preferred = 8
     // CHECK: size = 8
     "test.data_layout_query"() : () -> !llvm.ptr<5>
     return
   }
 }
 
 // -----
 
 module attributes { dlti.dl_spec = #dlti.dl_spec<
-  #dlti.dl_entry<!llvm.ptr<i8>, dense<[32, 32, 64]> : vector<3xi32>>,
-  #dlti.dl_entry<!llvm.ptr<i8, 5>, dense<[64, 64, 64]> : vector<3xi32>>,
+  #dlti.dl_entry<!llvm.ptr, dense<[32, 32, 64]> : vector<3xi32>>,
+  #dlti.dl_entry<!llvm.ptr<5>, dense<[64, 64, 64]> : vector<3xi32>>,
   #dlti.dl_entry<!llvm.ptr<4>, dense<[32, 64, 64]> : vector<3xi32>>,
   #dlti.dl_entry<"dlti.alloca_memory_space", 5 : ui32>
 >} {
   // CHECK: @spec
   func.func @spec() {
     // CHECK: alignment = 4
     // CHECK: alloca_memory_space = 5
     // CHECK: bitsize = 32
     // CHECK: preferred = 8
     // CHECK: size = 4
-    "test.data_layout_query"() : () -> !llvm.ptr<i8>
+    "test.data_layout_query"() : () -> !llvm.ptr
     // CHECK: alignment = 4
     // CHECK: alloca_memory_space = 5
     // CHECK: bitsize = 32
     // CHECK: preferred = 8
     // CHECK: size = 4
-    "test.data_layout_query"() : () -> !llvm.ptr<i32>
-    // CHECK: alignment = 4
-    // CHECK: alloca_memory_space = 5
-    // CHECK: bitsize = 32
-    // CHECK: preferred = 8
-    // CHECK: size = 4
-    "test.data_layout_query"() : () -> !llvm.ptr<bf16>
-    // CHECK: alignment = 4
-    // CHECK: alloca_memory_space = 5
-    // CHECK: bitsize = 32
-    // CHECK: preferred = 8
-    // CHECK: size = 4
-    "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
-    // CHECK: alignment = 4
-    // CHECK: alloca_memory_space = 5
-    // CHECK: bitsize = 32
-    // CHECK: preferred = 8
-    // CHECK: size = 4
-    "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
+    "test.data_layout_query"() : () -> !llvm.ptr<3>
     // CHECK: alignment = 8
     // CHECK: alloca_memory_space = 5
     // CHECK: bitsize = 64
     // CHECK: preferred = 8
     // CHECK: size = 8
-    "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+    "test.data_layout_query"() : () -> !llvm.ptr<5>
     // CHECK: alignment = 4
     // CHECK: alloca_memory_space = 5
     // CHECK: bitsize = 32
     // CHECK: preferred = 8
     // CHECK: size = 4
     "test.data_layout_query"() : () -> !llvm.ptr<3>
     // CHECK: alignment = 8
     // CHECK: alloca_memory_space = 5
     // CHECK: bitsize = 32
     // CHECK: preferred = 8
     // CHECK: size = 4
   "test.data_layout_query"() : () -> !llvm.ptr<4>
     return
   }
 }
 
 // -----
 
-// expected-error@below {{unexpected layout attribute for pointer to 'i32'}}
-module attributes { dlti.dl_spec = #dlti.dl_spec<
-  #dlti.dl_entry<!llvm.ptr<i32>, dense<[64, 64, 64]> : vector<3xi32>>
->} {
-  func.func @pointer() {
-    return
-  }
-}
-
-// -----
-
-// expected-error@below {{expected layout attribute for '!llvm.ptr<i8>' to be a dense integer elements attribute with 3 or 4 elements}}
+// expected-error@below {{expected layout attribute for '!llvm.ptr' to be a dense integer elements attribute with 3 or 4 elements}}
 module attributes { dlti.dl_spec = #dlti.dl_spec<
-  #dlti.dl_entry<!llvm.ptr<i8>, dense<[64.0, 64.0, 64.0]> : vector<3xf32>>
+  #dlti.dl_entry<!llvm.ptr, dense<[64.0, 64.0, 64.0]> : vector<3xf32>>
 >} {
   func.func @pointer() {
     return
   }
 }
 
 // -----
 
 // expected-error@below {{preferred alignment is expected to be at least as large as ABI alignment}}
 module attributes { dlti.dl_spec = #dlti.dl_spec<
-  #dlti.dl_entry<!llvm.ptr<i8>, dense<[64, 64, 32]> : vector<3xi32>>
+  #dlti.dl_entry<!llvm.ptr, dense<[64, 64, 32]> : vector<3xi32>>
 >} {
   func.func @pointer() {
     return
   }
 }
 
 // -----
 

mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir

+// RUN: mlir-opt %s -split-input-file -verify-diagnostics | FileCheck %s
+
+// CHECK-LABEL: @nvvm_wmma_load_tf32
+func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr<i32>, %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}
+    : (!llvm.ptr<i32>) -> !llvm.struct<(i32, i32, i32, i32)>
+  llvm.return %0 : !llvm.struct<(i32, i32, i32, i32)>
+}
+
+// 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 : !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} : !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
+llvm.func @ld_matrix(%arg0: !llvm.ptr<i32, 3>) {
+  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 1 : i32} : (!llvm.ptr<i32, 3>) -> i32
+  %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
+  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 2 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+  %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 4 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
+  %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
+  llvm.return
+}
+
+// CHECK-LABEL: llvm.func @redux_sync
+llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {
+  // CHECK: nvvm.redux.sync  add %{{.*}}
+  %r1 = nvvm.redux.sync add %value, %offset : i32 -> i32
+  // CHECK: nvvm.redux.sync  max %{{.*}}
+  %r2 = nvvm.redux.sync max %value, %offset : i32 -> i32
+  // CHECK: nvvm.redux.sync  min %{{.*}}
+  %r3 = nvvm.redux.sync min %value, %offset : i32 -> i32
+  // CHECK: nvvm.redux.sync  umax %{{.*}}
+  %r5 = nvvm.redux.sync umax %value, %offset : i32 -> i32
+  // CHECK: nvvm.redux.sync  umin %{{.*}}
+  %r6 = nvvm.redux.sync umin %value, %offset : i32 -> i32
+  // CHECK: nvvm.redux.sync  and %{{.*}}
+  %r7 = nvvm.redux.sync and %value, %offset : i32 -> i32
+  // CHECK: nvvm.redux.sync  or %{{.*}}
+  %r8 = nvvm.redux.sync or %value, %offset : i32 -> i32
+  // CHECK: nvvm.redux.sync  xor %{{.*}}
+  %r9 = nvvm.redux.sync xor %value, %offset : i32 -> i32
+  llvm.return %r1 : i32
+}

mlir/test/Dialect/LLVMIR/nvvm.mlir

%0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]
     {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,
      multiplicandAPtxType = #nvvm.mma_type<s4>, multiplicandBPtxType = #nvvm.mma_type<s4>,
      intOverflowBehavior=#nvvm.mma_int_overflow<wrapped>,
      shape = #nvvm.shape<m = 16, n = 8, k = 32>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>
   llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>
 }
 
 // CHECK-LABEL: @nvvm_wmma_load_tf32
-func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr<i32>, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
+func.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}
-    : (!llvm.ptr<i32>) -> !llvm.struct<(i32, i32, i32, i32)>
+    : (!llvm.ptr) -> !llvm.struct<(i32, i32, i32, i32)>
   llvm.return %0 : !llvm.struct<(i32, i32, i32, i32)>
 }
 
 // CHECK-LABEL: @nvvm_wmma_mma
 func.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)> {
   // CHECK: nvvm.wmma.mma {{.*}} {eltypeA = #nvvm.mma_type<tf32>, eltypeB = #nvvm.mma_type<f32>, k = 8 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
   %r = nvvm.wmma.mma %0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15
     {eltypeA = #nvvm.mma_type<tf32>, eltypeB = #nvvm.mma_type<f32>, k = 8 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
     : (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>) {
+llvm.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {
 // CHECK:  nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16
-  nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+  nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr<3>, !llvm.ptr<1>
 // CHECK:  nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16 {bypass_l1}
-  nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+  nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr<3>, !llvm.ptr<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
-llvm.func @ld_matrix(%arg0: !llvm.ptr<i32, 3>) {
-  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 1 : i32} : (!llvm.ptr<i32, 3>) -> i32
-  %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
-  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 2 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
-  %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
-  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 4 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
-  %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
+llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {
+  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 1 : i32} : (!llvm.ptr<3>) -> i32
+  %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> i32
+  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 2 : i32} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>
+  %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>
+  // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 4 : i32} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
+  %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
   llvm.return
 }
 
 // CHECK-LABEL: llvm.func @redux_sync
-llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {  
+llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {
   // CHECK: nvvm.redux.sync  add %{{.*}}
   %r1 = nvvm.redux.sync add %value, %offset : i32 -> i32
   // CHECK: nvvm.redux.sync  max %{{.*}}
   %r2 = nvvm.redux.sync max %value, %offset : i32 -> i32
   // CHECK: nvvm.redux.sync  min %{{.*}}
   %r3 = nvvm.redux.sync min %value, %offset : i32 -> i32
   // CHECK: nvvm.redux.sync  umax %{{.*}}
   %r5 = nvvm.redux.sync umax %value, %offset : i32 -> i32
   // CHECK: nvvm.redux.sync  umin %{{.*}}
   %r6 = nvvm.redux.sync umin %value, %offset : i32 -> i32
   // CHECK: nvvm.redux.sync  and %{{.*}}
-  %r7 = nvvm.redux.sync and %value, %offset : i32 -> i32  
+  %r7 = nvvm.redux.sync and %value, %offset : i32 -> i32
   // CHECK: nvvm.redux.sync  or %{{.*}}
-  %r8 = nvvm.redux.sync or %value, %offset : i32 -> i32  
+  %r8 = nvvm.redux.sync or %value, %offset : i32 -> i32
   // CHECK: nvvm.redux.sync  xor %{{.*}}
   %r9 = nvvm.redux.sync xor %value, %offset : i32 -> i32
   llvm.return %r1 : i32
 }
 
 
 // -----
 
 // expected-error@below {{attribute attached to unexpected op}}
 func.func private @expected_llvm_func() attributes { nvvm.kernel }

mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir

+// RUN: mlir-opt %s -split-input-file -verify-diagnostics
+
+// Argument attributes
+
+// expected-error@below {{"llvm.sret" attribute attached to LLVM pointer argument of different type}}
+llvm.func @invalid_sret_attr_type(%0 : !llvm.ptr<f32> {llvm.sret = !llvm.struct<(i32)>})

mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir

 
 // -----
 
 // expected-error@below {{"llvm.sret" attribute attached to non-pointer LLVM type}}
 llvm.func @invalid_sret_arg_type(%0 : i32 {llvm.sret = !llvm.struct<(i32)>})
 
 // -----
 
-// expected-error@below {{"llvm.sret" attribute attached to LLVM pointer argument of different type}}
-llvm.func @invalid_sret_attr_type(%0 : !llvm.ptr<f32> {llvm.sret = !llvm.struct<(i32)>})
-
-// -----
-
 // expected-error@below {{"llvm.byval" attribute attached to non-pointer LLVM type}}
 llvm.func @invalid_byval_arg_type(%0 : i32 {llvm.byval = !llvm.struct<(i32)>})
 
 // -----
 
 // expected-error@below {{"llvm.byval" attribute attached to LLVM pointer argument of different type}}
 llvm.func @invalid_byval_attr_type(%0 : !llvm.ptr<!llvm.struct<(f32)>> {llvm.byval = !llvm.struct<(i32)>})
 

mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir

+// RUN: mlir-opt --allow-unregistered-dialect -split-input-file -verify-diagnostics %s
+
+func.func @void_pointer() {
+  // expected-error @+1 {{invalid pointer element type}}
+  "some.op"() : () -> !llvm.ptr<void>
+}
+
+// -----
+
+func.func @repeated_struct_name() {
+  "some.op"() : () -> !llvm.struct<"a", (ptr<struct<"a">>)>
+  // expected-error @+1 {{identified type already used with a different body}}
+  "some.op"() : () -> !llvm.struct<"a", (i32)>
+}
+
+// -----
+
+func.func @dynamic_vector() {
+  // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
+  "some.op"() : () -> !llvm.vec<? x ptr<f32>>
+}
+
+// -----
+
+func.func @dynamic_scalable_vector() {
+  // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
+  "some.op"() : () -> !llvm.vec<?x? x ptr<f32>>
+}
+
+// -----
+
+func.func @unscalable_vector() {
+  // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
+  "some.op"() : () -> !llvm.vec<4x4 x ptr<i32>>
+}
+
+// -----
+
+func.func @zero_vector() {
+  // expected-error @+1 {{the number of vector elements must be positive}}
+  "some.op"() : () -> !llvm.vec<0 x ptr<i32>>
+}

mlir/test/Dialect/LLVMIR/types-invalid.mlir

 
 func.func @function_taking_function() {
   // expected-error @+1 {{invalid function argument type}}
   "some.op"() : () -> !llvm.func<void (func<void ()>)>
 }
 
 // -----
 
-func.func @void_pointer() {
-  // expected-error @+1 {{invalid pointer element type}}
-  "some.op"() : () -> !llvm.ptr<void>
-}
-
-// -----
-
 func.func @repeated_struct_name() {
-  "some.op"() : () -> !llvm.struct<"a", (ptr<struct<"a">>)>
+  "some.op"() : () -> !llvm.struct<"a", (ptr)>
   // expected-error @+1 {{identified type already used with a different body}}
   "some.op"() : () -> !llvm.struct<"a", (i32)>
 }
 
 // -----
 
 func.func @repeated_struct_name_packed() {
   "some.op"() : () -> !llvm.struct<"a", packed (i32)>
func.func @identified_struct_with_void() {
   // expected-error @+1 {{invalid LLVM structure element type}}
   "some.op"() : () -> !llvm.struct<"a", (void)>
 }
 
 // -----
 
 func.func @dynamic_vector() {
   // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
-  "some.op"() : () -> !llvm.vec<? x ptr<f32>>
+  "some.op"() : () -> !llvm.vec<? x ptr>
 }
 
 // -----
 
 func.func @dynamic_scalable_vector() {
   // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
-  "some.op"() : () -> !llvm.vec<?x? x ptr<f32>>
+  "some.op"() : () -> !llvm.vec<?x? x ptr>
 }
 
 // -----
 
 func.func @unscalable_vector() {
   // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
-  "some.op"() : () -> !llvm.vec<4x4 x ptr<i32>>
+  "some.op"() : () -> !llvm.vec<4x4 x ptr>
 }
 
 // -----
 
 func.func @zero_vector() {
   // expected-error @+1 {{the number of vector elements must be positive}}
-  "some.op"() : () -> !llvm.vec<0 x ptr<i32>>
+  "some.op"() : () -> !llvm.vec<0 x ptr>
 }
 
 // -----
 
 func.func @nested_vector() {
   // expected-error @+1 {{invalid vector element type}}
   "some.op"() : () -> !llvm.vec<2 x vector<2xi32>>
 }

mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir

 // RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file | mlir-opt -allow-unregistered-dialect | FileCheck %s
 
-// CHECK-LABEL: @primitive
-func.func @primitive() {
-  // CHECK: !llvm.void
-  "some.op"() : () -> !llvm.void
-  // CHECK: !llvm.ppc_fp128
-  "some.op"() : () -> !llvm.ppc_fp128
-  // CHECK: !llvm.x86_mmx
-  "some.op"() : () -> !llvm.x86_mmx
-  // CHECK: !llvm.token
-  "some.op"() : () -> !llvm.token
-  // CHECK: !llvm.label
-  "some.op"() : () -> !llvm.label
-  // CHECK: !llvm.metadata
-  "some.op"() : () -> !llvm.metadata
-  return
-}
-
-// CHECK-LABEL: @func
-func.func @func() {
-  // CHECK: !llvm.func<void ()>
-  "some.op"() : () -> !llvm.func<void ()>
-  // CHECK: !llvm.func<void (i32)>
-  "some.op"() : () -> !llvm.func<void (i32)>
-  // CHECK: !llvm.func<i32 ()>
-  "some.op"() : () -> !llvm.func<i32 ()>
-  // CHECK: !llvm.func<i32 (f16, bf16, f32, f64)>
-  "some.op"() : () -> !llvm.func<i32 (f16, bf16, f32, f64)>
-  // CHECK: !llvm.func<i32 (i32, i32)>
-  "some.op"() : () -> !llvm.func<i32 (i32, i32)>
-  // CHECK: !llvm.func<void (...)>
-  "some.op"() : () -> !llvm.func<void (...)>
-  // CHECK: !llvm.func<void (i32, i32, ...)>
-  "some.op"() : () -> !llvm.func<void (i32, i32, ...)>
-  return
-}
-
-// CHECK-LABEL: @integer
-func.func @integer() {
-  // CHECK: i1
-  "some.op"() : () -> i1
-  // CHECK: i8
-  "some.op"() : () -> i8
-  // CHECK: i16
-  "some.op"() : () -> i16
-  // CHECK: i32
-  "some.op"() : () -> i32
-  // CHECK: i64
-  "some.op"() : () -> i64
-  // CHECK: i57
-  "some.op"() : () -> i57
-  // CHECK: i129
-  "some.op"() : () -> i129
-  return
-}
-
 // CHECK-LABEL: @ptr
 func.func @ptr() {
   // CHECK: !llvm.ptr<i8>
   "some.op"() : () -> !llvm.ptr<i8>
   // CHECK: !llvm.ptr<f32>
   "some.op"() : () -> !llvm.ptr<f32>
   // CHECK: !llvm.ptr<ptr<i8>>
   "some.op"() : () -> !llvm.ptr<ptr<i8>>
func.func @array() {
   "some.op"() : () -> !llvm.array<8 x f32>
   // CHECK: !llvm.array<10 x ptr<i32, 4>>
   "some.op"() : () -> !llvm.array<10 x ptr<i32, 4>>
   // CHECK: !llvm.array<10 x array<4 x f32>>
   "some.op"() : () -> !llvm.array<10 x array<4 x f32>>
   return
 }
 
-// CHECK-LABEL: @literal_struct
-func.func @literal_struct() {
-  // CHECK: !llvm.struct<()>
-  "some.op"() : () -> !llvm.struct<()>
-  // CHECK: !llvm.struct<(i32)>
-  "some.op"() : () -> !llvm.struct<(i32)>
-  // CHECK: !llvm.struct<(f32, i32)>
-  "some.op"() : () -> !llvm.struct<(f32, i32)>
-  // CHECK: !llvm.struct<(struct<(i32)>)>
-  "some.op"() : () -> !llvm.struct<(struct<(i32)>)>
-  // CHECK: !llvm.struct<(i32, struct<(i32)>, f32)>
-  "some.op"() : () -> !llvm.struct<(i32, struct<(i32)>, f32)>
-
-  // CHECK: !llvm.struct<packed ()>
-  "some.op"() : () -> !llvm.struct<packed ()>
-  // CHECK: !llvm.struct<packed (i32)>
-  "some.op"() : () -> !llvm.struct<packed (i32)>
-  // CHECK: !llvm.struct<packed (f32, i32)>
-  "some.op"() : () -> !llvm.struct<packed (f32, i32)>
-  // CHECK: !llvm.struct<packed (f32, i32)>
-  "some.op"() : () -> !llvm.struct<packed (f32, i32)>
-  // CHECK: !llvm.struct<packed (struct<(i32)>)>
-  "some.op"() : () -> !llvm.struct<packed (struct<(i32)>)>
-  // CHECK: !llvm.struct<packed (i32, struct<(i32, i1)>, f32)>
-  "some.op"() : () -> !llvm.struct<packed (i32, struct<(i32, i1)>, f32)>
-
-  // CHECK: !llvm.struct<(struct<packed (i32)>)>
-  "some.op"() : () -> !llvm.struct<(struct<packed (i32)>)>
-  // CHECK: !llvm.struct<packed (struct<(i32)>)>
-  "some.op"() : () -> !llvm.struct<packed (struct<(i32)>)>
-  return
-}
-
 // CHECK-LABEL: @identified_struct
 func.func @identified_struct() {
   // CHECK: !llvm.struct<"empty", ()>
   "some.op"() : () -> !llvm.struct<"empty", ()>
   // CHECK: !llvm.struct<"opaque", opaque>
   "some.op"() : () -> !llvm.struct<"opaque", opaque>
   // CHECK: !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
   "some.op"() : () -> !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
func.func @identified_struct() {
   "some.op"() : () -> !llvm.struct<"struct-of-arrays", (array<10 x i32>)>
   // CHECK: !llvm.array<10 x struct<"array-of-structs", (i32)>>
   "some.op"() : () -> !llvm.array<10 x struct<"array-of-structs", (i32)>>
   // CHECK: !llvm.ptr<struct<"ptr-to-struct", (i8)>>
   "some.op"() : () -> !llvm.ptr<struct<"ptr-to-struct", (i8)>>
   return
 }
 
-func.func @verbose() {
-  // CHECK: !llvm.struct<(i64, struct<(f32)>)>
-  "some.op"() : () -> !llvm.struct<(i64, !llvm.struct<(f32)>)>
-  return
-}
-
 // CHECK-LABEL: @ptr_elem_interface
 // CHECK-COUNT-3: !llvm.ptr<!test.smpla>
 // CHECK: llvm.mlir.undef : !llvm.ptr<!test.smpla>
 func.func @ptr_elem_interface(%arg0: !llvm.ptr<!test.smpla>) {
   %0 = llvm.load %arg0 : !llvm.ptr<!test.smpla>
   llvm.store %0, %arg0 : !llvm.ptr<!test.smpla>
   llvm.mlir.undef : !llvm.ptr<!test.smpla>
   return

mlir/test/Dialect/LLVMIR/types.mlir

func.func @integer() {
   "some.op"() : () -> i57
   // CHECK: i129
   "some.op"() : () -> i129
   return
 }
 
 // CHECK-LABEL: @ptr
 func.func @ptr() {
-  // CHECK: !llvm.ptr<i8>
-  "some.op"() : () -> !llvm.ptr<i8>
-  // CHECK: !llvm.ptr<f32>
-  "some.op"() : () -> !llvm.ptr<f32>
-  // CHECK: !llvm.ptr<ptr<i8>>
-  "some.op"() : () -> !llvm.ptr<ptr<i8>>
-  // CHECK: !llvm.ptr<ptr<ptr<ptr<ptr<i8>>>>>
-  "some.op"() : () -> !llvm.ptr<ptr<ptr<ptr<ptr<i8>>>>>
-  // CHECK: !llvm.ptr<i8>
-  "some.op"() : () -> !llvm.ptr<i8, 0>
-  // CHECK: !llvm.ptr<i8, 1>
-  "some.op"() : () -> !llvm.ptr<i8, 1>
-  // CHECK: !llvm.ptr<i8, 42>
-  "some.op"() : () -> !llvm.ptr<i8, 42>
-  // CHECK: !llvm.ptr<ptr<i8, 42>, 9>
-  "some.op"() : () -> !llvm.ptr<ptr<i8, 42>, 9>
   // CHECK: !llvm.ptr
   "some.op"() : () -> !llvm.ptr
+  // CHECK: !llvm.ptr
+  "some.op"() : () -> !llvm.ptr<0>
   // CHECK: !llvm.ptr<42>
   "some.op"() : () -> !llvm.ptr<42>
+  // CHECK: !llvm.ptr<ptr<42>, 9>
+  "some.op"() : () -> !llvm.ptr<ptr<42>, 9>
   return
 }
 
 // CHECK-LABEL: @vec
 func.func @vec() {
   // CHECK: vector<4xi32>
   "some.op"() : () -> vector<4xi32>
   // CHECK: vector<4xf32>
   "some.op"() : () -> vector<4xf32>
   // CHECK: !llvm.vec<? x 4 x i32>
   "some.op"() : () -> !llvm.vec<? x 4 x i32>
   // CHECK: !llvm.vec<? x 8 x f16>
   "some.op"() : () -> !llvm.vec<? x 8 x f16>
-  // CHECK: !llvm.vec<4 x ptr<i8>>
-  "some.op"() : () -> !llvm.vec<4 x ptr<i8>>
+  // CHECK: !llvm.vec<4 x ptr>
+  "some.op"() : () -> !llvm.vec<4 x ptr>
   return
 }
 
 // CHECK-LABEL: @array
 func.func @array() {
   // CHECK: !llvm.array<10 x i32>
   "some.op"() : () -> !llvm.array<10 x i32>
   // CHECK: !llvm.array<8 x f32>
   "some.op"() : () -> !llvm.array<8 x f32>
-  // CHECK: !llvm.array<10 x ptr<i32, 4>>
-  "some.op"() : () -> !llvm.array<10 x ptr<i32, 4>>
+  // CHECK: !llvm.array<10 x ptr<4>>
+  "some.op"() : () -> !llvm.array<10 x ptr<4>>
   // CHECK: !llvm.array<10 x array<4 x f32>>
   "some.op"() : () -> !llvm.array<10 x array<4 x f32>>
   return
 }
 
 // CHECK-LABEL: @literal_struct
 func.func @literal_struct() {
   // CHECK: !llvm.struct<()>

 }
 
 // CHECK-LABEL: @identified_struct
 func.func @identified_struct() {
   // CHECK: !llvm.struct<"empty", ()>
   "some.op"() : () -> !llvm.struct<"empty", ()>
   // CHECK: !llvm.struct<"opaque", opaque>
   "some.op"() : () -> !llvm.struct<"opaque", opaque>
-  // CHECK: !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
-  "some.op"() : () -> !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
-  // CHECK: !llvm.struct<"self-recursive", (ptr<struct<"self-recursive">>)>
-  "some.op"() : () -> !llvm.struct<"self-recursive", (ptr<struct<"self-recursive">>)>
+  // CHECK: !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr)>
+  "some.op"() : () -> !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr)>
   // CHECK: !llvm.struct<"unpacked", (i32)>
   "some.op"() : () -> !llvm.struct<"unpacked", (i32)>
   // CHECK: !llvm.struct<"packed", packed (i32)>
   "some.op"() : () -> !llvm.struct<"packed", packed (i32)>
   // CHECK: !llvm.struct<"name with spaces and !^$@$#", packed (i32)>
   "some.op"() : () -> !llvm.struct<"name with spaces and !^$@$#", packed (i32)>
-
-  // CHECK: !llvm.struct<"mutually-a", (ptr<struct<"mutually-b", (ptr<struct<"mutually-a">, 3>)>>)>
-  "some.op"() : () -> !llvm.struct<"mutually-a", (ptr<struct<"mutually-b", (ptr<struct<"mutually-a">, 3>)>>)>
-  // CHECK: !llvm.struct<"mutually-b", (ptr<struct<"mutually-a", (ptr<struct<"mutually-b">>)>, 3>)>
-  "some.op"() : () -> !llvm.struct<"mutually-b", (ptr<struct<"mutually-a", (ptr<struct<"mutually-b">>)>, 3>)>
-  // CHECK: !llvm.struct<"referring-another", (ptr<struct<"unpacked", (i32)>>)>
-  "some.op"() : () -> !llvm.struct<"referring-another", (ptr<struct<"unpacked", (i32)>>)>
-
+  // CHECK: !llvm.struct<"outer", (struct<"nested", ()>)>
+  "some.op"() : () -> !llvm.struct<"outer", (struct<"nested", ()>)>
+  // CHECK: !llvm.struct<"referring-another", (ptr)>
+  "some.op"() : () -> !llvm.struct<"referring-another", (ptr)>
   // CHECK: !llvm.struct<"struct-of-arrays", (array<10 x i32>)>
   "some.op"() : () -> !llvm.struct<"struct-of-arrays", (array<10 x i32>)>
   // CHECK: !llvm.array<10 x struct<"array-of-structs", (i32)>>
   "some.op"() : () -> !llvm.array<10 x struct<"array-of-structs", (i32)>>
-  // CHECK: !llvm.ptr<struct<"ptr-to-struct", (i8)>>
-  "some.op"() : () -> !llvm.ptr<struct<"ptr-to-struct", (i8)>>
   return
 }
 
 func.func @verbose() {
   // CHECK: !llvm.struct<(i64, struct<(f32)>)>
   "some.op"() : () -> !llvm.struct<(i64, !llvm.struct<(f32)>)>
   return
 }
 
-// CHECK-LABEL: @ptr_elem_interface
-// CHECK-COUNT-3: !llvm.ptr<!test.smpla>
-// CHECK: llvm.mlir.undef : !llvm.ptr<!test.smpla>
-func.func @ptr_elem_interface(%arg0: !llvm.ptr<!test.smpla>) {
-  %0 = llvm.load %arg0 : !llvm.ptr<!test.smpla>
-  llvm.store %0, %arg0 : !llvm.ptr<!test.smpla>
-  llvm.mlir.undef : !llvm.ptr<!test.smpla>
-  return
-}
-
 // -----
 
 // Check that type aliases can be used inside LLVM dialect types. Note that
 // currently they are _not_ printed back as this would require
 // DialectAsmPrinter to have a mechanism for querying the presence and
 // usability of an alias outside of its `printType` method.
 
 !baz = i64
 !qux = !llvm.struct<(!baz)>
 
-!rec = !llvm.struct<"a", (ptr<struct<"a">>)>
-
 // CHECK: aliases
 llvm.func @aliases() {
   // CHECK: !llvm.struct<(i32, f32, struct<(i64)>)>
   "some.op"() : () -> !llvm.struct<(i32, f32, !qux)>
-  // CHECK: !llvm.struct<"a", (ptr<struct<"a">>)>
-  "some.op"() : () -> !rec
   llvm.return
 }