[mlir][llvm] Move LLVMPointerType to TypeDef

[Mogball]

Summary

Depends on D136485

[rriddle]

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

125

I'm becoming more and more convinced that this should just be the default...

mlir/lib/Dialect/LLVMIR/IR/LLVMTypeSyntax.cpp

154–155

Can we delete?

mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp

284–285

Can you just inline this into the .td file?

[Mogball]

review comments

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

let llvmBuilder = [{
     llvm::Function *fn = llvm::Intrinsic::getDeclaration(
         module, llvm::Intrinsic::}] # llvmEnumName # [{, {}] #
         !interleave(ListIntSubst<LLVM_IntrPatterns.operand, [0]>.lst, ", ")
         # [{});
     builder.CreateCall(fn, {builder.getInt64(op.getSizeAttr().getInt()),
                             moduleTranslation.lookupValue(op.getPtr())});
   }];
 
-  let assemblyFormat = "$size `,` $ptr attr-dict `:` type($ptr)";
+  let assemblyFormat = "$size `,` $ptr attr-dict `:` qualified(type($ptr))";
 }
 
 def LLVM_LifetimeStartOp : LLVM_LifetimeBaseOp<"lifetime.start"> {
   // Custom builder to convert the size argument to an attribute.
   string mlirBuilder = [{
     $_builder.create<LLVM::LifetimeStartOp>(
       $_location, $_int_attr($size), $ptr);
   }];

mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td

let extraClassDeclaration = [{
     /// Return the llvm.mlir.global operation that defined the value referenced
     /// here.
     GlobalOp getGlobal(SymbolTableCollection &symbolTable);
 
     /// Return the llvm.func operation that is referenced here.
     LLVMFuncOp getFunction(SymbolTableCollection &symbolTable);
   }];
 
-  let assemblyFormat = "$global_name attr-dict `:` type($res)";
+  let assemblyFormat = "$global_name attr-dict `:` qualified(type($res))";
 }
 
 def LLVM_MetadataOp : LLVM_Op<"metadata", [
    NoRegionArguments, SymbolTable, Symbol
 ]> {
   let arguments = (ins
     SymbolNameAttr:$sym_name
   );
let description = [{
 
     // Null pointer to a function with signature void().
     %1 = llvm.mlir.null : !llvm.ptr<func<void ()>>
     ```
   }];
 
   let results = (outs LLVM_AnyPointer:$res);
   let builders = [LLVM_OneResultOpBuilder];
-  let assemblyFormat = "attr-dict `:` type($res)";
+  let assemblyFormat = "attr-dict `:` qualified(type($res))";
 }
 
 def LLVM_UndefOp : LLVM_Op<"mlir.undef", [Pure]>,
                    LLVM_Builder<"$res = llvm::UndefValue::get($_resultType);"> {
   let summary = "Creates an undefined value of LLVM dialect type.";
   let description = [{
     Unlike LLVM IR, MLIR does not have first-class undefined values. Such values
     must be created as SSA values using `llvm.mlir.undef`. This operation has no

mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.h

-//===- LLVMDialect.h - MLIR LLVM dialect types ------------------*- C++ -*-===//
+//===- LLVMTypes.h - MLIR LLVM dialect types --------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the types for the LLVM dialect in MLIR. These MLIR types

 DEFINE_TRIVIAL_LLVM_TYPE(LLVMX86MMXType);
 DEFINE_TRIVIAL_LLVM_TYPE(LLVMTokenType);
 DEFINE_TRIVIAL_LLVM_TYPE(LLVMLabelType);
 DEFINE_TRIVIAL_LLVM_TYPE(LLVMMetadataType);
 
 #undef DEFINE_TRIVIAL_LLVM_TYPE
 
 //===----------------------------------------------------------------------===//
-// LLVMPointerType.
-//===----------------------------------------------------------------------===//
-
-/// LLVM dialect pointer type. This type typically represents a reference to an
-/// object in memory. Pointers may be opaque or parameterized by the element
-/// type. Both opaque and non-opaque pointers are additionally parameterized by
-/// the address space.
-class LLVMPointerType
-    : public Type::TypeBase<
-          LLVMPointerType, Type, detail::LLVMPointerTypeStorage,
-          DataLayoutTypeInterface::Trait, SubElementTypeInterface::Trait> {
-public:
-  /// Inherit base constructors.
-  using Base::Base;
-
-  /// Checks if the given type can have a pointer type pointing to it.
-  static bool isValidElementType(Type type);
-
-  /// Gets or creates an instance of LLVM dialect pointer type pointing to an
-  /// object of `pointee` type in the given address space. The pointer type is
-  /// created in the same context as `pointee`. If the pointee is not provided,
-  /// creates an opaque pointer in the given context and address space.
-  static LLVMPointerType get(MLIRContext *context, unsigned addressSpace = 0);
-  static LLVMPointerType get(Type pointee, unsigned addressSpace = 0);
-  static LLVMPointerType
-  getChecked(function_ref<InFlightDiagnostic()> emitError, Type pointee,
-             unsigned addressSpace = 0);
-  static LLVMPointerType
-  getChecked(function_ref<InFlightDiagnostic()> emitError, MLIRContext *context,
-             unsigned addressSpace = 0);
-
-  /// Returns the pointed-to type. It may be null if the pointer is opaque.
-  Type getElementType() const;
-
-  /// Returns `true` if this type is the opaque pointer type, i.e., it has no
-  /// pointed-to type.
-  bool isOpaque() const;
-
-  /// Returns the address space of the pointer.
-  unsigned getAddressSpace() const;
-
-  /// Verifies that the type about to be constructed is well-formed.
-  static LogicalResult verify(function_ref<InFlightDiagnostic()> emitError,
-                              Type pointee, unsigned);
-  static LogicalResult verify(function_ref<InFlightDiagnostic()> emitError,
-                              MLIRContext *context, unsigned) {
-    return success();
-  }
-
-  /// Hooks for DataLayoutTypeInterface. Should not be called directly. Obtain a
-  /// DataLayout instance and query it instead.
-  unsigned getTypeSizeInBits(const DataLayout &dataLayout,
-                             DataLayoutEntryListRef params) const;
-  unsigned getABIAlignment(const DataLayout &dataLayout,
-                           DataLayoutEntryListRef params) const;
-  unsigned getPreferredAlignment(const DataLayout &dataLayout,
-                                 DataLayoutEntryListRef params) const;
-  bool areCompatible(DataLayoutEntryListRef oldLayout,
-                     DataLayoutEntryListRef newLayout) const;
-  LogicalResult verifyEntries(DataLayoutEntryListRef entries,
-                              Location loc) const;
-
-  void walkImmediateSubElements(function_ref<void(Attribute)> walkAttrsFn,
-                                function_ref<void(Type)> walkTypesFn) const;
-  Type replaceImmediateSubElements(ArrayRef<Attribute> replAttrs,
-                                   ArrayRef<Type> replTypes) const;
-};
-
-//===----------------------------------------------------------------------===//
 // LLVMStructType.
 //===----------------------------------------------------------------------===//
 
 /// LLVM dialect structure type representing a collection of different-typed
 /// elements manipulated together. Structured can optionally be packed, meaning
 /// that their elements immediately follow each other in memory without
 /// accounting for potential alignment.
 ///

mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td

let extraClassDeclaration = [{
     /// Returns the number of arguments to the function.
     unsigned getNumParams() const { return getParams().size(); }
 
     /// Returns `i`-th argument of the function. Asserts on out-of-bounds.
     Type getParamType(unsigned i) { return getParams()[i]; }
   }];
 }
 
+//===----------------------------------------------------------------------===//
+// LLVMPointerType
+//===----------------------------------------------------------------------===//
+
+def LLVMPointerType : LLVMType<"LLVMPointer", "ptr", [
+    DeclareTypeInterfaceMethods<DataLayoutTypeInterface, [
+      "areCompatible", "verifyEntries"]>,
+    DeclareTypeInterfaceMethods<SubElementTypeInterface>]> {
+  let summary = "LLVM pointer type";
+  let description = [{
+    The `!llvm.ptr` type is an LLVM pointer type. This type typically represents
+    a reference to an object in memory. Pointers may be opaque or parameterized
+    by the element type. Both opaque and non-opaque pointers are additionally
+    parameterized by the address space.
+
+    Example:
+
+    ```mlir
+    !llvm.ptr<i8>
+    !llvm.ptr
+    ```
+  }];
+
+  let parameters = (ins DefaultValuedParameter<"Type", "Type()">:$elementType,
+                        DefaultValuedParameter<"unsigned", "0">:$addressSpace);
+  let assemblyFormat = [{
+    (`<` custom<Pointer>($elementType, $addressSpace)^ `>`)?
+  }];
+
+  let genVerifyDecl = 1;
+
+  let builders = [
+    TypeBuilderWithInferredContext<(ins "Type":$elementType,
+                                         CArg<"unsigned", "0">:$addressSpace)>,
+    TypeBuilder<(ins CArg<"unsigned", "0">:$addressSpace), [{
+      return $_get($_ctxt, Type(), addressSpace);
+    }]>
+  ];
+
+  let extraClassDeclaration = [{
+    /// Returns `true` if this type is the opaque pointer type, i.e., it has no
+    /// pointed-to type.
+    bool isOpaque() const { return !getElementType(); }
+
+    /// Checks if the given type can have a pointer type pointing to it.
+    static bool isValidElementType(Type type);
+  }];
+}
+
 #endif // LLVMTYPES_TD

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

def NVVM_VoteBallotOp :
   Arguments<(ins LLVM_Type:$mask, LLVM_Type:$pred)> {
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder,
             llvm::Intrinsic::nvvm_vote_ballot_sync, {$mask, $pred});
   }];
   let hasCustomAssemblyFormat = 1;
 }
 
-def NVVM_SyncWarpOp : 
+def NVVM_SyncWarpOp :
   NVVM_Op<"bar.warp.sync">,
   Arguments<(ins LLVM_Type:$mask)> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_bar_warp_sync, {$mask});
   }];
   let assemblyFormat = "$mask attr-dict `:` type($mask)";
 }
 
bit ret = !cond(
     !and(!eq(geom, "m16n8k8"),
          !ne(c_type, d_type)): false,
 
     // All other are OK.
     true: true
   );
 }
 
-// Returns a list of operation suffixes corresponding to possible b1 
-// multiply-and-accumulate operations for all fragments which have a 
-// b1 type. For all other fragments, the list returned holds a list 
+// Returns a list of operation suffixes corresponding to possible b1
+// multiply-and-accumulate operations for all fragments which have a
+// b1 type. For all other fragments, the list returned holds a list
 // containing the empty string.
 class NVVM_MMA_B1OPS<list<WMMA_REGS> frags> {
   list<string> ret = !cond(
     !eq(frags[0].ptx_elt_type, "b1") : ["xor_popc", "and_popc"],
     true: [""]
   );
 }
 
 /// Generate enum value of the mma.sync intrinsic.
 class MMA_SYNC_NAME<string ALayout, string BLayout, string b1op, int Satfinite,
                WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
   string signature = MMA_SIGNATURE<A, B, C, D>.ret;
   string id = "llvm::Intrinsic::nvvm_mma"
                 # !if(!ne(b1op, ""), "_" # b1op, "")
                 # "_" # A.geom
                 # "_" # ALayout
                 # "_" # BLayout
                 # !if(Satfinite, "_satfinite", "")
-                # signature;  
+                # signature;
 }
 
 /// Helper to create the mapping between the configuration and the mma.sync
 /// intrinsic enum value.
 class MMA_SYNC_INTR {
   list<list<list<list<list<string>>>>> cond0 =
     !foreach(op, NVVM_MMA_OPS.all_mma_sync_ops,
       !foreach(layoutA, ["row", "col"],
         !foreach(layoutB, ["row", "col"],
           !foreach (sat, [0, 1],
             !foreach (b1op, NVVM_MMA_B1OPS<op>.ret,
               !if(NVVM_MMA_SUPPORTED<[op[0], op[1], op[2], op[3]],
                                      layoutA, layoutB, sat>.ret,
       "if (layoutA == \"" # layoutA #  "\" && layoutB == \"" # layoutB #  "\" && "
       "    m == " # op[0].m # " && n == " # op[0].n # " && k == " # op[0].k #
       "    && \"" # op[0].ptx_elt_type # "\" == eltypeA && \""
-       # op[1].ptx_elt_type # "\" == eltypeB && " 
+       # op[1].ptx_elt_type # "\" == eltypeB && "
        # " \"" # op[2].ptx_elt_type # "\" == eltypeC && "
        # " \"" # op[3].ptx_elt_type # "\" == eltypeD "
        # " && (sat.has_value()  ? " # sat # " == static_cast<int>(*sat) : true)"
        # !if(!ne(b1op, ""), " && (b1Op.has_value() ? MMAB1Op::" # b1op # " == b1Op.value() : true)", "") # ")\n"
        # "  return " #
-       MMA_SYNC_NAME<layoutA, layoutB, b1op, sat, op[0], op[1], op[2], op[3]>.id # ";", 
+       MMA_SYNC_NAME<layoutA, layoutB, b1op, sat, op[0], op[1], op[2], op[3]>.id # ";",
           "") // if supported
           ) // b1op
         ) // sat
       ) // layoutB
     ) // layoutA
   ); // all_mma_sync_ops
   list<list<list<string>>> f1 = !foldl([[[""]]],
-                                  !foldl([[[[""]]]], cond0, acc, el, 
+                                  !foldl([[[[""]]]], cond0, acc, el,
                                       !listconcat(acc, el)),
                                     acc1, el1, !listconcat(acc1, el1));
   list<list<string>> f2 = !foldl([[""]], f1, acc1, el1, !listconcat(acc1, el1));
   list<string> f3 = !foldl([""], f2, acc, el, !listconcat(acc, el));
   string id = !foldl("", f3, acc, el, acc # "\n" # el);
 }
 
 def MMALayoutRow : I32EnumAttrCase<"row", 0>;
string baseDescription = [{
 
     ```mlir
     nvvm.wmma.store %0, %1, %2, %3, %4, %5
       {eltype = "f16", k = 16 : i32, layout = "row", m = 16 : i32, n = 16 : i32}
       : !llvm.ptr<i32, 3>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
     ```
   }];
 
-  let assemblyFormat = "$ptr `,` $stride `,` $args attr-dict `:` type($ptr) `,` type($args)";
+  let assemblyFormat = [{
+    $ptr `,` $stride `,` $args attr-dict `:` qualified(type($ptr)) `,`
+    type($args)
+  }];
   let hasVerifier = 1;
 }
 
 // Base class for all the variants of WMMA mmaOps that may be defined.
 def NVVM_WMMAMmaOp : NVVM_Op<"wmma.mma">,
   Results<(outs LLVM_AnyStruct:$res)>,
   Arguments<(ins I32Attr:$m, I32Attr:$n, I32Attr:$k, MMALayoutAttr:$layoutA,
              MMALayoutAttr:$layoutB, MMATypesAttr:$eltypeA,

 }
 
 def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
 
   let summary = "cooperative matrix-multiply and accumulate";
 
   let description = [{
     The `nvvm.mma.sync` operation collectively performs the operation
-    `D = matmul(A, B) + C` using all threads in a warp.    
+    `D = matmul(A, B) + C` using all threads in a warp.
 
     All the threads in the warp must execute the same `mma.sync` operation.
 
     For each possible multiplicand PTX data type, there are one or more possible
     instruction shapes given as "mMnNkK". The below table describes the posssibilities
-    as well as the types required for the operands. Note that the data type for 
-    C (the accumulator) and D (the result) can vary independently when there are 
+    as well as the types required for the operands. Note that the data type for
+    C (the accumulator) and D (the result) can vary independently when there are
     multiple possibilities in the "C/D Type" column.
 
     When an optional attribute cannot be immediately inferred from the types of
     the operands and the result during parsing or validation, an error will be
     raised.
 
-    `b1Op` is only relevant when the binary (b1) type is given to 
+    `b1Op` is only relevant when the binary (b1) type is given to
     `multiplicandDataType`. It specifies how the multiply-and-acumulate is
     performed and is either `xor_popc` or `and_poc`. The default is `xor_popc`.
 
-    `intOverflowBehavior` is only relevant when the `multiplicandType` attribute 
+    `intOverflowBehavior` is only relevant when the `multiplicandType` attribute
     is one of `u8, s8, u4, s4`, this attribute describes how overflow is handled
     in the accumulator. When the attribute is `satfinite`, the accumulator values
     are clamped in the int32 range on overflow. This is the default behavior.
-    Alternatively, accumulator behavior `wrapped` can also be specified, in 
+    Alternatively, accumulator behavior `wrapped` can also be specified, in
     which case overflow wraps from one end of the range to the other.
 
-    `layoutA` and `layoutB` are required and should generally be set to 
+    `layoutA` and `layoutB` are required and should generally be set to
     `#nvvm.mma_layout<row>` and `#nvvm.mma_layout<col>` respectively, but other
     combinations are possible for certain layouts according to the table below.
 
     ```
     | A/B Type | Shape     | ALayout | BLayout | A Type   | B Type   | C/D Type          |
     |----------|-----------|---------|---------|----------|----------|-------------------|
     | f64      | .m8n8k4   | row     | col     | 1x f64   | 1x f64   | 2x f64            |
     | f16      | .m8n8k4   | row/col | row/col | 2x f16x2 | 2x f16x2 | 4x f16x2 or 8xf32 |
let description = [{
     | b1       | m8n8k128  | row     | col     | 1x i32   | 1x i32   | 2x i32            |
     |          | m16n8k128 | row     | col     | 2x i32   | 1x i32   | 4x i32            |
     ```
 
 
     Example:
     ```mlir
 
-    %128 = nvvm.mma.sync A[%120, %121, %122, %123]  
-                         B[%124, %125]  
-                         C[%126, %127]  
-                         {layoutA = #nvvm.mma_layout<row>, 
-                          layoutB = #nvvm.mma_layout<col>, 
-                          shape = {k = 16 : i32, m = 16 : i32, n = 8 : i32}} 
+    %128 = nvvm.mma.sync A[%120, %121, %122, %123]
+                         B[%124, %125]
+                         C[%126, %127]
+                         {layoutA = #nvvm.mma_layout<row>,
+                          layoutB = #nvvm.mma_layout<col>,
+                          shape = {k = 16 : i32, m = 16 : i32, n = 8 : i32}}
         : (vector<2xf16>, vector<2xf16>, vector<2xf16>)
            -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
     ```
   }];
 
   let results = (outs LLVM_AnyStruct:$res);
   let arguments = (ins NVVM_MMAShapeAttr:$shape,
-             OptionalAttr<MMAB1OpAttr>:$b1Op, 
+             OptionalAttr<MMAB1OpAttr>:$b1Op,
              OptionalAttr<MMAIntOverflowAttr>:$intOverflowBehavior,
              MMALayoutAttr:$layoutA,
              MMALayoutAttr:$layoutB,
              OptionalAttr<MMATypesAttr>:$multiplicandAPtxType,
              OptionalAttr<MMATypesAttr>:$multiplicandBPtxType,
              Variadic<LLVM_Type>:$operandA,
              Variadic<LLVM_Type>:$operandB,
-             Variadic<LLVM_Type>:$operandC); 
+             Variadic<LLVM_Type>:$operandC);
 
   let extraClassDeclaration = !strconcat([{
       static llvm::Intrinsic::ID getIntrinsicID(
             int64_t m, int64_t n, uint64_t k,
-            llvm::Optional<MMAB1Op> b1Op, 
+            llvm::Optional<MMAB1Op> b1Op,
             llvm::Optional<MMAIntOverflow> sat,
             mlir::NVVM::MMALayout layoutAEnum, mlir::NVVM::MMALayout layoutBEnum,
             mlir::NVVM::MMATypes eltypeAEnum, mlir::NVVM::MMATypes eltypeBEnum,
             mlir::NVVM::MMATypes eltypeCEnum, mlir::NVVM::MMATypes eltypeDEnum) {
         llvm::StringRef layoutA = stringifyEnum(layoutAEnum);
         llvm::StringRef layoutB = stringifyEnum(layoutBEnum);
         llvm::StringRef eltypeA = stringifyEnum(eltypeAEnum);
         llvm::StringRef eltypeB = stringifyEnum(eltypeBEnum);
         llvm::StringRef eltypeC = stringifyEnum(eltypeCEnum);
         llvm::StringRef eltypeD = stringifyEnum(eltypeDEnum);
         }],
         MMA_SYNC_INTR<>.id, [{
         return 0;
       }
 
       static Optional<mlir::NVVM::MMATypes> inferOperandMMAType(Type operandElType,
         bool isAccumulator);
 
       MMATypes accumPtxType();
       MMATypes resultPtxType();
     }]);
 
   let builders = [
-      OpBuilder<(ins  "Type":$resultType, "ValueRange":$operandA, 
+      OpBuilder<(ins  "Type":$resultType, "ValueRange":$operandA,
         "ValueRange":$operandB, "ValueRange":$operandC,
         "ArrayRef<int64_t>":$shape, "Optional<MMAB1Op>":$b1Op,
         "Optional<MMAIntOverflow>":$intOverflow,
         "Optional<std::array<MMATypes, 2>>":$multiplicandPtxTypes,
         "Optional<std::array<MMALayout, 2>>":$multiplicandLayouts)>
     ];
 
   string llvmBuilder = [{
     auto operands = moduleTranslation.lookupValues(opInst.getOperands());
     auto intId = mlir::NVVM::MmaOp::getIntrinsicID(
-        $shape.getM(), $shape.getN(), $shape.getK(), 
+        $shape.getM(), $shape.getN(), $shape.getK(),
         $b1Op, $intOverflowBehavior,
         $layoutA, $layoutB,
-        $multiplicandAPtxType.value(), 
+        $multiplicandAPtxType.value(),
         $multiplicandBPtxType.value(),
-        op.accumPtxType(), 
+        op.accumPtxType(),
         op.resultPtxType());
 
     $res = createIntrinsicCall(
       builder, intId, operands);
   }];
 
   let hasCustomAssemblyFormat = 1;
   let hasVerifier = 1;
 }
 
 #endif // NVVMIR_OPS

mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp

void LLVMDialect::initialize() {
 
   // clang-format off
   addTypes<LLVMVoidType,
            LLVMPPCFP128Type,
            LLVMX86MMXType,
            LLVMTokenType,
            LLVMLabelType,
            LLVMMetadataType,
-           LLVMPointerType,
            LLVMFixedVectorType,
            LLVMScalableVectorType,
            LLVMStructType>();
   // clang-format on
   registerTypes();
 
   addOperations<
 #define GET_OP_LIST

mlir/lib/Dialect/LLVMIR/IR/LLVMTypeSyntax.cpp

 void mlir::LLVM::detail::printType(Type type, AsmPrinter &printer) {
   if (!type) {
     printer << "<<NULL-TYPE>>";
     return;
   }
 
   printer << getTypeKeyword(type);
 
-  if (auto ptrType = type.dyn_cast<LLVMPointerType>()) {
-    if (ptrType.isOpaque()) {
-      if (ptrType.getAddressSpace() != 0)
-        printer << '<' << ptrType.getAddressSpace() << '>';
-      return;
-    }
-
-    printer << '<';
-    dispatchPrint(printer, ptrType.getElementType());
-    if (ptrType.getAddressSpace() != 0)
-      printer << ", " << ptrType.getAddressSpace();
-    printer << '>';
-    return;
-  }
+  if (auto ptrType = type.dyn_cast<LLVMPointerType>())
+    return ptrType.print(printer);
 
   if (auto arrayType = type.dyn_cast<LLVMArrayType>())
     return arrayType.print(printer);
   if (auto vectorType = type.dyn_cast<LLVMFixedVectorType>())
     return printVectorType(printer, vectorType);
 
   if (auto vectorType = type.dyn_cast<LLVMScalableVectorType>()) {
     printer << "<? x " << vectorType.getMinNumElements() << " x ";
if (auto funcType = type.dyn_cast<LLVMFunctionType>())
     return funcType.print(printer);
 }
 
 //===----------------------------------------------------------------------===//
 // Parsing.
 //===----------------------------------------------------------------------===//
 
 static ParseResult dispatchParse(AsmParser &parser, Type &type);
 
-/// Parses an LLVM dialect pointer type.
-///   llvm-type ::= `ptr<` llvm-type (`,` integer)? `>`
-///               | `ptr` (`<` integer `>`)?
-static LLVMPointerType parsePointerType(AsmParser &parser) {
-  SMLoc loc = parser.getCurrentLocation();
-  Type elementType;
-  if (parser.parseOptionalLess()) {
-    return parser.getChecked<LLVMPointerType>(loc, parser.getContext(),
-                                              /*addressSpace=*/0);
-  }
-
-  unsigned addressSpace = 0;
-  OptionalParseResult opr = parser.parseOptionalInteger(addressSpace);
-  if (opr.has_value()) {
-    if (failed(*opr) || parser.parseGreater())
-      return LLVMPointerType();
-    return parser.getChecked<LLVMPointerType>(loc, parser.getContext(),
-                                              addressSpace);
-  }
-
-  if (dispatchParse(parser, elementType))
-    return LLVMPointerType();
-
-  if (succeeded(parser.parseOptionalComma()) &&
-      failed(parser.parseInteger(addressSpace)))
-    return LLVMPointerType();
-  if (failed(parser.parseGreater()))
-    return LLVMPointerType();
-  return parser.getChecked<LLVMPointerType>(loc, elementType, addressSpace);
-}
-
 /// Parses an LLVM dialect vector type.
 ///   llvm-type ::= `vec<` `? x`? integer `x` llvm-type `>`
 /// Supports both fixed and scalable vectors.
 static Type parseVectorType(AsmParser &parser) {
   SmallVector<int64_t, 2> dims;
   SMLoc dimPos, typePos;
   Type elementType;
   SMLoc loc = parser.getCurrentLocation();
   if (parser.parseLess() || parser.getCurrentLocation(&dimPos) ||
static Type dispatchParse(AsmParser &parser, bool allowAny = true) {
   return StringSwitch<function_ref<Type()>>(key)
       .Case("void", [&] { return LLVMVoidType::get(ctx); })
       .Case("ppc_fp128", [&] { return LLVMPPCFP128Type::get(ctx); })
       .Case("x86_mmx", [&] { return LLVMX86MMXType::get(ctx); })
       .Case("token", [&] { return LLVMTokenType::get(ctx); })
       .Case("label", [&] { return LLVMLabelType::get(ctx); })
       .Case("metadata", [&] { return LLVMMetadataType::get(ctx); })
       .Case("func", [&] { return LLVMFunctionType::parse(parser); })
-      .Case("ptr", [&] { return parsePointerType(parser); })
+      .Case("ptr", [&] { return LLVMPointerType::parse(parser); })
       .Case("vec", [&] { return parseVectorType(parser); })
       .Case("array", [&] { return LLVMArrayType::parse(parser); })
       .Case("struct", [&] { return parseStructType(parser); })
       .Default([&] {
         parser.emitError(keyLoc) << "unknown LLVM type: " << key;
         return Type();
       })();
 }

mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp

+//===- LLVMTypes.cpp - MLIR LLVM dialect types ------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the types for the LLVM dialect in MLIR. These MLIR types
if (isVarArg) {
     if (!params.empty())
       p << ", ";
     p << "...";
   }
   p << ')';
 }
 
 //===----------------------------------------------------------------------===//
+// custom<Pointer>
+//===----------------------------------------------------------------------===//
+
+static ParseResult parsePointer(AsmParser &p, FailureOr<Type> &elementType,
+                                FailureOr<unsigned> &addressSpace) {
+  addressSpace = 0;
+  // `<` addressSpace `>`
+  OptionalParseResult result = p.parseOptionalInteger(*addressSpace);
+  if (result.has_value()) {
+    if (failed(result.value()))
+      return failure();
+    elementType = Type();
+    return success();
+  }
+
+  if (parsePrettyLLVMType(p, elementType))
+    return failure();
+  if (succeeded(p.parseOptionalComma()))
+    return p.parseInteger(*addressSpace);
+
+  return success();
+}
+
+static void printPointer(AsmPrinter &p, Type elementType,
+                         unsigned addressSpace) {
+  if (elementType)
+    printPrettyLLVMType(p, elementType);
+  if (addressSpace != 0) {
+    if (elementType)
+      p << ", ";
+    p << addressSpace;
+  }
+}
+
+//===----------------------------------------------------------------------===//
 // ODS-Generated Definitions
 //===----------------------------------------------------------------------===//
 
 /// These are unused for now.
 /// TODO: Move over to these once more types have been migrated to TypeDef.
 LLVM_ATTRIBUTE_UNUSED static OptionalParseResult
 generatedTypeParser(AsmParser &parser, StringRef *mnemonic, Type &value);
 LLVM_ATTRIBUTE_UNUSED static LogicalResult

 }
 
 Type LLVMFunctionType::replaceImmediateSubElements(
     ArrayRef<Attribute> replAttrs, ArrayRef<Type> replTypes) const {
   return get(replTypes.front(), replTypes.drop_front(), isVarArg());
 }
 
 //===----------------------------------------------------------------------===//
-// Pointer type.
+// LLVMPointerType
 //===----------------------------------------------------------------------===//
 
 bool LLVMPointerType::isValidElementType(Type type) {
   if (!type)
     return true;
   return isCompatibleOuterType(type)
              ? !type.isa<LLVMVoidType, LLVMTokenType, LLVMMetadataType,
                          LLVMLabelType>()
              : type.isa<PointerElementTypeInterface>();
 }
 
 LLVMPointerType LLVMPointerType::get(Type pointee, unsigned addressSpace) {
   assert(pointee && "expected non-null subtype, pass the context instead if "
                     "the opaque pointer type is desired");
   return Base::get(pointee.getContext(), pointee, addressSpace);
 }
 
-LLVMPointerType LLVMPointerType::get(MLIRContext *context,
-                                     unsigned addressSpace) {
-  return Base::get(context, Type(), addressSpace);
-}
-
-LLVMPointerType
-LLVMPointerType::getChecked(function_ref<InFlightDiagnostic()> emitError,
-                            Type pointee, unsigned addressSpace) {
-  return Base::getChecked(emitError, pointee.getContext(), pointee,
-                          addressSpace);
-}
-
-LLVMPointerType
-LLVMPointerType::getChecked(function_ref<InFlightDiagnostic()> emitError,
-                            MLIRContext *context, unsigned addressSpace) {
-  return Base::getChecked(emitError, context, Type(), addressSpace);
-}
-
-Type LLVMPointerType::getElementType() const { return getImpl()->pointeeType; }
-
-bool LLVMPointerType::isOpaque() const { return !getImpl()->pointeeType; }
-
-unsigned LLVMPointerType::getAddressSpace() const {
-  return getImpl()->addressSpace;
-}
-
 LogicalResult
 LLVMPointerType::verify(function_ref<InFlightDiagnostic()> emitError,
                         Type pointee, unsigned) {
   if (!isValidElementType(pointee))
     return emitError() << "invalid pointer element type: " << pointee;
   return success();
 }
 
+//===----------------------------------------------------------------------===//
+// DataLayoutTypeInterface
+
 constexpr const static unsigned kDefaultPointerSizeBits = 64;
 constexpr const static unsigned kDefaultPointerAlignment = 8;
 
 Optional<unsigned> mlir::LLVM::extractPointerSpecValue(Attribute attr,
                                                        PtrDLEntryPos pos) {
   auto spec = attr.cast<DenseIntElementsAttr>();
   auto idx = static_cast<unsigned>(pos);
   if (idx >= spec.size())
if (extractPointerSpecValue(values, PtrDLEntryPos::Abi) >
         extractPointerSpecValue(values, PtrDLEntryPos::Preferred)) {
       return emitError(loc) << "preferred alignment is expected to be at least "
                                "as large as ABI alignment";
     }
   }
   return success();
 }
 
+//===----------------------------------------------------------------------===//
+// SubElementTypeInterface
+
 void LLVMPointerType::walkImmediateSubElements(
     function_ref<void(Attribute)> walkAttrsFn,
     function_ref<void(Type)> walkTypesFn) const {
   walkTypesFn(getElementType());
 }
 
 Type LLVMPointerType::replaceImmediateSubElements(
     ArrayRef<Attribute> replAttrs, ArrayRef<Type> replTypes) const {

mlir/lib/Dialect/LLVMIR/IR/TypeDetail.h

private:
   unsigned keySizeAndFlags = 0;
 
   /// Number of the types contained in an identified struct combined with
   /// mutable flags. Must only be used through the Mutable* bitfields.
   unsigned identifiedBodySizeAndFlags = 0;
 };
 
 //===----------------------------------------------------------------------===//
-// LLVMPointerTypeStorage.
-//===----------------------------------------------------------------------===//
-
-/// Storage type for LLVM dialect pointer types. These are uniqued by a pair of
-/// element type and address space. The element type may be null indicating that
-/// the pointer is opaque.
-struct LLVMPointerTypeStorage : public TypeStorage {
-  using KeyTy = std::tuple<Type, unsigned>;
-
-  LLVMPointerTypeStorage(const KeyTy &key)
-      : pointeeType(std::get<0>(key)), addressSpace(std::get<1>(key)) {}
-
-  static LLVMPointerTypeStorage *construct(TypeStorageAllocator &allocator,
-                                           const KeyTy &key) {
-    return new (allocator.allocate<LLVMPointerTypeStorage>())
-        LLVMPointerTypeStorage(key);
-  }
-
-  bool operator==(const KeyTy &key) const {
-    return std::make_tuple(pointeeType, addressSpace) == key;
-  }
-
-  Type pointeeType;
-  unsigned addressSpace;
-};
-
-//===----------------------------------------------------------------------===//
 // LLVMTypeAndSizeStorage.
 //===----------------------------------------------------------------------===//
 
 /// Common storage used for LLVM dialect types that need an element type and a
 /// number: arrays, fixed and scalable vectors. The actual semantics of the
 /// type is defined by its kind.
 struct LLVMTypeAndSizeStorage : public TypeStorage {
   using KeyTy = std::tuple<Type, unsigned>;

mlir/test/Dialect/LLVMIR/invalid.mlir

func.func @invalid_vector_type_5(%a : vector<4xf32>, %idx : i32) -> vector<4xf32> {
   // expected-error@+1 {{failed to verify that argument type matches vector element type}}
   %b = "llvm.insertelement"(%a, %a, %idx) : (vector<4xf32>, vector<4xf32>, i32) -> vector<4xf32>
   return %b : vector<4xf32>
 }
 
 // -----
 
 func.func @null_non_llvm_type() {
-  // expected-error@+1 {{custom op 'llvm.mlir.null' invalid kind of type specified}}
+  // expected-error@+1 {{'llvm.mlir.null' op result #0 must be LLVM pointer type, but got 'i32'}}
   llvm.mlir.null : i32
 }
 
 // -----
 
 func.func @nvvm_invalid_shfl_pred_1(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i32) {
   // expected-error@+1 {{expected return type to be a two-element struct with i1 as the second element}}
   %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> i32

mlir/test/Dialect/LLVMIR/layout.mlir

func.func @spec() {
     // CHECK: size = 4
     "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
     // CHECK: alignment = 8
     // CHECK: bitsize = 64
     // CHECK: preferred = 8
     // CHECK: size = 8
     "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
     // CHECK: alignment = 4
-  // CHECK: bitsize = 32
+    // CHECK: bitsize = 32
     // CHECK: preferred = 8
     // CHECK: size = 4
     "test.data_layout_query"() : () -> !llvm.ptr<3>
     // CHECK: alignment = 8
-  // CHECK: bitsize = 32
+    // CHECK: bitsize = 32
     // CHECK: preferred = 8
     // CHECK: size = 4
   "test.data_layout_query"() : () -> !llvm.ptr<4>
     return
   }
 }
 
 // -----